北京极感科技一面
面试岗位是高性能计算，包括Opencv的算子，AI推理加速 1.自我介绍
-
实习经历，做了哪些事情，你怎么用CUDA加速的，怎么才能来写出效率更高的核函数，用到了哪些进一步的优化呢，使用共享内存的时候要注意哪些问题，怎么把HWC排列的图像转换为一维呢，问了一下CUDA编程方面的知识
-
考察八股，C++、操作系统等方面的知识
（1）说一下static关键字是干什么的
（2）线程和进程的区别，局部性原理是什么
答案：进程是资源分配的基本单位。每个进程拥有独立的地址空间、文件描述符、环境变量、信号处理等资源。
线程：是CPU调度和执行的基本单位，一个进程可以包含多个线程，这些线程共享进程的资源，但各自拥有独立的栈、寄存器状态和程序计数器。
| 进程 | 线程 | |
|---|---|---|
| 地址空间 | 独立 | 共享所属进程的地址空间 |
| 全局变量/堆 | 不共享 | 共享 |
| 栈 | 各自有独立栈 | 各自有独立栈 |
| 文件描述符 | 独立 | 共享 |
-
线程创建与切换的开销比进程小，进程切换需要刷新TLB、切换地址空间、代价高；线程切换只需要保存寄存器和指针，更快。高并发场景下更适合使用多线程。
-
进程之间完全隔离，一个进程的崩溃不会影响其他进程，线程之间无隔离，一个线程访问非法内存或者死锁，可能导致整个进程崩溃。
-
通信方式：进程间通信需要通过管道、消息队列、共享内存、Socket等显式机制； 线程间通信是通过读写共享变量，配合互斥锁、条件变量等同步原语即可。
（3）vector是浅拷贝还是深拷贝，怎么禁止拷贝构造函数
答：深拷贝和浅拷贝是对象复制时的两种不同策略，核心区别在于，是否复制对象所指向的底层资源
- 浅拷贝：
定义：只复制对象的直接成员变量，对于指针或者引用类型的成员，仅复制地址，而不复制指针所指向的数据。
结果：原对象和拷贝对象内存共享一块堆内存或其他资源
风险：修改一方会影响另一方，当两个对象析构时，可能重复释放同一块内存，导致程序崩溃，C++的默认拷贝构造函数和复制运算就是浅拷贝。
- 深拷贝：
定义：不仅复制对象的成员值，还为指针成员分配新的内存，并将原数据完整的复制过去。
结果：原对象和拷贝对象拥有各自独立的资源，互不影响。
安全：各自管理自己的内存，析构时不会冲突。
| 特性 | 浅拷贝 | 深拷贝 |
|---|---|---|
| 指针成员处理 | 复制地址 | 复制地址+复制指向的数据 |
| 内存是否独立 | 否，内存共享 | 是，各自拥有独立的空间 |
| 默认行为 | C++默认拷贝是浅拷贝 | 需要动手实现 |
| 安全性 | 低 | 高 |
| 性能开销 | 小 | 大，需要额外内存分配和复制 |
浅拷贝是是表面复制，共享资源；深拷贝是彻底克隆，资源独立。C++默认浅拷贝，管理资源时必须显示实现深拷贝或者禁止浅拷贝。
vector默认执行的是深拷贝，当元素的是值类型的时候，自定义类且正确实现拷贝语义时，vector会为每个元素调用其拷贝构造函数，新的vector拥有完全独立的副本，修改原来的vector不会影响拷贝后的对象，用delete关键字来禁止拷贝构造函数。
（4）CUDA编程中为什么线程分化会降低性能，block的大小是怎么影响性能的
答：CUDA编程中，线程以warp为单位调度（1个warp = 32个连续的线程），同一个warp内的所有线程，在同一时间内必须执行相同的指令，当同一个warp中的线程因为条件分支，走向了不同路径的时候，就发生了线程分化。
__global__ void kernel(float* data) {
int idx = threadIdx.x;
if (idx % 2 == 0) {
data[idx] *= 2; // 偶数线程走这里
} else {
data[idx] += 1; // 奇数线程走这里
}
}
对此，GPU处理的分化方法是，“串行执行所有分支路径”，先让满足if条件的线程执行，其他线程停顿，然后再让满足else的线程执行，之前的线程停顿。最后所有的线程在}处汇合。导致本来一个周期能完成的操作，编程两个周期，只有warp内部分化影响性能，不同warps之间的分支互不影响。
避免或者减轻分化的方法：尽量让warp内部线程走相同的路径：warp 对齐数据（32 的倍数）；避免基于 threadIdx.x 的随机分支；使用 predication（谓词执行）：对于简单短分支，编译器可能自动转换为条件赋值（无分化）；重构算法：如将控制流转为数据流（查表、掩码操作）
Block大小，直接影响GPU资源利用率和并行度，需要权衡多个因素 1.warp对齐与分化：block大小最好是32的倍数，否则最后一个warp不满，浪费计算资源。2.占有率：占有率=实际活跃的warps数/SM最大支持warps数，高 occupancy 可隐藏内存延迟（当一个 warp 等待内存时，切换到另一个 warp 执行）3.资源竞争与并行粒度，block太小：并行度不足，无法充分利用数据，启动开销占比高；block太大：单个block执行时间长，负载不均衡，更容易受到单个warp分化影响整体性能
（5）CPU流水线并行是什么，分支预测是什么
（6）解释一下什么是虚拟内存
答案：虚拟内存是现代操作系统提供的一种内存管理机制，它的核心思想是：为每个进程提供一个独立、连续、私有的地址空间（称为虚拟地址空间），使得程序在运行时无需关心物理内存的实际布局和大小。关键点在于地址抽象与隔离，扩展可用内存，按需加载与内存保护，虚拟内存通过硬件（MMU、TLB）和软件（操作系统页表管理）协同，实现了内存的抽象、隔离、扩展和保护，是现代多任务操作系统的基础之一。
（7）智能指针
（8）介绍一下抽象类
答：C++中的抽象类，是一种不能被实例化的类，它至少包含一个纯虚函数，抽象类的主要作用作为一组相关类定义统一的接口
//纯虚函数通过 = 0 声明，没有函数体（或即使有也不能直接调用）
class Shape {
public:
virtual void draw() = 0; // 纯虚函数
virtual ~Shape() = default; // 建议提供虚析构函数
};
//不能创建对象，不能实例化
Shape s; // ❌ 编译错误：不能实例化抽象类
Shape* p = new Shape(); // ❌ 同样错误
C++没有单独的interface关键字，抽象类承担了接口角色，如果一个类只有一个纯虚函数（无成员变量，无非虚函数） ，就相当于一个接口。
抽象类是包含至少一个纯虚函数的类，用于定义接口规范、支持多态，并禁止直接实例化；它是实现面向对象“开闭原则”和“依赖倒置”的重要工具。
(9) 内联函数的优点和缺点，写了内联关键字一定能百分百内联成功吗，怎么判断函数是否内联成功
答：内联函数是C++中一种通过inline关键字建议编译器将函数调用改为展开函数体的代码机制，目的是为了消除普通函数调用的开销（如压栈，跳转，返回等），在编译阶段，将函数调用处直接替换为函数体代码，类似于宏展开，但具有类型检查和作用域安全。
inline int max(int a, int b) {
return a > b ? a : b;
}
int x = max(3, 5); // 编译后可能变为：int x = (3 > 5 ? 3 : 5);
内联函数的优点：减少函数调用的开销，提升CPU流水线效率（避免分支跳转打断流水线），便于编译器优化 内联函数的缺点：代码膨胀，指令缓存命中率下降；更多的内存占用，在极端环境下反而降低性能；增加编译的时间和链接复杂度；不适用于负责函数：递归，大函数，含循环或者异常处理的函数通常不适合内联。
写了inline关键字不一定百分百内联，最终是否内联有编译器更具优化策略决定。以下情况可能内联失败
函数体过大或者逻辑复杂；包含递归调用，包含静态局部变量，调用频率极低，编译时未开启优化
判断是否内联成功的方法如下：
查看汇编代码；使用编译器诊断选项；性能分析
（10）CUDA编程中使用共享内存时要注意哪些事情
| 特性 | 说明 |
|---|---|
| 作用域 | 仅限于同一个线程块内的线程共享 |
| 生命周期 | 从 block 启动到 block 结束 |
| 大小限制 | 通常每个 SM 有 48KB~96KB（取决于 GPU 架构），需多个 blocks 共享 |
| 访问速度 | 几乎与寄存器相当（远快于全局内存） |
| 同步要求 | 写入后需用 __syncthreads() 同步，确保所有线程可见 |
必须显式声明为__shared__;需要线程同步__synthreads();避免banck conflict；不要超出共享内存容量；初始化共享内存需要谨慎；
共享内存的应用场景包括：分块矩阵乘法，归约，卷积计算，排序扫描，可以使用Nsight Compute来查看是否发生banck conflict
4.手撕代码，题目并不很难，然后还有开放性思维问题
