c - 在不减少线程的情况下在 CUDA 中使用共享内存

Question

查看 Mark Harris 的缩减示例，我正在尝试查看是否可以让线程在不进行缩减操作的情况下存储中间值:

例如 CPU 代码:

for(int i = 0; i < ntr; i++)
{
    for(int j = 0; j < pos* posdir; j++)
    {
        val = x[i] * arr[j];
        if(val > 0.0)
        {
            out[xcount] = val*x[i];
            xcount += 1;
        }
    }
}

等效的 GPU 代码:

const int threads = 64; 
num_blocks = ntr/threads;

__global__ void test_g(float *in1, float *in2, float *out1, int *ct, int posdir, int pos)
{
    int tid = threadIdx.x + blockIdx.x*blockDim.x;
    __shared__ float t1[threads];
    __shared__ float t2[threads];

    int gcount  = 0;

    for(int i = 0; i < posdir*pos; i += 32) {
        if (threadIdx.x < 32) {
            t1[threadIdx.x] = in2[i%posdir];
        }
       __syncthreads();

        for(int i = 0; i < 32; i++)
        {
            t2[i] = t1[i] * in1[tid];
                if(t2[i] > 0){
                    out1[gcount] = t2[i] * in1[tid];
                    gcount = gcount + 1;
                }
        }
    }        
    ct[0] = gcount;
}

我在这里尝试做的是以下步骤:

(1)将in2的32个值存入共享内存变量t1，

(2)对于i和in1[tid]的每一个值，计算t2[i]，

(3)如果 t2[i] > 0 对于 i 的特定组合，将 t2[i]*in1[tid] 写入 out1[ gcount]

但是我的输出全错了。我什至无法统计 t2[i] 大于 0 的所有时间。

关于如何保存每个 i 和 tid 的 gcount 值的任何建议？在调试时，我发现对于 block (0,0,0) 和线程 (0,0,0)，我可以按顺序看到 t2 的值已更新。 CUDA 内核将焦点切换到 block(0,0,0) 和 thread(32,0,0) 后，out1[0] 的值再次被重写。如何获取/存储每个线程的 out1 值并将其写入输出？

到目前为止，我尝试了两种方法:(由 NVIDIA 论坛上的@paseolatis 建议)

(1) 定义offset=tid*32;并将 out1[gcount] 替换为 out1[offset+gcount],

(2) 定义

__device__ int totgcount=0; // this line before main()
atomicAdd(&totgcount,1);
out1[totgcount]=t2[i] * in1[tid];

int *h_xc = (int*) malloc(sizeof(int) * 1);
cudaMemcpyFromSymbol(h_xc, totgcount, sizeof(int)*1, cudaMemcpyDeviceToHost);
printf("GPU: xcount = %d\n", h_xc[0]); // Output looks like this: GPU: xcount = 1928669800

有什么建议吗？提前致谢 !

Answer 1

好的，让我们将您对代码应该做什么的描述与您发布的内容进行比较(有时称为 rubber duck debugging)。

1. 将 in2 的 32 个值存储在共享内存变量 t1

   你的内核包含这个:

   if (threadIdx.x < 32) {
       t1[threadIdx.x] = in2[i%posdir];
   }
   

   这有效地将 相同的值 从 in2 加载到 t1 的每个值中。我怀疑你想要更像这样的东西:

   if (threadIdx.x < 32) {
       t1[threadIdx.x] = in2[i+threadIdx.x];
   }
   

2. 对于i和in1[tid]的每个值，计算t2[i]，

   这部分没问题，但为什么共享内存中根本不需要t2？它只是一个中间结果，可以在内部迭代完成后丢弃。你可以很容易地拥有类似的东西:

   float inval = in1[tid];
   .......
   for(int i = 0; i < 32; i++)
   {
        float result = t1[i] * inval;
        ......
   

3. 如果 t2[i] > 0 对于 i 的特定组合，写t2[i]*in1[tid] 到 out1[gcount]

   这才是问题真正开始的地方。在这里你这样做:

           if(t2[i] > 0){
               out1[gcount] = t2[i] * in1[tid];
               gcount = gcount + 1;
           }
   

   这是一场内存竞赛。 gcount 是线程局部变量，因此每个线程都会在不同的时间用自己的值覆盖任何给定的 out1[gcount]。要让这段代码按照编写的方式正确工作，您必须拥有的是将 gcount 作为全局内存变量，并使用原子内存更新来确保每个线程使用唯一的 gcount< 值 每次输出一个值。但请注意，如果经常使用原子内存访问，它会非常昂贵(这就是为什么我在评论中询问每个内核启动有多少个输出点)。

生成的内核可能看起来像这样:

__device__ int gcount; // must be set to zero before the kernel launch

__global__ void test_g(float *in1, float *in2, float *out1, int posdir, int pos)
{
    int tid = threadIdx.x + blockIdx.x*blockDim.x;
    __shared__ float t1[32];

    float ival = in1[tid];

    for(int i = 0; i < posdir*pos; i += 32) {
        if (threadIdx.x < 32) {
            t1[threadIdx.x] = in2[i+threadIdx.x];
        }
        __syncthreads();

        for(int j = 0; j < 32; j++)
        {
            float tval = t1[j] * ival;
            if(tval > 0){
                int idx = atomicAdd(&gcount, 1);
                out1[idx] = tval * ival
            }
        }
    }        
}

免责声明:用浏览器编写，从未编译或测试，使用风险自负......

请注意，您对 ct 的写入也是一场内存竞争，但是现在 gcount 是一个全局值，您可以在内核之后读取该值而不需要 ct .

---

编辑:看来您在运行内核之前将 gcount 清零时遇到了一些问题。为此，您需要使用 cudaMemcpyToSymbol 或 cudaGetSymbolAddress 和 cudaMemset 之类的东西。它可能看起来像:

const int zero = 0;
cudaMemcpyToSymbol("gcount", &zero, sizeof(int), 0, cudaMemcpyHostToDevice);

同样，通常的免责声明:用浏览器编写，从未编译或测试，使用风险自负......