长沙诚拙微电子面经

面试官人很和蔼,其中一个还是校友,但我还是太菜了,感觉实习这么久后期有点赛博打螺丝了,稍微吐槽一下目前实习的公司(合肥半导体公司),我都实习差不多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 的同一列,地址不连续。

这会导致全局内存写入性能极差。

解决方案就是使用共享内存分块转置

核心思想:

  1. 分块(Tile):将大矩阵分成小块(如 32x32)。
  2. 协作加载:一个线程块负责一个 tile。
  3. 读取时合并:从全局内存合并地读取一个 tile 到共享内存。
  4. 在共享内存中转置:线程将数据“转置”地写入共享内存。
  5. 写回时合并:从共享内存合并地写回全局内存

关键:共享内存没有合并访存的概念!它是一个低延迟的片上内存,我们可以在其中任意顺序读写。真正的合并访存优化发生在全局内存的读和写

#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),把所有参数打包称结构体,未设置字段使用默认值

全部评论

相关推荐

评论
点赞
收藏
分享

创作者周榜

更多
牛客网
牛客网在线编程
牛客网题解
牛客企业服务