代码之家 › 专栏 › 技术社区 › Pete

CUDA矢量缩减以处理长度小于512的矢量?

gpu-programming cuda c++

Pete · 技术社区 · 6 年前

我正在研究NVIDIA的并行向量降阶算法,用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);
}

0 回复 | 直到 6 年前

nick 6 年前

你的代码是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__ 内存以减少所有块中的值。然而, γ粒子 内存大约比(片上)慢10-20倍 第二股 内存,所以一旦你开始使用大量的块,这将变得非常低效。

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