c++ - CUDA tex1Dfetch() 错误行为

Question

我对 CUDA 编程非常陌生，我面临着一个让我发疯的问题。这是怎么回事: 我有一个非常简单的程序(仅用于学习目的)，其中创建了一张输入图像和一张输出图像 16x16。输入图像被初始化为 0..255 的值，然后绑定(bind)到纹理。 CUDA 内核只是将输入图像复制到输出图像。输入图像值是通过调用 tex1Dfetch() 获得的，在某些情况下它会返回非常奇怪的值。请参阅下面的代码、内核中的注释以及程序的输出。代码完整且可编译，因此您可以在 VC 中创建 CUDA 项目并将代码粘贴到主“.cu”文件中。

请帮助我!我做错了什么？

我正在使用 VS 2013 社区和 VS 2013 的 CUDA SDK 6.5 + CUDA 集成。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>

texture<unsigned char> tex;

cudaError_t testMyKernel(unsigned char * inputImg, unsigned char * outputImg, int width, int height);

__global__ void myKernel(unsigned char *outImg, int width)
{
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int idx = row*width + col;
    __shared__ unsigned char input;
    __shared__ unsigned char input2;
    unsigned char *outPix = outImg + idx;

    //It fetches strange value, for example, when the idx==0 then the input is 51. 
    //But I expect that input==idx (according to the input image initialization).   
    input = tex1Dfetch(tex, idx);
    printf("Fetched for idx=%d: %d\n", idx, input);
    *outPix = input;

    //Very strange is that when I test the following code then the tex1Dfetch() returns correct values.
    if (idx == 0)
    {   
        printf("\nKernel test print:\n");
        for (int i = 0; i < 256; i++)
        {
            input2 = tex1Dfetch(tex, i);
            printf("%d,", input2);
        }
    }
}

int main()
{
    const int width = 16;
    const int height = 16;
    const int count = width * height;
    unsigned char imgIn[count];
    unsigned char imgOut[count];

    for (int i = 0; i < count; i++)
    {
        imgIn[i] = i;
    }

    cudaError_t cudaStatus = testMyKernel(imgIn, imgOut, width, height);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "testMyKernel failed!");
        return 1;
    }

    printf("\n\nOutput values:\n");
    for (int i = 0; i < height; i++)
    {
        for (int j = 0; j < width; j++)
        {
            printf("%d,", imgOut[i * width + j]);
        }
    }
    printf("\n");

    cudaStatus = cudaDeviceReset();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceReset failed!");
        return 1;
    }

    getchar();
    return 0;
}


cudaError_t testMyKernel(unsigned char * inputImg, unsigned char * outputImg, int width, int height)
{
    unsigned char * dev_in;
    unsigned char * dev_out;

    size_t size = width * height * sizeof(unsigned char);
    cudaError_t cudaStatus;

    cudaStatus = cudaSetDevice(0);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
        goto Error;
    }

    // input data
    cudaStatus = cudaMalloc((void**)&dev_in, size);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }
    cudaStatus = cudaMemcpy(dev_in, inputImg, size, cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }
    cudaStatus = cudaBindTexture(NULL, tex, dev_in, size);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaBindTexture failed!");
        goto Error;
    }

    // output data
    cudaStatus = cudaMalloc((void**)&dev_out, size);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    dim3 threadsPerBlock(4, 4);
    int blk_x = width / threadsPerBlock.x;  
    int blk_y = height / threadsPerBlock.y;
    dim3 numBlocks(blk_x, blk_y);

    // Launch a kernel on the GPU with one thread for each element.
    myKernel<<<numBlocks, threadsPerBlock>>>(dev_out, width);

    cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "myKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
        goto Error;
    }

    cudaStatus = cudaDeviceSynchronize();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching myKernel!\n", cudaStatus);
        goto Error;
    }

    //copy output image to host
    cudaStatus = cudaMemcpy(outputImg, dev_out, size, cudaMemcpyDeviceToHost);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

Error:
    cudaUnbindTexture(tex);
    cudaFree(dev_in);
    cudaFree(dev_out);

    return cudaStatus;
}

这是程序的输出(被截断一点):

Fetched for idx=0: 51
Fetched for idx=1: 51
Fetched for idx=2: 51
Fetched for idx=3: 51
Fetched for idx=16: 51
Fetched for idx=17: 51
Fetched for idx=18: 51
Fetched for idx=19: 51
Fetched for idx=32: 51
Fetched for idx=33: 51
Fetched for idx=34: 51
Fetched for idx=35: 51
Fetched for idx=48: 51
Fetched for idx=49: 51
Fetched for idx=50: 51
Fetched for idx=51: 51
Fetched for idx=192: 243
Fetched for idx=193: 243
Fetched for idx=194: 243
Fetched for idx=195: 243
Fetched for idx=208: 243
Fetched for idx=209: 243
Fetched for idx=210: 243
Fetched for idx=211: 243
Fetched for idx=224: 243
etc... (output truncated.. see the Output values)

Kernel test print:
0,1,2,3,4,5,6,7,8,9,10,11,12,13,14,15,16,17,18,19,20,21,22,23,24,25,26,27,28,29,
30,31,32,33,34,35,36,37,38,39,40,41,42,43,44,45,46,47,48,49,50,51,52,53,54,55,56
etc...(correct values)

Output values:
51,51,51,51,55,55,55,55,59,59,59,59,63,63,63,63,51,51,51,51,55,55,55,55,59,59,59
,59,63,63,63,63,51,51,51,51,55,55,55,55,59,59,59,59,63,63,63,63,51,51,51,51,55,55,
etc.. (wrong values)

Answer 1

内核的这一行

input = tex1Dfetch(tex, idx);

导致 block 的线程之间出现竞争状况。 block 中的所有线程都试图同时从纹理获取值到 __shared__ 变量 input 中，从而导致未定义的行为。您应该以 __shared__ 数组的形式为 block 中的每个线程分配单独的共享内存空间。

对于您当前的情况，可能类似于

__shared__ unsigned char input[16]; //4 x 4 block size

内核的其余部分应该类似于:

int idx_local = threadIdx.y * blockDim.x + threadIdx.x; //local id of thread in a block
input[idx_local] = tex1Dfetch(tex, idx);
printf("Fetched for idx=%d: %d\n", idx, input[idx_local]);
*outPix = input[idx_local];

内核末尾条件内的代码工作正常，因为由于指定的条件 if (idx == 0)，只有第一个 block 的第一个线程会执行所有操作串行处理，而所有其他线程将保持空闲，因此由于不存在竞争条件，问题将消失。