长沙诚拙微电子面经
面试官人很和蔼,其中一个还是校友,但我还是太菜了,感觉实习这么久后期有点赛博打螺丝了,稍微吐槽一下目前实习的公司(合肥半导体公司),我都实习差不多8个月了,哪怕签三方也要一直实习到毕业,这一点真的有点坑了,相当于一直拿实习工资干活,还有点耽误我秋招,所以我一直没有背八股文,长沙这家面试还要下班了才有时间面试。
1.自我介绍,简要介绍项目经验以及工作内容
2.工作中内容提到的ISP图像处理使用CUDA加速,那图像捕获帧处理帧是怎么实现同步的,怎么处理丢帧问题的(因为我其实不负责这块,所以回答的不是很好,我只知道消费者生产者问题,然后说了可以用信号量实现同步,这里可以联系到线程同步互斥的相关知识,可以使用互斥锁,条件变量、信号量等)
3.在工作中你是用什么标准来衡量一个核函数的好坏呢?(我回答的是主要更具时间来判断,这一点确实我回答的太片面了,公司也没有人对我有这方面的指导,而且公司就主要关注时间开销,我们组就我一个人干这个方向,感觉很心累,CUDA都是进公司现学的,但其实有很多标准来衡量核函数的性能,包括执行时间,硬件资源利用、内存访问、计算效率、并行度)
4.在矩阵转置中,怎么用共享内存实现,转置过程中,必然有线程索引不连续,那你怎么解决,还能合并访存吗?
如下代码:
// 错误做法:直接转置 __global__ void bad_transpose(float* A, float* B, int N) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x < N && y < N) { B[y * N + x] = A[x * N + y]; // ❌ 写入B时,y变化,x固定 → 列访问 → 非合并! } }
矩阵在内存中是行优先存储,这里B或者A肯定有一个不是按连续的地址访问的,假设N=10,x=0,y=0,1,2,3,4....9。那就会出现
B[0] = A[0],B[10] = A[1],B[20] = A[2].......
这里我们就没有连续的访问B,要注意到在一个wrap中threadIdx.y固定,threadIdx.x连续
- 读取
A[x][y]
:是合并的(行优先存储,x连续)。 - 写入
B[y][x]
:是非合并的!因为多个线程的y
不同,x
相同,访问的是B
的同一列,地址不连续。
这会导致全局内存写入性能极差。
解决方案就是使用共享内存分块转置
核心思想:
- 分块(Tile):将大矩阵分成小块(如 32x32)。
- 协作加载:一个线程块负责一个 tile。
- 读取时合并:从全局内存合并地读取一个 tile 到共享内存。
- 在共享内存中转置:线程将数据“转置”地写入共享内存。
- 写回时合并:从共享内存合并地写回全局内存
关键:共享内存没有合并访存的概念!它是一个低延迟的片上内存,我们可以在其中任意顺序读写。真正的合并访存优化发生在全局内存的读和写。
#define TILE_SIZE 32 __global__ void tiled_transpose(float* A, float* B, int N) { // 共享内存:用于存储一个 tile __shared__ float tile[TILE_SIZE][TILE_SIZE]; int x = blockIdx.x * TILE_SIZE + threadIdx.x; int y = blockIdx.y * TILE_SIZE + threadIdx.y; // Step 1: 从A中合并地读取一个 tile 到共享内存 // 注意:我们按“行”读取,所以是合并访存 if (x < N && y < N) { tile[threadIdx.y][threadIdx.x] = A[y * N + x]; } // 【同步】确保所有线程都完成了加载 __syncthreads(); // Step 2: 计算转置后的全局坐标 int tx = blockIdx.y * TILE_SIZE + threadIdx.x; // 注意:blockIdx.y → x方向 int ty = blockIdx.x * TILE_SIZE + threadIdx.y; // 注意:blockIdx.x → y方向 // Step 3: 从共享内存中“转置”地读取,并合并地写回B // 现在我们写入 B[ty][tx],而 ty 是连续变化的(因为 threadIdx.y 连续) if (tx < N && ty < N) { B[ty * N + tx] = tile[threadIdx.x][threadIdx.y]; // 注意索引交换 } // 注意:这里不需要再次 __syncthreads(),因为写回后没有后续操作 }
1. 读取阶段(从 A 读取)
- 线程
(tx, ty)
读取A[y * N + x]
- 在一个 warp 中(
threadIdx.y
固定,threadIdx.x
连续 0~31):访问的地址:y*N + 0, y*N + 1, ..., y*N + 31连续地址 → 完美合并访存 ✅
2. 写回阶段(写入 B)
- 线程
(tx, ty)
写入B[ty * N + tx]
- 在一个 warp 中(
threadIdx.y
固定,threadIdx.x
连续):tx 连续增加(0,1,2,...)ty 固定(因为 threadIdx.y 固定)访问的地址:ty*N + 0, ty*N + 1, ..., ty*N + 31连续地址 → 完美合并访存。
4.C++了解多少,什么是多态,什么是C++缺省,怎么用C语言实现
c++多态是指通同一个接口,多种实现。即用父类的指针或者引用,去调用子类的函数,程序在运行时自动决定该调用哪个版本的函数。实现方式:虚函数(virtual)+继承。
#include <iostream> // 基类(抽象接口) class Animal { public: // 声明为虚函数,表示可以被子类重写 virtual void makeSound() const { std::cout << "Animal makes a sound" << std::endl; } // 虚析构函数很重要!确保正确释放子类对象 virtual ~Animal() = default; }; // 子类 Dog class Dog : public Animal { public: void makeSound() const override { // override 明确表示重写 std::cout << "Woof! Woof!" << std::endl; } }; // 子类 Cat class Cat : public Animal { public: void makeSound() const override { std::cout << "Meow~ Meow~" << std::endl; } }; int main() { Dog dog; Cat cat; // 使用基类指针指向不同子类对象 Animal* ptr1 = &dog; Animal* ptr2 = &cat; // 调用同一个接口,但执行不同行为 ptr1->makeSound(); // 输出: Woof! Woof! ptr2->makeSound(); // 输出: Meow~ Meow~ return 0; }
C++缺省参数,在C语言中可以用宏定义实现,或者使用结构体(struct),把所有参数打包称结构体,未设置字段使用默认值