c++ - CUDA Vector Reduction 处理长度小于 512 的 vector ？

Question

我正在研究 NVIDIA 的并行 vector_reduction 算法教程，以使用 CUDA C++ API 实现该算法。我已经实现了该算法，但它仅适用于固定为 512 的 vector 长度。我无法弄清楚如何让它适用于小于 512 的 vector ？我希望它适用于任意大小，即 324、123、23。

#include <stdio.h>

#define NUM_ELEMENTS 512

__global__ void reduction(float *g_data, int n)
{
    __shared__ float partialSum[NUM_ELEMENTS];

    int tx = threadIdx.x;
    int i = tx + blockIdx.x * blockDim.x;

    if (i < n) {
        partialSum[tx] = g_data[i];
    }

    int stride;
    for (stride = blockDim.x/2; stride > 0;  stride >>= 1) {
        __syncthreads();
        if (tx < stride) {
           partialSum[tx] += partialSum[tx + stride];
        }
    }

    if (tx == 0) {
        g_data[blockIdx.x] = partialSum[tx];
    }
}

float computeOnDevice(float* h_data, int num_elements)
{
    float* d_data = NULL;
    float result;

    // Memory allocation on device side
    cudaMalloc((void**)&d_data, sizeof(float)*num_elements);

    // Copy from host memory to device memory
    cudaMemcpy(d_data, h_data, num_elements * sizeof(float), cudaMemcpyHostToDevice );

    dim3 blockSize, gridSize;

    // Number of threads in each thread block
    blockSize = dim3(num_elements, 1, 1);

    // Number of thread blocks in grid
    gridSize = dim3(1, 1, 1);

    // Invoke the kernel
    reduction<<<gridSize, blockSize>>>(d_data, num_elements);

    // Copy from device memory back to host memory
    cudaMemcpy(&result, d_data, sizeof(float), cudaMemcpyDeviceToHost);
    cudaFree(d_data);
    cudaDeviceReset();
    return result;
}

int main() {

    float *data = new float[NUM_ELEMENTS];
    for (int i = 0; i < NUM_ELEMENTS; i++) data[i] = 1;
    float r = computeOnDevice(data, NUM_ELEMENTS);
    printf(" result = %f\n" , r);
}

Answer 1

您的代码 100% 正确。问题是您的位移不考虑数组的最后一部分。您可以通过人为地将数组扩展到 2 的下一个幂来轻松解决此问题。这样您的整个数组就会减少，额外的“元素”(它们实际上并不存在)将被忽略。

#include <math.h>

__global__ void reduction(float *g_data, int n){
    // figure out exponent of next larger power of 2
    int exponent = ceilf(log2f(n));
    // calculate next larger power of 2
    int size = (int)powf(2, exponent);
    __shared__ float partialSum[NUM_ELEMENTS];

    int tx = threadIdx.x;
    int i = tx + blockIdx.x * blockDim.x;

    if (i < n){
        partialSum[tx] = g_data[i];
    }

    for (int stride = size / 2; stride > 0; stride >>= 1){
        __syncthreads();

        if (tx < stride) {
            // all threads that run out of bounds do nothing
            // equivalent to adding 0
            if((tx + stride) < n)
                partialSum[tx] += partialSum[tx + stride];
        }
    }

    if (tx == 0){
        g_data[blockIdx.x] = partialSum[tx];
    }
}

编辑

关于您的评论，这种归约方法永远不适用于在多个 block 中归约的数组。因此，对于计算能力 1.0-1.3，您可以减少的最大数组是 512 个元素，对于 >1.3 的计算能力，您最多可以减少 1024 个元素，这是每个 block 的最大线程数。

这是因为__shared__ 内存在线程之间共享而不是 block 。因此，要减少分散在多个 block 上的数组，您需要对数组进行分区，以便每个 block 减少一个 block ，然后利用 __global__ 内存来减少所有 block 的值。但是，__global__ 内存比(片上)__shared__ 内存慢大约 10-20 倍，因此一旦开始使用大量 block ，这将变得非常低效.

另一种方法是让每个线程处理多个索引，但是，最终您的 partialSum 数组将不再适合共享内存并溢出到全局内存中。这种方法也意味着您永远不能使用超过 512(或 1024)个线程，这违背了使用 CUDA 的目的，CUDA 依赖于运行大量线程来隐藏延迟并使从主机到设备的昂贵内存传输值得.