CUDA实现RGB转灰度图异常：从设备端取回的数组全为0

你的代码出现结果全为0的问题，核心是存在几个关键错误，以下是问题拆解和修复方案：

关键错误点 #

• 设备输出数组未初始化：cudaMalloc分配的设备内存是随机值，直接累加会导致结果不可控，且初始值不为0，最终结果必然错误。
• 变量名拼写错误：核函数中定义的变量是grayScaleAddition，但累加时误写为未定义的grayScale，编译报错会导致核函数无法正确执行。
• 回传数据字节数错误：h_out仅需存储9个float值（对应ARRAY_BYTES/3），但代码中用ARRAY_BYTES作为回传大小，会导致内存越界，无法正确获取结果。
• 缺少CUDA错误检查：所有CUDA API调用和核函数启动都未做错误校验，无法及时定位问题。
• 不必要的类型转换：将float类型的d_in元素赋值给int temp，虽当前场景数值正确，但易引发后续扩展问题。

修正后的完整代码 #

#include <stdio.h>

#define CHECK_CUDA_ERROR(err) \
    if (err != cudaSuccess) { \
        fprintf(stderr, "CUDA Error: %s at line %d\n", cudaGetErrorString(err), __LINE__); \
        exit(EXIT_FAILURE); \
    }

__global__ void RGBToGrayScale(float *d_out, float *d_in, int rowCount, int colCount) {
    int pixel_idx = threadIdx.x * colCount + threadIdx.y;
    float temp = d_in[blockIdx.x * rowCount * colCount + pixel_idx];
    float grayScaleAddition;

    switch(blockIdx.x) {
        case 0:
            grayScaleAddition = 0.299 * temp;
            break;
        case 1:
            grayScaleAddition = 0.587 * temp;
            break;
        case 2:
            grayScaleAddition = 0.114 * temp;
            break;
        default:
            grayScaleAddition = 0.0f;
            break;
    }
    atomicAdd(&d_out[pixel_idx], grayScaleAddition);
}

int main() {
    int image[3][3][3] = {
        {{1, 3, 2}, {4, 5, 6}, {7, 8, 9}},
        {{10, 11, 12}, {13, 14, 15}, {16, 17, 18}},
        {{19, 20, 21}, {22, 23, 24}, {25, 26, 27}}
    };

    const int IMAGE_ROW_COUNT = 3;
    const int IMAGE_COLUMN_COUNT = 3;
    const int TOTAL_PIXELS = IMAGE_ROW_COUNT * IMAGE_COLUMN_COUNT;
    const int ARRAY_BYTES = TOTAL_PIXELS * 3 * sizeof(float);
    const int OUTPUT_BYTES = TOTAL_PIXELS * sizeof(float);

    float* h_in = (float*)malloc(ARRAY_BYTES);
    float* h_out = (float*)malloc(OUTPUT_BYTES);

    if (!h_in || !h_out) {
        fprintf(stderr, "Host memory allocation failed\n");
        exit(EXIT_FAILURE);
    }

    for (int i = 0; i < IMAGE_ROW_COUNT; ++i) {
        for (int j = 0; j < IMAGE_COLUMN_COUNT; ++j) {
            for (int k = 0; k < 3; ++k) {
                h_in[k * TOTAL_PIXELS + i * IMAGE_COLUMN_COUNT + j] = (float)image[i][j][k];
            }
        }
    }

    float* d_in;
    float* d_out;
    cudaError_t err;

    err = cudaMalloc((void **) &d_in, ARRAY_BYTES);
    CHECK_CUDA_ERROR(err);

    err = cudaMalloc((void **) &d_out, OUTPUT_BYTES);
    CHECK_CUDA_ERROR(err);

    err = cudaMemset(d_out, 0, OUTPUT_BYTES);
    CHECK_CUDA_ERROR(err);

    err = cudaMemcpy(d_in, h_in, ARRAY_BYTES, cudaMemcpyHostToDevice);
    CHECK_CUDA_ERROR(err);

    RGBToGrayScale<<<3, dim3(IMAGE_ROW_COUNT, IMAGE_COLUMN_COUNT)>>>(d_out, d_in, IMAGE_ROW_COUNT, IMAGE_COLUMN_COUNT);
    err = cudaGetLastError();
    CHECK_CUDA_ERROR(err);
    err = cudaDeviceSynchronize();
    CHECK_CUDA_ERROR(err);

    err = cudaMemcpy(h_out, d_out, OUTPUT_BYTES, cudaMemcpyDeviceToHost);
    CHECK_CUDA_ERROR(err);

    for (int i = 0; i < TOTAL_PIXELS; ++i) {
        printf("%.4f", h_out[i]);
        printf(((i % 3) != 2) ? "\t" : "\n");
    }

    cudaFree(d_in);
    cudaFree(d_out);
    free(h_in);
    free(h_out);
    return 0;
}

额外优化说明 #

• 加入atomicAdd：三个block的线程会同时对d_out的同一位置累加，普通+=会引发竞争条件，atomicAdd保证操作原子性。
• 新增错误检查宏CHECK_CUDA_ERROR，快速定位CUDA操作中的问题。
• 简化索引计算，用pixel_idx统一管理像素位置，提升代码可读性。

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