c++ - CUDA 线程在全局内存中的私有(private)位置写入错误结果

Question

编辑 3:我需要每个线程在全局内存中写入和读取一个私有(private)位置。下面我发布了一个显示我的问题的工作代码。下面，我将列出涉及的主要变量和结构。

变量:

• srcArr_h (host) --> srcArr_d (device) : [0, COLORLEVELS] 范围内的随机 float 组，维度由 ARRDIM
给出
• auxD (device) : 维度数组 ARRDIM * ARRDIM 将最终结果保存在设备中
• auxH (host) : 维度数组 ARRDIM * ARRDIM 将最终结果保存在 host
• c_glob_d(设备):为每个线程保留一个私有(private)位置 COLORLEVELS 的数组，大小由 num_threads * 色阶
• idx (device) : 当前线程的标识号

我的问题:在内核中，我为每个值 ic 更新了 c_glob[idx](ic ∈ [0, COLORLEVELS]), 即 c_glob[idx][ic]。我使用 c_glob[idx][COLORLEVELS] 来计算存储在 auxD 中的最终结果 g0。我的问题是我的最终结果是错误的。复制到 auxH 的结果显示我得到的数字至少比预期大一个数量级，甚至是奇怪的数字，表明我的操作可能会溢出。
帮助:我做错了什么？如何让每个线程写入和读取全局内存中的每个私有(private)位置？现在我正在使用 ARRDIM = 512 进行调试，但我的目标是使其适用于 ARRDIM~ 10^4，从而创建一个 c_glob 10^4*10^4 个线程的数组)。我想每次运行允许的线程总数会有问题。所以我想知道您是否可以针对我的问题提出任何其他解决方案。
谢谢。

#include <string>
#include <stdint.h>
#include <iostream>
#include <stdio.h>
#include "cuPrintf.cu"
using namespace std;

#define ARRDIM 512
#define COLORLEVELS 4

__global__ void gpuKernel
(
    float *sa, float *aux,
    size_t memPitchAux, int w,
    float *c_glob
)
{
    float sc_loc[COLORLEVELS];

    float g0=0.0f;

    int tidx = blockIdx.x * blockDim.x + threadIdx.x; 
    int tidy = blockIdx.y * blockDim.y + threadIdx.y; 

    int idx  = tidy * memPitchAux/4 + tidx;

    for(int ic=0; ic<COLORLEVELS; ic++)
    {
        sc_loc[ic] = ((float)(ic*ic));
    }

    for(int is=0; is<COLORLEVELS; is++)
    {
        int ic = fabs(sa[tidy*w +tidx]);
        c_glob[tidy * COLORLEVELS + tidx + ic] += 1.0f;
    }

    for(int ic=0; ic<COLORLEVELS; ic++)
    {
        g0 += c_glob[tidy * COLORLEVELS + tidx + ic]*sc_loc[ic];
    }

    aux[idx] = g0;
}

int main(int argc, char* argv[])
{
    /*
     * array src host and device
     */
    int heightSrc = ARRDIM;
    int widthSrc = ARRDIM;
    cudaSetDevice(0);

    float *srcArr_h, *srcArr_d;
    size_t nBytesSrcArr = sizeof(float)*heightSrc * widthSrc;

    srcArr_h = (float *)malloc(nBytesSrcArr); // Allocate array on host
    cudaMalloc((void **) &srcArr_d, nBytesSrcArr); // Allocate array on device
    cudaMemset((void*)srcArr_d,0,nBytesSrcArr); // set to zero

    int totArrElm = heightSrc*widthSrc;

    for(int ic=0; ic<totArrElm; ic++)
    {
        srcArr_h[ic] = (float)(rand() % COLORLEVELS);
    }

    cudaMemcpy( srcArr_d, srcArr_h,nBytesSrcArr,cudaMemcpyHostToDevice);

    /*
     * auxiliary buffer auxD to save final results
     */
    float *auxD;
    size_t auxDPitch;
    cudaMallocPitch((void**)&auxD,&auxDPitch,widthSrc*sizeof(float),heightSrc);
    cudaMemset2D(auxD, auxDPitch, 0, widthSrc*sizeof(float), heightSrc);

    /*
     * auxiliary buffer auxH allocation + initialization on host
     */
    size_t auxHPitch;
    auxHPitch = widthSrc*sizeof(float);
    float *auxH = (float *) malloc(heightSrc*auxHPitch);

    /*
     * kernel launch specs
     */
    int thpb_x = 16;
    int thpb_y = 16;

    int blpg_x = (int) widthSrc/thpb_x;
    int blpg_y = (int) heightSrc/thpb_y;
    int num_threads = blpg_x * thpb_x + blpg_y * thpb_y;

    /* 
     * c_glob: array that reserves a private location of COLORLEVELS floats for each thread
     */
    int cglob_w = COLORLEVELS;
    int cglob_h = num_threads;

    float *c_glob_d;
    size_t c_globDPitch;
    cudaMallocPitch((void**)&c_glob_d,&c_globDPitch,cglob_w*sizeof(float),cglob_h);
    cudaMemset2D(c_glob_d, c_globDPitch, 0, cglob_w*sizeof(float), cglob_h);

    /*
    * kernel launch
    */
    dim3 dimBlock(thpb_x,thpb_y, 1);
    dim3 dimGrid(blpg_x,blpg_y,1);

    gpuKernel<<<dimGrid,dimBlock>>>(srcArr_d,auxD, auxDPitch, widthSrc, c_glob_d);

    cudaThreadSynchronize();

    cudaMemcpy2D(auxH,auxHPitch, 
                 auxD,auxDPitch,  
                 auxHPitch, heightSrc,
                 cudaMemcpyDeviceToHost);
    cudaThreadSynchronize();

    float min = auxH[0];
    float max = auxH[0];
    float f;
    string str;

    for(int i=0; i<widthSrc*heightSrc; i++)
    {

        if(min > auxH[i])
            min = auxH[i];
        if(max < auxH[i])
            max = auxH[i];
    }
    cudaFree(srcArr_d);
    cudaFree(auxD);
    cudaFree(c_glob_d);

}

Answer 1

您决定既不显示完整代码，也不显示重现您问题的缩小尺寸。因此，无法对以下可能的解决方案进行测试和验证。

我认为您已经找到问题的根源:多个线程正试图并行写入相同的内存位置。这是导致竞争条件的情况。例如，请参阅演示文稿的第四张幻灯片 "CUDA C: race conditions, atomics, locks, mutex, and warps" .

竞争条件有一个蛮力解决方案:原子函数。它们在 CUDA C 编程指南的第 B.12 节中进行了描述。所以你可以尝试通过更改行来解决你的问题

c[ic] += 1.0f;

到

atomicAdd(&c[ic],1);

您将为此修复付出性能代价:原子操作序列化代码以避免竞争条件。

我已经提到原子函数是解决您问题的一种蛮力解决方案，因为通过适本地重新考虑实现，您可以找到避免它们的方法。但由于您提供的细节非常少，目前还无法确定。