0

我 tr 分配一个 cuda 全局内存数组。我的求和内核为：

__device__ float R_d = 0;
__global__ void perform_summation(float* A, int N){
    int idx = blockDim.x*blockIdx.x+threadIdx.x;
    extern __shared__ float sharedArray [];
    float result[]; //THIS IS THE THING i TRIED TO CREATE
    if(idx < N){
        sharedArray[threadIdx.x] = A[idx];
        //  }else{
        //      sharedArray[threadIdx.x] = 0 ;
        //  }

        for (unsigned int stride = 1; stride < blockDim.x; stride *= 2) {
            __syncthreads();
            if(threadIdx.x % (2*stride) == 0){
                sharedArray[threadIdx.x]+=sharedArray[threadIdx.x+stride];
            }
        }
    }
    if(idx % blockDim.x == 0){
//      R_d += sharedArray[threadIdx.x];
        result[blockIdx.x] = sharedArray[threadIdx.x];
    }
    for (int i = 0; i < gridDim.x; ++i) {
        R_d += result[i];
    }
}

作为总结，y内核采用一个数组并通过map reduce方法找到元素的总和。每个块将相关元素放入共享内存并将其中的所有数据相加，然后将结果放入我尝试创建的全局数组中。最后，我会将全局数组的所有数字相加，以找到最后的答案。

作为第一种方法，我没有使用全局数组来收集每个块的结果，我只是将块的结果汇总到变量中R_d，但它不起作用并且只显示来自最后一个块的值作为结果。我想因为我没有同步。块之间最后一个块覆盖最后的所有值。这是我在内核末尾第一次尝试时所做的

f(idx < N){
        sharedArray[threadIdx.x] = A[idx];
        //  }else{
        //      sharedArray[threadIdx.x] = 0 ;
        //  }

        for (unsigned int stride = 1; stride < blockDim.x; stride *= 2) {
            __syncthreads();
            if(threadIdx.x % (2*stride) == 0){
                sharedArray[threadIdx.x]+=sharedArray[threadIdx.x+stride];
            }
        }

        if(threadIdx.x == 0){
            R_d += sharedArray[threadIdx.x];
        }
    }

所以我实际上有两个问题。如何为我提出的第一个解决方案在设备内存中定义一个全局内存数组，对于仅使用该R_d变量的第二个解决方案是否有任何解决方案？

4

1 回答 1

1

您可以通过cudaMalloc在全局设备内存中分配数组：

cudaMalloc((void **)&ptr, size);

但是您不想在内核内部执行此操作，而是在调用内核并将指针传递给内核之前执行此操作。

至于减少，看看这些 nVidia 幻灯片，它解释得很好。基本上，这取决于您使用了多少块和线程。可以说有几个块。所以在共享内存中定义一个数组：

__shared__ float cache[BLOCK_THREADS];

为每个块分配共享内存，因此我们将每个块中的值与cache.

__syncthreads();

int step = (BLOCK_THREADS >> 1); //the same result as BLOCK_THREADS/2
while(step > 0) {
    if (threadInBlock < step) {
        cache[threadInBlock] += cache[threadInBlock + step];
    }
    __syncthreads();
    step = (step >> 1);
}

因此，这将每个块中的所有元素相加到cache[0]. 现在我们可以再次使用归约，或者我们可以使用原子操作将每个块的所有总和相加。如果每个块的块数明显少于线程数，这将是可以的。

__syncthreads();
if (threadInBlock == 0) {
    atomicAdd(result, cache[0]);
}

请注意，它result是指向全局内存中单个值的指针。另请注意，这仅BLOCK_THREADS在 2 的幂时才有效 - 这很常见，因为每个块的线程数应该是 32 的倍数（与扭曲对齐）。

于 2013-04-14T12:55:42.157 回答