CUDA实现1到100整数、浮点数求和及代码报错排查

原代码问题点 #

• 核函数中所有线程都会执行sum = 0操作，会反复重置累加结果，导致原子加的结果被覆盖
• threadIdx.x取值范围是0_{99，要计算1}100的和需要加threadIdx.x + 1
• 设备端的__device__变量不能直接在主机端用printf读取，需要通过cudaMemcpyFromSymbol接口拷贝到主机侧内存后再输出
• 核函数启动是异步操作，主机端需要调用cudaDeviceSynchronize()等待核函数执行完成，再读取结果

修复后代码 #

#include <stdio.h>
#include <cuda.h>

__device__ int d_sum;

__global__ void kernel(){
    // 只让第一个线程做初始化，避免反复重置
    if (threadIdx.x == 0) {
        d_sum = 0;
    }
    // 所有线程同步，确保初始化完成再累加
    __syncthreads();
    atomicAdd(&d_sum, threadIdx.x + 1);
}

int main(){
    int h_sum = 0;
    kernel<<<1,100>>>();
    // 等待核函数执行完成
    cudaDeviceSynchronize();
    // 将设备端变量拷贝到主机端
    cudaMemcpyFromSymbol(&h_sum, d_sum, sizeof(int));
    printf("1~100整数求和结果：%d\n", h_sum);
    return 0;
}

首先明确：CUDA从计算能力2.0开始已经原生支持float类型的atomicAdd，如果你的GPU是2010年之后发布的，直接替换int为float即可使用。如果需要兼容旧设备，或者需要实现double类型的原子加，可以通过atomicCAS接口实现自旋锁版本的原子加，示例代码如下：

#include <stdio.h>
#include <cuda.h>

__device__ float d_sum_float;

// 兼容旧设备的float原子加实现
__device__ float atomicAddFloat(float* addr, float value) {
    unsigned int* addr_as_uint = (unsigned int*)addr;
    unsigned int old = *addr_as_uint, assumed;
    do {
        assumed = old;
        old = atomicCAS(addr_as_uint, assumed, __float_as_uint(value + __uint_as_float(assumed)));
    } while (assumed != old);
    return __uint_as_float(old);
}

__global__ void float_sum_kernel(){
    if (threadIdx.x == 0) {
        d_sum_float = 0.0f;
    }
    __syncthreads();
    atomicAddFloat(&d_sum_float, (float)(threadIdx.x + 1));
}

int main(){
    float h_sum_float = 0.0f;
    float_sum_kernel<<<1,100>>>();
    cudaDeviceSynchronize();
    cudaMemcpyFromSymbol(&h_sum_float, d_sum_float, sizeof(float));
    printf("1~100浮点数求和结果：%.1f\n", h_sum_float);
    return 0;
}

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