CUDA减少许多小的，不相等大小的阵列

Question

我想知道是否有人可以建议最好的方法来计算大量相对较小的平均值/标准偏差，但是在CUDA中的不同大小的阵列？

SDK中平行下降示例工作一个非常大的阵列上，似乎大小为方便每块的线程数的倍数，但我的情况有点不同：

概念，然而，我具有大量的对象，每个对象包含两个组件，分别为upper和lower，并且这些组件中的每一个都具有x和y坐标。即

upper.x, lower.x, upper.y, lower.y 

这些阵列中的每一个的长度为大约800但它的对象之间变化（未在对象内），例如

Object1.lower.x = 1.1, 2.2, 3.3 
Object1.lower.y = 4.4, 5.5, 6.6 
Object1.upper.x = 7.7, 8.8, 9.9 
Object1.upper.y = 1.1, 2.2, 3.3 

Object2.lower.x = 1.0, 2.0, 3.0, 4.0, 5.0 
Object2.lower.y = 6.0, 7.0, 8.0, 9.0, 10.0 
Object2.upper.x = 11.0, 12.0, 13.0, 14.0, 15.0 
Object2.upper.y = 16.0, 17.0, 18.0, 19.0, 20.0 

请注意以上仅仅是我代表阵列的方式和我的数据没有存储在C结构或类似的东西：数据可以以任何方式，我需要进行组织。重点是，对于每个阵列，需要计算均值，标准偏差和最终直方图，并且在一个特定对象内，需要计算数组之间的比率和差异。

我应该如何将这些数据发送到GPU设备并组织我的线程块层次结构？我想到的一个想法是将所有数组填充为零，以使它们具有相同的长度，并且在每个对象上都有一组块，但是如果它能够工作，似乎存在各种各样的问题。

在此先感谢

Answer 1

如果我理解正确的，你想Object1.lower.x减少到一个结果，Object1.lower.y另一个结果等。对于任何给定的对象，都有四个数组要缩小，所有长度都相等（对于该对象）。

有很多可能的方法来解决这个问题，其中一个影响因素是系统中的对象总数。我会假设这个数字很大。

为了获得最佳性能，您希望获得最佳内存访问模式，并且希望避免发散。因为一致数组的数量是4，所以如果你采用了每个线程在下面做一个数组的简单方法，那么你不仅会遭受较差的内存访问，而且h/w还需要检查每次迭代的线程数warp需要执行循环 - 那些不会被禁用的可能是低效的（特别是如果一个数组比其他数据长得多）。

for (int i = 0 ; i < myarraylength ; i++) 
    sum += myarray[i]; 

相反，如果你每经总结一个数组，那么不仅会更加高效，而且你的内存访问模式会更好，因为相邻的线程将读取的相邻元素[1]。

for (int i = tidwithinwarp ; i < warparraylength ; i += warpsize) 
{ 
    mysum += warparray[i]; 
} 
mysum = warpreduce(mysum); 

你也应该采取阵列的排列考虑，最好是64字节边界上对齐但如果你是在计算能力1.2或更高版本开发，那么这是不是很上了年纪的GPU一样重要。

在这个例子中，你将启动每个块四个warp，即128个线程和尽可能多的块。

[1]你的确可以说你可以选择你喜欢的任何内存布局，通常交错数组使得array [0] [0]紧挨array [1] [0]是有用的，因为这会意味着相邻的线程可以在相邻的阵列上运行并获得合并访问。但是由于数组的长度不是常数，这可能很复杂，需要填充较短的数组。

Answer 2

作为汤姆回答的后续，我想提一下，减量可以通过CUB轻松实现。

这里是一个工作实例：

#include <cub/cub.cuh> 
#include <cuda.h> 

#include "Utilities.cuh" 

#include <iostream> 

#define WARPSIZE 32 
#define BLOCKSIZE 256 

const int N = 1024; 

/*************************/ 
/* WARP REDUCTION KERNEL */ 
/*************************/ 
__global__ void sum(const float * __restrict__ indata, float * __restrict__ outdata) { 

    unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x; 

    unsigned int warp_id = threadIdx.x/WARPSIZE; 

    // --- Specialize WarpReduce for type float. 
    typedef cub::WarpReduce<float, WARPSIZE> WarpReduce; 

    // --- Allocate WarpReduce shared memory for (N/WARPSIZE) warps 
    __shared__ typename WarpReduce::TempStorage temp_storage[BLOCKSIZE/WARPSIZE]; 

    float result; 
    if(tid < N) result = WarpReduce(temp_storage[warp_id]).Sum(indata[tid]); 

    if(tid % WARPSIZE == 0) outdata[tid/WARPSIZE] = result; 
} 

/********/ 
/* MAIN */ 
/********/ 
int main() { 

    // --- Allocate host side space for 
    float *h_data  = (float *)malloc(N * sizeof(float)); 
    float *h_result  = (float *)malloc((N/WARPSIZE) * sizeof(float)); 

    float *d_data;  gpuErrchk(cudaMalloc(&d_data, N * sizeof(float))); 
    float *d_result; gpuErrchk(cudaMalloc(&d_result, (N/WARPSIZE) * sizeof(float))); 

    for (int i = 0; i < N; i++) h_data[i] = (float)i; 

    gpuErrchk(cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice)); 

    sum<<<iDivUp(N, BLOCKSIZE), BLOCKSIZE>>>(d_data, d_result); 
    gpuErrchk(cudaPeekAtLastError()); 
    gpuErrchk(cudaDeviceSynchronize()); 

    gpuErrchk(cudaMemcpy(h_result, d_result, (N/WARPSIZE) * sizeof(float), cudaMemcpyDeviceToHost)); 

    std::cout << "output: "; 
    for(int i = 0; i < (N/WARPSIZE); i++) std::cout << h_result[i] << " "; 
    std::cout << std::endl; 

    gpuErrchk(cudaFree(d_data)); 
    gpuErrchk(cudaFree(d_result)); 

    return 0; 
} 

在该示例中，创建长度N的阵列，并且其结果是连续32元素的总和。所以

result[0] = data[0] + ... + data[31]; 
result[1] = data[32] + ... + data[63]; 
....