c - 在CUDA中使用共享内存而不减少线程

看看Mark Harris的reduction示例，我想看看是否可以让线程在不执行reduction操作的情况下存储中间值：
例如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）对于i的特定组合，将if t2[i] > 0写入t2[i]*in1[tid]
但我的输出全错了。我甚至无法计算出t2[I]大于0的所有次数。
对于如何保存每个i和tid的gcount值有何建议？？当我调试时，我发现对于block（0,0,0）和thread（0,0,0），我可以依次看到t2的值被更新。在CUDA内核将焦点切换到块（0,0,0）和线程（32,0,0）之后，out1[0]的值将重新写入。如何获取/存储每个线程的out1值并将其写入输出？
到目前为止，我尝试了两种方法：（由NVIDIA论坛上的@paseolatis建议）
（1）定义out1[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

有什么建议吗？提前谢谢！

最佳答案

好吧，让我们比较一下您对代码应该做什么的描述和您发布的内容（有时称为rubber duck debugging）。
在共享内存变量中存储32个in2值t1
您的内核包含以下内容：

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

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

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

对于i和in1[tid]的每个值，计算t2[i]，
这部分还可以，但是为什么共享内存中需要t2？它只是一个中间结果，可以在内部迭代完成后丢弃。你可以很容易地得到这样的东西：

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

如果我的特定组合是
t2[i] > 0到t2[i]*in1[tid]
这才是问题真正开始的地方。在这里您可以执行以下操作：

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

这是一场记忆竞赛。out1[gcount]是一个线程局部变量，因此每个线程将在不同的时间用自己的值覆盖任何给定的gcount。要使代码在编写时正常工作，必须将out1[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
            }
        }
    }
}

免责声明：在浏览器中编写，从未经过编译或测试，使用风险自负。。。。。
注意，您对gcount的写入也是一个内存竞赛，但是现在gcount是一个全局值，您可以在内核之后读取该值，而不需要ct。
编辑：在运行内核之前，您似乎有一些归零ct的问题。为此，您需要使用gcount或cudaMemcpyToSymbol和cudaGetSymbolAddress之类的内容。可能看起来像：

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

再次重申，通常的免责声明：写在浏览器中，从未被编译或测试，使用风险自负。。。。。

关于c - 在CUDA中使用共享内存而不减少线程，我们在Stack Overflow上找到一个类似的问题：https://stackoverflow.com/questions/10285718/