c - atomicAdd 导致错误无法启动/执行内核

Question

我有以下 CUDA C 代码:

 int i = threadIdx.x + blockIdx.x*blockDim.x;
 int stride = blockDim.x*gridDim.x;
 while(i < size)
 {
     atomicAdd(&(histo_private[buffer[i]]),1);
     i+=stride;
 }

这导致我的程序崩溃并出现错误:“无法启动/执行内核”

这里 buffer 是这个 size 元素函数的整数输入数组，histo_private 是 histo_size 元素共享内存中的整数数组。我知道这不是索引越界错误，因为当我使用代码时:

int i = threadIdx.x + blockIdx.x*blockDim.x;
int stride = blockDim.x*gridDim.x;
while(i < size)
{
     int a = histo_private[buffer[i]];
     i+=stride;
}

所以我推测atomicAdd函数和/或这个32位int数组的内存地址有问题。

kernel.cu 文件包含以下代码:

// Define your kernels in this file you may use more than one kernel if you
// need to

// INSERT KERNEL(S) HERE

__global__ void histo_kernel(unsigned int* buffer, unsigned int size, int* histo, unsigned int histo_size)
{
    extern __shared__ int histo_private[];
    if(threadIdx.x < histo_size)
        histo_private[threadIdx.x] = 0;
    __syncthreads();

    // compute block's histogram
    int i = threadIdx.x + blockIdx.x*blockDim.x;
    int stride = blockDim.x*gridDim.x;
    while(i < size)
    {
        //int a = histo_private[buffer[i]];
        atomicAdd(&(histo_private[buffer[i]]),1);
        i+=stride;
    }

    // store to global histogram
    __syncthreads();
    //if(threadIdx.x < histo_size)
    //  atomicAdd(&(histo[threadIdx.x]),histo_private[threadIdx.x]);
}   

// ensures that no bins contains more than 255 elements
__global__ void enforce_saturation(int* histo, unsigned int histo_size)
{
    int i = threadIdx.x + blockIdx.x*blockDim.x;
    if(i < histo_size)
    {
        if(histo[i] > 255)  // this will be necessary to prevent data loss
            histo[i] = 255; // when converting from int to uint8_t                          
    }
}

__global__ void construct_histo(uint8_t* histo_unpacked, int* histo, unsigned int histo_size)
{
    int i = threadIdx.x + blockIdx.x*blockDim.x;
    if(i < histo_size)
        histo_unpacked[i] = histo[i];
}

// unpacks the input array into an output array with 'spaces'
__global__ void unpack(uint8_t* in, uint8_t* out, unsigned int size)
{
    int i = threadIdx.x + blockIdx.x*blockDim.x;
    if(i < size)
    {
        out[4*i] = in[i];
        out[4*i+1] = 0;
        out[4*i+2] = 0;
        out[4*i+3] = 0;
    }
}

// converts the input uint8_t array to an int array
__global__ void convert(uint8_t* in, int* out, unsigned int size)
{
    int i = threadIdx.x + blockIdx.x*blockDim.x;
    if(i < size)
    {
        out[i] = (int) in[4*i];
    }
}

// converts the input int array to a uint8_t array 
__global__ void convert_back(int* in, uint8_t* out, unsigned int size)
{
    int i = threadIdx.x + blockIdx.x*blockDim.x;
    if(i < size)
    {
        out[i] = (uint8_t) in[i];
    }
}



void histogram(unsigned int* input, uint8_t* bins, unsigned int num_elements, unsigned int num_bins) 
{

    int BLOCK_SIZE = (int) num_bins;
    BLOCK_SIZE = 512;
    dim3 dim_grid, dim_block;
    dim_block.x = BLOCK_SIZE; dim_block.y = dim_block.z = 1;
        dim_grid.x = 1+(num_elements-1)/BLOCK_SIZE; dim_grid.y = dim_grid.z = 1;

    // create an array of uint8_t to be converted into an array of int
    uint8_t* bins_unpacked;
    cudaMalloc((void**)&bins_unpacked, 4 * num_bins * sizeof(uint8_t));

    // unpack the input uint8_t array
    unpack<<<dim_grid,dim_block>>>(bins, bins_unpacked, num_bins);

    // need an int version of bins_d
    int* bins_int_d;
    cudaMalloc((void**)&bins_int_d, num_bins * sizeof(int));

    // convert the uint8_t array to an int array
    convert<<<dim_grid,dim_block>>>(bins_unpacked, bins_int_d, num_bins);   

    // run kernel and enforce saturation requirements
    int histo_private_size = num_bins;
    histo_kernel<<<dim_grid,dim_block,histo_private_size>>>(input, num_elements, bins_int_d, num_bins);
    enforce_saturation<<<dim_grid,dim_block>>>(bins_int_d,num_bins);

    // convert the int array back to uint8_t
    convert_back<<<dim_grid,dim_block>>>(bins_int_d, bins, num_bins);
}       

虽然调用最后一个直方图函数的函数位于 main.cu 中(我没有制作第二个文件——它是提供给我的——而且，我一直在通过 make test 编译来测试一致的数据)模式):

#include <stdio.h>
#include <stdint.h>

#include "support.h"
#include "kernel.cu"

int main(int argc, char* argv[])
{
    Timer timer;

    // Initialize host variables ----------------------------------------------

    #if TEST_MODE
    printf("\n***Running in test mode***\n"); fflush(stdout);
    #endif

    printf("\nSetting up the problem..."); fflush(stdout);
    startTime(&timer);

    unsigned int *in_h;
    uint8_t* bins_h;
    unsigned int *in_d;
    uint8_t* bins_d;
    unsigned int num_elements, num_bins;
    cudaError_t cuda_ret;

    if(argc == 1) {
        num_elements = 1000000;
        num_bins = 4096;
    } else if(argc == 2) {
        num_elements = atoi(argv[1]);
        num_bins = 4096;
    } else if(argc == 3) {
        num_elements = atoi(argv[1]);
        num_bins = atoi(argv[2]);
    } else {
        printf("\n    Invalid input parameters!"
           "\n    Usage: ./histogram            # Input: 1,000,000, Bins: 4,096"
           "\n    Usage: ./histogram <m>        # Input: m, Bins: 4,096"
           "\n    Usage: ./histogram <m> <n>    # Input: m, Bins: n"
           "\n");
        exit(0);
    }
    initVector(&in_h, num_elements, num_bins);
    bins_h = (uint8_t*) malloc(num_bins*sizeof(uint8_t));

    // TESTING
    for(unsigned int i = 0; i < num_bins; ++i) 
    {
        bins_h[i] = i;
        //printf("uint8_t Element %u: is %u \n", i, bins_h[i]);
    }



    stopTime(&timer); printf("%f s\n", elapsedTime(timer));
    printf("    Input size = %u\n    Number of bins = %u\n", num_elements,
        num_bins);

    // Allocate device variables ----------------------------------------------

    printf("Allocating device variables..."); fflush(stdout);
    startTime(&timer);

    cuda_ret = cudaMalloc((void**)&in_d, num_elements * sizeof(unsigned int));
    if(cuda_ret != cudaSuccess) FATAL("Unable to allocate device memory");
    cuda_ret = cudaMalloc((void**)&bins_d, num_bins * sizeof(uint8_t));
    if(cuda_ret != cudaSuccess) FATAL("Unable to allocate device memory");

    cudaDeviceSynchronize();
    stopTime(&timer); printf("%f s\n", elapsedTime(timer));

    // Copy host variables to device ------------------------------------------

    printf("Copying data from host to device..."); fflush(stdout);
    startTime(&timer);

    cuda_ret = cudaMemcpy(in_d, in_h, num_elements * sizeof(unsigned int),
        cudaMemcpyHostToDevice);
    if(cuda_ret != cudaSuccess) FATAL("Unable to copy memory to the device");

    cuda_ret = cudaMemset(bins_d, 0, num_bins * sizeof(uint8_t));
    if(cuda_ret != cudaSuccess) FATAL("Unable to set device memory");

    // TESTING
    //cuda_ret = cudaMemcpy(bins_d, bins_h, num_bins * sizeof(uint8_t),
    //    cudaMemcpyHostToDevice);
    //if(cuda_ret != cudaSuccess) FATAL("Unable to copy memory to the device");



    cudaDeviceSynchronize();
    stopTime(&timer); printf("%f s\n", elapsedTime(timer));

    // Launch kernel ----------------------------------------------------------
    printf("Launching kernel..."); fflush(stdout);
    startTime(&timer);

    histogram(in_d, bins_d, num_elements, num_bins);
    cuda_ret = cudaDeviceSynchronize();
    if(cuda_ret != cudaSuccess) FATAL("Unable to launch/execute kernel");

    stopTime(&timer); printf("%f s\n", elapsedTime(timer));

    // Copy device variables from host ----------------------------------------

    printf("Copying data from device to host..."); fflush(stdout);
    startTime(&timer);

    cuda_ret = cudaMemcpy(bins_h, bins_d, num_bins * sizeof(uint8_t),
        cudaMemcpyDeviceToHost);
    if(cuda_ret != cudaSuccess) FATAL("Unable to copy memory to host");

    cudaDeviceSynchronize();
    stopTime(&timer); printf("%f s\n", elapsedTime(timer));

    #if TEST_MODE
    printf("\nResult:\n");
    for(unsigned int binIdx = 0; binIdx < num_bins; ++binIdx) {
       printf("Bin %u: %u elements\n", binIdx, bins_h[binIdx]);
    }

    printf("\nElements Vec:\n");
    for(unsigned int i = 0; i < num_elements; ++i) {
        printf("Element %u: %u  is \n", i, in_h[i]);
    }



    #endif

    // Verify correctness -----------------------------------------------------

    printf("Verifying results..."); fflush(stdout);

    verify(in_h, bins_h, num_elements, num_bins);

    // Free memory ------------------------------------------------------------

    cudaFree(in_d); cudaFree(bins_d);
    free(in_h); free(bins_h);

    return 0;
}

Answer 1

事实证明这只是一个索引越界错误。元素 buffer[i] 大于 histo_private 的长度。正如另一位发帖者提到的，由于 C 编译器的以下工件，这一点并不明显:

编译器可以假设每次访问都在范围内。如果访问在范围内，我的测试代码行将不执行任何操作，因此允许编译器假设该行代码不执行任何操作。因此它不需要访问，因此测试代码的成功运行会产生误导。一旦该行更改为在 buffer[i] 处修改变量 hist_private 的位置，就会出现运行时错误。