0

我尝试自己使用 CUDA 实现向量和缩减，但遇到了一个我可以修复但不了解实际问题的错误。

我在下面实现了内核，它与 NVIDIA 的示例中使用的内核几乎相同。

__global__ 
void reduce0(int *input, int *output)
{
    extern __shared__ int s_data[];

    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    s_data[tid] = input[i];
    __syncthreads();

    for( int s=1; s < blockDim.x; s *= 2) {
        if((tid % 2*s) == 0) {
            s_data[tid] += s_data[tid + s];
        }

        __syncthreads();
    }

    if(tid == 0) {
        output[blockIdx.x] = s_data[0];
    }
}

此外，我在主机端计算共享内存空间如下

int sharedMemSize = numberOfValues * sizeof(int);

如果使用了超过 1 个线程块，则代码运行良好。仅使用 1 个块会导致上述索引超出范围错误。通过将我的主机代码与我发现以下行的示例之一进行比较来查找我的错误：

int smemSize = (threads <= 32) ? 2 * threads * sizeof(T) : threads * sizeof(T);

稍微玩一下我的块/网格设置，我得到了以下结果：

  1. 块，任意数量的线程 => 代码崩溃
  2. >2 块，任意数量的线程 => 代码运行良好
  3. 1 块，任意数量的线程，共享内存大小 2*#threads => 代码运行良好

尽管考虑了几个小时，但我不明白为什么在使用太少数量的线程或块时会出现越界错误。

更新：主机代码按要求调用内核

int numberOfValues = 1024 ;
int numberOfThreadsPerBlock = 32;
int numberOfBlocks = numberOfValues / numberOfThreadsPerBlock;

int memSize = sizeof(int) * numberOfValues;

int *values = (int *) malloc(memSize);
int *result = (int *) malloc(memSize);

int *values_device, *result_device;
cudaMalloc((void **) &values_device, memSize);
cudaMalloc((void **) &result_device, memSize);

for(int i=0; i < numberOfValues ; i++) {
    values[i] = i+1;
}

cudaMemcpy(values_device, values, memSize, cudaMemcpyHostToDevice);

dim3 dimGrid(numberOfBlocks,1);
dim3 dimBlock(numberOfThreadsPerBlock,1);
int sharedMemSize = numberOfThreadsPerBlock * sizeof(int);

reduce0 <<< dimGrid, dimBlock, sharedMemSize >>>(values_device, result_device);

if (cudaSuccess != cudaGetLastError())
        printf( "Error!\n" );

cudaMemcpy(result, result_device, memSize, cudaMemcpyDeviceToHost);
4

1 回答 1

1

你的问题可能是模数和乘法的优先顺序吗？ tid % 2*s等于(tid % s)*2但你想要tid % (s*2)

您需要int smemSize = (threads <= 32) ? 2 * threads * sizeof(T) : threads * sizeof(T)用于少量线程的原因是超出范围的索引。发生这种情况的一个示例是当您启动 29 个线程时。何时tid=28和s=2分支将被采用，28 % (2*2) == 0并且您将索引到s_data[28+2]但您只为 29 个线程分配了共享内存。

于 2013-01-04T10:35:05.057 回答