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

Posted 2023-02-16

【中文标题】atomicAdd 导致错误无法启动/执行内核

【英文标题】：atomicAdd causing error Unable to Launch/Execute Kernel

【发布时间】：2013-11-07 04:16:54

【问题描述】：

我有以下 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;

【问题讨论】：

你不知道这不是索引越界错误。您的测试代码允许编译器假设索引没有超出范围，如果不是，则不会产生任何后果，而如果索引超出范围，您的实际代码将崩溃。

你在哪个 GPU 上运行它？ “无法启动/执行内核”必须是您在程序中创建的消息。我们能看到生成该消息的代码和相关的错误检查吗？您能否提供用于编译程序的nvcc compile 命令行？您提供的代码 sn-p 没有任何明显错误，这让我回到了有关 SSCCE.org 代码的讨论中。是否有某些原因您不能围绕此处显示的代码编写一个简单的程序并发布整个程序？

@Erroldactyl 多读几遍我的解释，直到你理解为止。了解您是否想成为一名称职的 C 程序员至关重要。这是关键部分：“允许编译器假设每次访问都在界限内。如果访问在界限内，那行代码什么都不做。因此，编译器可以假设这行代码什么都不做。因此它没有”不需要访问权限。”

David 是说，由于您发布的代码与a 无关，因此实际上允许编译器通过消除它来优化该代码行 - 并消除访问权限。因此，您的代码基本上什么都不做，也不能证明您的访问模式。当您使用 cuda-memcheck 运行代码时会发生什么？

这个int histo_private_size = num_bins;不应该是这个：int histo_private_size = num_bins * sizeof(int);吗？

【参考方案1】：

原来这只是一个索引越界错误。元素缓冲区[i] 大于 histo_private 的长度。正如另一张海报所提到的，由于 c 编译器的以下工件，这并不明显：

允许编译器假定每次访问都在界限内。如果访问在范围内，我的测试代码行什么也不做，因此允许编译器假设该代码行什么都不做。因此它不需要访问，因此测试代码的成功运行具有误导性。一旦该行更改为在 buffer[i] 处修改变量 hist_private 的位置，就会出现运行时错误。

【讨论】：

编译器怎么能做其他事情？所有有问题的内存都是动态分配的，因此编译器无法知道大小和强制边界，即使存在这样的代码检查工具。

我要打败一匹死马，因为这对我来说是一个热键，以防你没有注意到。根据您最初的发布，任何人都无法发现这个越界问题。这是一个强烈的动机，建议提出此类问题的人提供一个完整的复制者，即 SSCCE.org，这正是 SO 的既定期望。

@talonmies 它可以在运行时轻松完成，这大概是 OP 所期望的。 （并且有一些方法可以让 C 和 C++ 代码这样做。）