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 结构或类似的东西:数据可以组织什么我所需要的。这一点是，每个阵列，平均值，标准偏差并最终的直方图需要计算的，并在一个特定的对象，比率之间的差异阵列需要加以计算。

我应该如何去关于发送这种数据的技术设备和组织我的线块层次结构？ 一个想法我有的是零垫我所有阵列，以便他们有相同的长度，并有一组块的工作对每个对象，但似乎有各种各样的问题，方法是否会在所有的工作。

在此先感谢

Answer 1

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

有许多可能的途径来此，一个影响因素将是您的系统对象的总数。我假设数目是大的。

为了获得最佳性能要最佳存储器访问模式和要避免发散。由于全等阵列的数量为四，如果你把做每线程一个数组的简单方法，下面，你不仅会从具有存储器存取欠佳，而且H / W将需要检查在每次迭代该线程经需要执行循环 - 。那些不将被禁用，其可能是低效的（特别是如果一个阵列比其它更长，例如）

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作为重要的发展。

在这个例子中，你将推出每块中的四个经纱，即128个线程，并且尽可能多的块作为你有对象。

[1]你说，你可以选择您喜欢的任何存储器装置，常常也可以是交错的阵列，使得数组[0] [0]是旁边阵列有用[1] [0]，因为这将意味着相邻线程可以在相邻的阵列操作并且得到聚结的访问。然而，由于阵列的长度不是恒定的，这是可能复杂，需要更短的阵列的填充。

Answer 2

作为后续行动的汤姆的答案，我想提及的是 经减少 可以很容易实现的 幼崽.

这里是工作的例子：

#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];
....