CUDA：使用线性化二维共享内存对数组中的所有元素求和

Question

我是 CUDA 和一般算法的新手。有人可以告诉我我这样做是否正确，或者是否有更好的方法。一个问题是代码的输入和输出应该在 GPU 上，这样主机和设备之间就没有内存复制。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <stdint.h>

#include <iostream>

#define TILE_WIDTH 8

__global__ void gpu_sumElements(int height, int width, float *in, float *out){

    extern __shared__ float cache[];

    int w = blockIdx.x * blockDim.x + threadIdx.x; // Col // width
    int h = blockIdx.y * blockDim.y + threadIdx.y;

    int index = h * width + w;
    int cacheIndex = threadIdx.y * blockDim.x + threadIdx.x;

    float temp = 0;

    if ((w < width) && (h < height)){
        temp += in[index];
        //index += (height * width);
    }

    cache[cacheIndex] = temp;
    __syncthreads();

    int i = (blockDim.x * blockDim.y) / 2;
    while (i != 0){
        if (cacheIndex < i)
            cache[cacheIndex] += cache[cacheIndex + i];
        __syncthreads();
        i /= 2;
    }

    if (cacheIndex == 0)
        out[blockIdx.y * gridDim.x + blockIdx.x] = cache[0];
}


int main(){
                                                                        // Initial Parameters       
    int width = 2363;
    int height = 781;
    float my_sum = 0;

    int block_height = (height - 1) / TILE_WIDTH + 1;
    int block_width = (width - 1) / TILE_WIDTH + 1;

    dim3 dimGrid(block_width, block_height, 1);
    dim3 dimBlock(TILE_WIDTH, TILE_WIDTH, 1);

    int sharedMemSize = TILE_WIDTH * TILE_WIDTH * sizeof(float);
                                                                        // Initialize host arrays
    float *test_array = new float[height * width];
    float *out = new float[height * width];
    for (int i = 0; i < (height * width); i++)
        test_array[i] = 1.0f;
                                                                        // Initialize device arrays 
    float *gpu_temp_array;
    float *gpu_out;
    cudaMalloc((void **)&gpu_temp_array, (height * width * sizeof(float)));
    cudaMalloc((void **)&gpu_out, (height * width * sizeof(float)));
    cudaMemcpy(gpu_out, test_array, (height * width * sizeof(float)), cudaMemcpyHostToDevice);

                                                                        // Copy these, need them elsewhere
    float sum_height = height; 
    float sum_width = width ;
    dim3 sum_dimGrid = dimGrid;
    int i = (height * width); 

                                                                        // Launch kernel, get & print results
    while (i != 0){
        gpu_sumElements<<<sum_dimGrid, dimBlock, sharedMemSize>>>(sum_height, sum_width, gpu_out, gpu_temp_array);
        cudaMemcpy(gpu_out, gpu_temp_array, (sum_height * sum_width * sizeof(float)), cudaMemcpyDeviceToDevice);
        cudaMemset(gpu_temp_array, 0, (height * width * sizeof(float)));

        sum_height = ceil(sum_height/TILE_WIDTH);
        sum_width = ceil(sum_width/TILE_WIDTH);;
        sum_dimGrid.x = (sum_width - 1) / TILE_WIDTH + 1; 
        sum_dimGrid.y = (sum_height - 1) / TILE_WIDTH + 1;
        i /= TILE_WIDTH*TILE_WIDTH;
    }

    cudaMemcpy(out, gpu_out, (height * width * sizeof(float)), cudaMemcpyDeviceToHost);
    std::cout << out[0] << std::endl << std::endl;  

    delete[] test_array;
    delete[] out;
    cudaFree(gpu_out);
    cudaFree(gpu_temp_array);

    system("pause");
    return 0;
}

Answer 1

一般来说，通常不需要使用多个内核启动来产生一个（最终）结果的并行缩减。 cuda 示例代码和随附的 PDF 详细记录了生成组织良好的并行缩减的过程，该过程只需要为任意数据大小启动两次内核。

要创建仅使用单个内核启动的并行缩减，至少有两种常见方法：

1. 使用所谓的“threadfence reduction”方法。这也在 CUDA 示例代码 中捕获。在这种方法中，最后的减少阶段是通过跟踪“内核耗尽”来执行的。具体来说，每个线程块在完成其工作时更新一个“完成计数”变量（原子地）。由于启动的线程块的数量是已知的，因此线程块可以确定它是否是最后一个完成的线程块。如果是，则该线程块会将其他线程块产生的所有中间结果相加，这些结果现在将写入全局内存。 “threadfence”绰号是由于每个线程块必须确保其部分结果在退出之前在全局内存中可用（使用 threadfence intrinsic）。这种方法可以处理“任意”减少。

2. 让（单个线程中的）每个线程块自动更新最终的内核范围结果，使用它自己的部分结果。这仅适用于提供相应原子函数的归约，例如求和、求最大值、求最小值等

上述任何一种方法都将受益于 CUDA 并行缩减示例代码中涵盖的基本技术，特别是将线程块的数量减少到仍然允许充分利用 GPU 的最小值。这种优化允许最少数量的原子操作。考虑到这些优化，与相应的 2 内核或多内核缩减相比，缩减可以更快、更“简单”（例如，单个内核调用，中间结果没有太多主机管理）。