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

ryhaxcpt  于 2023-05-06  发布在  其他
关注(0)|答案(2)|浏览(226)

看看Mark Harris' reduction example,我想看看是否可以让线程存储中间值而不进行归约操作:
例如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中,
1.对于iin1[tid]的每个值,计算t2[i]

  1. if t2[i] > 0对于i的特定组合,将t2[i]*in1[tid]写入out1[gcount]
    但我的输出全错了。我甚至无法得到t2[i]大于0的所有次数的计数。
    对于如何为每个itid保存gcount的值,有什么建议吗?在调试时,我发现对于块(0,0,0)和线程(0,0,0),我可以顺序地看到t2的值被更新。CUDA内核将焦点切换到块(0,0,0)和线程(32,0,0)后,out1[0]的值将再次重写。如何获取/存储每个线程的out1值并将其写入输出?
    到目前为止,我尝试了两种方法:(由NVIDIA论坛上的@paseolatis建议)
    1.定义offset=tid*32;并用out1[offset+gcount]替换out1[gcount]
    1.定义的
__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

有什么建议吗?先谢谢你了!

41zrol4v

41zrol4v1#

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

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

这是有效地加载 * 相同的值 * 从in2t1的每个值。我猜你想要的东西更像这样:

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

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

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

1.如果对于i特定组合为t2[i] > 0,则将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进行归零时遇到了一些问题。要做到这一点,您需要使用像cudaMemcpyToSymbolcudaGetSymbolAddresscudaMemset这样的东西。它可能看起来像这样:

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

再一次,常见的免责声明:在浏览器中编写,从未编译或测试过,使用风险自担.....

sczxawaw

sczxawaw2#

一个更好的方法是为每个线程提供自己的输出,让它递增自己的count并输入值-这样,双for循环可以以任何顺序并行发生,这是GPU擅长的。输出是错误的,因为线程共享out 1数组,所以它们都将覆盖它。
您还应该将要复制到共享内存中的代码移动到一个单独的循环中,后面加上__syncthreads()。在__syncthreads()退出循环后,您应该会获得更好的性能-这意味着您的共享阵列必须是in 2的大小-如果这是一个问题,在这个答案的最后有一个更好的方法来处理这个问题。
您还应该将threadIdx.x < 32检查移到外部。所以你的代码看起来像这样:

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

for(int i = threadIdx.x; i < posdir*pos; i += 32) {
    for(int j = 0; j < 32; j++)
    {
         ...
    }
}

然后放入一个__syncthreads(),一个gcount += count的原子加法,以及一个从本地输出数组到全局输出数组的副本-这部分是顺序的,会影响性能。如果可以的话,我只需要为每个本地数组创建一个全局指针列表,然后将它们放在CPU上。
另一个变化是t2不需要共享内存--这对您没有帮助。你这样做的方式,似乎只有在你使用单个块时才有效。为了获得大多数NVIDIA GPU的良好性能,您应该将其划分为多个块。您可以根据共享内存约束来定制它。当然,块之间没有__syncthreads(),因此每个块中的线程必须遍历内部循环的整个范围,以及外部循环的一个分区。

相关问题