CUDA遍历数组每间隔M个元素取值的核函数实现咨询

实现方案 #

你想到的加i % M == 0判断的方案可以跑通，但不是性能最优的选择，有两种可选的实现思路：

1. 高效无空转方案（推荐） #

直接把步长逻辑映射到线程索引上，让每个线程只处理有效位置的元素，完全没有冗余判断和空转，性能最好：

__global__ void dec(const complex * signal, int N, int M, complex * output)
{
    // 线程id对应第t个有效元素
    int t = blockIdx.x * blockDim.x + threadIdx.x;
    // 计算总共有多少个需要处理的有效元素
    int total_valid = (N / M) + 1;
    if (t < total_valid) {
        // 映射到原数组的步长位置
        int i = t * M;
        output[i] = signal[i];
        // 如果你用的是自定义complex结构体不能直接赋值，再拆成x和y赋值
        // output[i].x = signal[i].x;
        // output[i].y = signal[i].y;
    }
}

核函数启动的时候，按照有效元素数计算grid size即可，示例：

const int block_size = 256;
int total_valid = (N / M) + 1;
int grid_size = (total_valid + block_size - 1) / block_size;
dec<<<grid_size, block_size>>>(d_signal, N, M, d_output);

2. 模运算判断方案（不推荐） #

如果图省事不想改启动参数，也可以直接在原有核函数上加判断，只是当M比较大的时候，大部分线程会因为不满足条件空转，浪费GPU计算资源：

__global__ void dec(const complex * signal, int N, int M, complex * output)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i <= N && i % M == 0) 
    {
        output[i].x = signal[i].x;
        output[i].y = signal[i].y;
    }
}

线程同步问题说明 #

这个场景下不需要做任何线程同步操作：

• 每个线程只负责自己对应的i位置的读写，不同线程访问的内存地址没有重叠，不存在竞态条件
• 没有线程间共享数据的交互，不需要加__syncthreads()这类块内同步指令
• 核函数执行完成后，只需要正常调用cudaDeviceSynchronize()确认核函数跑完，再把结果从显存拷回内存即可

注意事项 #

• 注意你原始C++代码的循环条件是i <= N，所以计算有效元素数的时候不要漏了i=N的边界情况，(N / M) + 1的整数除法逻辑刚好匹配原始循环的步长规则
• 如果用CUDA官方头文件cuComplex.h定义的复数类型，支持直接赋值，不需要单独读写x、y分量，代码更简洁

内容的提问来源于stack exchange，提问作者Yihan Hu