面试官面试问：你能写出十种矩阵加法的cuda程序吗？(二)

前言

在CUDA编程中，线程、线程块和网格是三个基本的并行计算单位：

1. 线程 (Thread)：线程是执行基本计算任务的最小单元，每个线程执行相同的代码但处理不同的数据。

2. 线程块 (Thread Block)：线程块是一组线程的集合，这些线程可以共享数据并通过同步机制协调工作，线程块的大小在执行时固定。

3. 网格 (Grid)：网格是多个线程块的集合，所有线程块并行执行任务，网格的大小同样在执行时固定。

01 一般编写

__global__ void addKernel(float *pA, float *pB, float *pC, int size)
{
    int index = blockIdx.x * blockDim.x + threadIdx.x; // 计算当前数组中的索引
    if (index >= size)
        return;

    pC[index] = pA[index] + pB[index];
}

addKernel<<<arr_len + 512 - 1, 512>>>(dev_a, dev_b, dev_c, arr_len);

02 使用线程

__global__ void addKernel_thread(float *pA, float *pB, float *pC, int size)
{
    int index = threadIdx.x; // 计算当前数组中的索引
    if (index >= size)
        return;

    pC[index] = pA[index] + pB[index];
}

addKernel_thread<<<1, arr_len>>>(dev_a, dev_b, dev_c, arr_len);

03  使用block块进行编程

__global__ void addKernel_block(float *pA, float *pB, float *pC, int size)
{
    int index = blockIdx.x; // 计算当前数组中的索引
    if (index >= size)
        return;

    pC[index] = pA[index] + pB[index];
}

addKernel_block<<<arr_len, 1>>>(dev_a, dev_b, dev_c, arr_len);

04 使用循环

__global__ void readOffsetUnroll2(float *A, float *B, float *C, const int n,
                                  int offset)
{
    unsigned int i = blockIdx.x * blockDim.x * 2 + threadIdx.x;
    unsigned int k = i + offset;

    if (k < n) C[i] = A[k] + B[k];
    if (k + blockDim.x < n) {
        C[i + blockDim.x] = A[k + blockDim.x] + B[k + blockDim.x];
    }
}

readOffsetUnroll2<<<grid.x/2, block>>>(d_A, d_B, d_C, nElem, offset);

05 基于grid2d_bock1d

// grid 2D block 1D
__global__ void sumMatrixOnGPUMix(float *MatA, float *MatB, float *MatC, int nx,
                                  int ny)
{
    unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;
    unsigned int iy = blockIdx.y;
    unsigned int idx = iy * nx + ix;
    printf("ix iy idx %d %d %d %d %d \n", ix, iy, idx, blockDim.x, blockDim.y);
    if (ix < nx && iy < ny)
        MatC[idx] = MatA[idx] + MatB[idx];
}

    int dimx = 3;
    dim3 block(dimx, 2);
    dim3 grid((nx + block.x - 1) / block.x, ny);

    iStart = seconds();
    sumMatrixOnGPUMix<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);

06 基于grid2d_block2d

// grid 2D block 2D
__global__ void sumMatrixOnGPU2D(float *MatA, float *MatB, float *MatC, int nx,
                                 int ny)
{
    unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;
    unsigned int iy = threadIdx.y + blockIdx.y * blockDim.y;
    unsigned int idx = iy * nx + ix;
    if (ix < nx && iy < ny) {
        MatC[idx] = MatA[idx] + MatB[idx];
        printf("ix iy idx %d %d %d \n", ix, iy, idx);
    }
}
    // invoke kernel at host side
    int dimx = 2;
    int dimy = 2;
    dim3 block(dimx, dimy);
    dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y);

    sumMatrixOnGPU2D<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    

你知道另外四种是如何编写吗？在评论区留下你的想法。

最后别忘了，帮忙点“在看”。  

您的点赞，在看，是我创作的动力。

关注我的公众号auto_driver_ai(Ai fighting), 第一时间获取更新内容。

AiFighing是全网第一且唯一以代码、项目的形式讲解自动驾驶感知方向的关键技术，关注我，一起学习自动驾驶感知技术。