1

查看 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)if t2[i] > 0对于 i 的特定组合，t2[i]*in1[tid]写入out1[gcount]

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

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

到目前为止，我尝试了两种方法：（@paseolatis 在 NVIDIA 论坛上建议）

(1) 定义offset=tid*32; and replace out1[gcount] with 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

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

4

2 回答 2

2

好的，让我们将您对代码应该做什么的描述与您发布的内容进行比较（这有时称为橡皮鸭调试）。

  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);

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

于 2012-04-23T20:40:07.390 回答
1

做你正在做的事情的更好方法是给每个线程自己的输出，让它增加自己的输出count并输入值 - 这样，双 for 循环可以以任何顺序并行发生，这就是 GPU 所做的出色地。输出是错误的，因为线程共享 out1 数组，所以它们都会覆盖它。

您还应该将要复制到共享内存中的代码移动到单独的循环中，并在__syncthreads()后面加上。使用__syncthreads()循环外，您应该获得更好的性能 - 这意味着您的共享数组必须是 in2 的大小 - 如果这是一个问题，那么在这个答案的末尾有一个更好的方法来处理这个问题。

您还应该将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()块之间的块，所以每个块中的线程必须跨越整个范围内循环和外循环的分区。

于 2012-04-23T20:14:01.007 回答