未修改 CUDA 核函数输出变量

Question

我正在尝试将对象传递给内核。这个对象基本上有两个变量，一个作为输入，另一个作为内核的输出。但是当我启动内核时，输出变量不会改变。但是当我向内核添加另一个变量并将输出值也分配给这个变量时，它突然对它们都有效。

我在另一个线程 (While loop fails in CUDA kernel) 中读到，如果内核不产生任何输出，编译器可以将内核评估为空以进行优化。

所以我作为唯一内核参数传递的这个输入/输出对象有可能没有被编译器以某种方式识别为输出？如果这是真的。有没有一种优雅的方法（我想避免添加另一个内核参数），例如可以防止这种情况的编译选项？

这是这个对象的类。

class Replica
{
    public :
        signed char gA[1024];
        int MA;
    __device__ __host__ Replica(){
    }   
};

这是一个基本上是求和的内核。

__global__ void sumKerA(Replica* Rd) 
{
    int t = threadIdx.x;
    int b = blockIdx.x;

    __shared__ signed short gAs[1024];
    gAs[t] = Rd[b].gA[t];

    for (unsigned int stride = 1024 >> 1; stride > 0; stride >>= 1){
        __syncthreads();
        if (t < stride){
            gAs[t] += gAs[t + stride];
        }
    }
    __syncthreads();

    if (t == 0){
        Rd[b].MA = gAs[0];
    }
}

最后是我的主机代码。

int main ()
{
    // replicas - array of objects
    Replica R[128];
    for (int i = 0; i < 128; ++i){
        for (int j = 0; j < 1024; ++j){
            R[i].gA[j] = 2*(rand() % 2) - 1;
        }
        R[i].MA = 0;
    }

    Replica* Rd;

    cudaSetDevice(0);

    cudaMalloc((void **)&Rd,128*sizeof(Replica));
    cudaMemcpy(Rd,R,128*sizeof(Replica),cudaMemcpyHostToDevice);

    dim3 DimBlock(1024,1,1);
    dim3 DimGridA(128,1,1);

    sumKerA <<< DimBlock, DimGridA >>> (Rd);
    cudaThreadSynchronize();

    cudaMemcpy(&R,Rd,128*sizeof(Replica),cudaMemcpyDeviceToHost);
    // cudaMemcpy(&M,Md,128*sizeof(int),cudaMemcpyDeviceToHost);
    for (int i = 0; i < 128; ++i){
        cout << R[i].MA << " ";
    }

    cudaFree(Rd);

    return 0;
}

Answer 1

根据您的缩减代码，您似乎打算在每个块中启动 1024 个线程。

在这种情况下，这是不正确的：

dim3 DimBlock(1024,1,1);
dim3 DimGridA(128,1,1);

sumKerA <<< DimBlock, DimGridA >>> (Rd);

第一个内核配置参数是网格的尺寸。第二个参数是线程块的维度。如果你希望每个块有 1024 个线程，同时启动 128 个块，你的内核启动应该是这样的：

sumKerA <<< DimGridA, DimBlock >>> (Rd);

如果您将proper cuda error checking 添加到您的代码中，我希望您会看到内核启动失败，因为使用块变量 (blockIdx.x) 来索引包含 128 个元素的 Rd 数组的索引将超出数组，在你原来的情况下。

如果您在内核中修改 Rd 指向的 Replica 对象，这是外部可见状态，因此任何修改这些对象的代码都不能被编译器“优化掉”。

另请注意，cudaThreadSynchronize() 已被弃用，取而代之的是 cudaDeviceSynchronize()（它们具有相同的行为。）