CUDA Vector Reduction可以处理长度小于512的向量？

Question

我正在研究NVIDIA的并行vector_reduction算法教程，以使用CUDA C ++ API实现该算法。我已经实现了算法，但它只适用于固定为512的向量长度。我无法弄清楚如何使它适用于小于512的向量？我希望它适用于任意大小，即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];
    }
}

编辑

关于你的评论，这种减少方法对于在多个块中减少的数组永远不会起作用。因此，对于计算能力1.0-1.3，可以减少的最大阵列是512个元素，对于计算能力> 1.3，最多可以执行1024个元素，这是每个块的最大线程数。

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

另一种方法是让每个线程处理多个索引，但是，最终你的partialSum数组将不再适合共享内存并且无论如何都会溢出到全局内存中。这种方法也意味着你永远不会使用超过512（或1024）个线程，这会破坏使用CUDA的目的，这取决于运行大量线程来隐藏延迟并使昂贵的内存从主机转移到设备值得。