4

我有一个函数，它可以拍摄彩色照片并返回它的灰色版本。如果我在主机上运行顺序代码，一切正常。如果我在设备上运行它，结果会略有不同（1000 中的一个像素与正确值相比是 +1 或 -1）。

我认为这与转换有关，但我不确定。这是我使用的代码：

    __global__ void rgb2gray_d (unsigned char *deviceImage, unsigned char *deviceResult, const int height, const int width){
    /* calculate the global thread id*/
    int threadsPerBlock  = blockDim.x * blockDim.y;
    int threadNumInBlock = threadIdx.x + blockDim.x * threadIdx.y;
    int blockNumInGrid   = blockIdx.x  + gridDim.x  * blockIdx.y;

    int globalThreadNum = blockNumInGrid * threadsPerBlock + threadNumInBlock;
    int i = globalThreadNum;

    float grayPix = 0.0f;
    float r = static_cast< float >(deviceImage[i]);
    float g = static_cast< float >(deviceImage[(width * height) + i]);
    float b = static_cast< float >(deviceImage[(2 * width * height) + i]);
    grayPix = (0.3f * r) + (0.59f * g) + (0.11f * b);

    deviceResult[i] = static_cast< unsigned char > (grayPix);
}

void rgb2gray(unsigned char *inputImage, unsigned char *grayImage, const int width, const int height, NSTimer &timer) {

    unsigned char *deviceImage;
    unsigned char *deviceResult;

    int initialBytes = width * height * 3;  
    int endBytes =  width * height * sizeof(unsigned char);

    unsigned char grayImageSeq[endBytes];

    cudaMalloc((void**) &deviceImage, initialBytes);
    cudaMalloc((void**) &deviceResult, endBytes);
    cudaMemset(deviceResult, 0, endBytes);
    cudaMemset(deviceImage, 0, initialBytes);

    cudaError_t err = cudaMemcpy(deviceImage, inputImage, initialBytes, cudaMemcpyHostToDevice);    

    // Convert the input image to grayscale 
    rgb2gray_d<<<width * height / 256, 256>>>(deviceImage, deviceResult, height, width);
    cudaDeviceSynchronize();

    cudaMemcpy(grayImage, deviceResult, endBytes, cudaMemcpyDeviceToHost);

    ////// Sequential
    for ( int y = 0; y < height; y++ ) {
             for ( int x = 0; x < width; x++ ) {
                   float grayPix = 0.0f;
                   float r = static_cast< float >(inputImage[(y * width) + x]);
                   float g = static_cast< float >(inputImage[(width * height) + (y * width) + x]);
                   float b = static_cast< float >(inputImage[(2 * width * height) + (y * width) + x]);

                   grayPix = (0.3f * r) + (0.59f * g) + (0.11f * b);
                   grayImageSeq[(y * width) + x] = static_cast< unsigned char > (grayPix);
              }
        }

    //compare sequential and cuda and print pixels that are wrong
    for (int i = 0; i < endBytes; i++)
    {
        if (grayImage[i] != grayImageSeq[i])
        cout << i << "-" << static_cast< unsigned int >(grayImage[i]) <<
                 " should be " << static_cast< unsigned int >(grayImageSeq[i]) << endl;
        }

    cudaFree(deviceImage);
    cudaFree(deviceResult);
}

我提到我为初始图像分配宽度 * 高度 * 3，因为初始图像是 CImg。

我在 GeForce GTX 480 上工作。

4

2 回答 2

5

最后我找到了答案。CUDA 会自动融合单精度和双精度乘加。使用下面的文档1，第 4.4 节，我设法修复它。而不是做

grayPix = (0.3f * r) + (0.59f * g) + (0.11f * b);

我现在正在做

grayPix = __fadd_rn(__fadd_rn(__fmul_rn(0.3f, r),__fmul_rn(0.59f, g)), __fmul_rn(0.11f, b));

这将禁用乘法的合并并添加到融合的乘法加法指令中。

NVIDIA GPU 的浮点和 IEEE 754 合规性

于 2013-01-19T10:58:26.280 回答
1

浮点数学可以在设备代码和主机代码中产生略微不同的结果。

为什么会出现这种情况有多种可能性。您必须考虑这两个函数是由两个不同的编译器编译成两个不同的二进制程序，在两个不同的浮点硬件实现上运行。

例如，如果浮点计算以不同的顺序执行，舍入误差可能会导致不同的结果。

此外，在 x86 架构 CPU 上使用 32 位（浮点）或 64 位（双精度）浮点表示运行浮点计算时，浮点数学由内部使用 80 位精度的 FPU 单元完成，结果为然后截断回 32 位浮点数据类型或 64 位双精度数据类型。

GPU 的 ALU 使用 32 位精度进行浮点数学运算（假设您使用的是浮点数据类型）。

可以在这里找到一篇讨论浮点表示和算术主题的优秀文章。

于 2013-01-18T20:53:26.497 回答