或者有没有更好的方法，只使用一个reduction函数，但能够返回Struct.X、Struct.Y或struct.Z的总和？
通常加速计算的主要焦点是速度。GPU代码的速度（性能）通常在很大程度上取决于数据存储和访问模式。因此，尽管正如您在问题中指出的那样，我们可以通过多种方式实现解决方案，但让我们专注于应该相对较快的事情。
像这样的缩减没有太多的算术/操作强度，所以我们对性能的关注主要围绕着数据存储来进行有效的访问。当访问全局内存时，GPU通常会以大块的方式进行访问-- 32字节或128字节的块。为了有效地利用内存子系统，我们希望在每个请求中使用所有被请求的32或128个字节。
但你的结构隐含的数据存储模式：

template <typename T>
struct planet {
    T x, y, z;
    T vx, vy, vz;
    T mass;
};

基本排除了这个可能性对于这个问题，您关心的是vx、vy和vz。这3个项目在给定的结构（元素）中应该是连续的，但是在这些结构的数组中，它们将被其他结构项目所需的存储空间分隔开，至少：

planet0:       T x
               T y
               T z               ---------------
               T vx      <--           ^
               T vy      <--           |
               T vz      <--       32-byte read
               T mass                  |
planet1:       T x                     |
               T y                     v
               T z               ---------------
               T vx      <--
               T vy      <--
               T vz      <--
               T mass
planet2:       T x
               T y
               T z
               T vx      <--
               T vy      <--
               T vz      <--
               T mass

(for例如，假设T是float）
这指出了GPU中结构阵列（AoS）存储格式的一个关键缺点。由于GPU的访问粒度（32字节），从连续结构访问相同元素是低效的。在这种情况下，通常的性能建议是将AoS存储转换为SoA（阵列结构）：

template <typename T>
struct planets {
    T x[N], y[N], z[N];
    T vx[N], vy[N], vz[N];
    T mass[N];
};

上面只是一个可能的例子，可能不是你实际使用的，因为结构没有什么用处，因为我们只有一个N行星的结构。关键是，现在当我访问连续行星的vx时，各个vx元素在内存中都是相邻的，因此32字节的读取给了我32字节的vx数据，没有浪费或未使用的元素。
通过这样的转换，从代码组织的Angular 来看，缩减问题再次变得相对简单。您可以使用与单个数组归约代码基本相同的代码，要么连续调用3次，要么直接扩展内核代码，基本上独立处理所有3个数组。一个“3合1”内核可能看起来像这样：

template <typename T>
__global__ void reduceSum(T *input_vx, T *input_vy, T *input_vz, T *outdata_vx, T *outdata_vy, T *outdata_vz, int size){
    extern __shared__ T sdata[];

    const int VX = 0;
    const int VY = blockDim.x;
    const int VZ = 2*blockDim.x;

    unsigned int tID = threadIdx.x;
    unsigned int i = tID + blockIdx.x * (blockDim.x * 2);
    sdata[tID+VX] = input_vx[i] + input_vx[i + blockDim.x];
    sdata[tID+VY] = input_vy[i] + input_vy[i + blockDim.x];
    sdata[tID+VZ] = input_vz[i] + input_vz[i + blockDim.x];
    __syncthreads();

    for (unsigned int stride = blockDim.x / 2; stride > 32; stride >>= 1)
    {
        if (tID < stride)
        {
            sdata[tID+VX] += sdata[tID+VX + stride];
            sdata[tID+VY] += sdata[tID+VY + stride];
            sdata[tID+VZ] += sdata[tID+VZ + stride];
        }
        __syncthreads();
    }

    if (tID < 32){ warpReduce(sdata+VX, tID); }
    if (tID < 32){ warpReduce(sdata+VY, tID); }
    if (tID < 32){ warpReduce(sdata+VZ, tID); }

    if (tID == 0)
    {
        outdata_vx[blockIdx.x] = sdata[VX];
        outdata_vy[blockIdx.x] = sdata[VY];
        outdata_vz[blockIdx.x] = sdata[VZ];
    }
}

（在浏览器中编码-未测试-仅仅是您所显示的“参考内核”的扩展）
上面的AoS -> SoA数据转换可能也会在代码的其他地方带来性能上的好处。由于建议的内核将一次处理3个数组，因此网格和块的尺寸应该与您在单数组情况下使用的参考内核完全相同。每个块的共享内存存储将需要增加（三倍）。

Robert Crovella给出了一个很好的答案，强调了AoS -> SoA布局转换的重要性，这通常会提高GPU的性能，我只是想提出一个可能更方便的中间立场。CUDA语言提供了一些向量类型，用于您所描述的目的（请参阅CUDA编程指南的本节）。
例如，CUDA定义了int 3，一种存储3个整数的数据类型。

struct int3
 {
    int x; int y; int z;
 };

类似的类型也存在于float，chars，double等。这些数据类型的优点在于，它们可以用一条指令加载，这可能会给您带来一些性能提升。参见this NVIDIA blog post以获得对此的讨论。在这种情况下，它也是一种更“自然”的数据类型，它可能会使代码的其他部分更容易使用。例如，您可以定义：

struct planets {
    float3 position[N];
    float3 velocity[N];
    int mass[N];
};

使用这种数据类型的约简内核可能看起来像这样（改编自Robert的）。

__inline__ __device__ void SumInt3(int3 const & input1, int3 const & input2, int3 & result)
{
    result.x = input1.x + input2.x;
    result.y = input1.y + input2.y;
    result.z = input1.z + input2.z;
}

__inline__ __device__ void WarpReduceInt3(int3 const & input, int3 & output, unsigned int const tID)
{
    output.x = WarpReduce(input.x, tID);
    output.y = WarpReduce(input.y, tID);
    output.z = WarpReduce(input.z, tID);    
}

__global__ void reduceSum(int3 * inputData, int3 * output, int size){
    extern __shared__ int3 sdata[];

    int3 temp;

    unsigned int tID = threadIdx.x;
    unsigned int i = tID + blockIdx.x * (blockDim.x * 2);

    // Load and sum two integer triplets, store the answer in temp.
    SumInt3(input[i], input[i + blockDim.x], temp);

    // Write the temporary answer to shared memory.
    sData[tID] = temp;

    __syncthreads();

    for (unsigned int stride = blockDim.x / 2; stride > 32; stride >>= 1)
    {
        if (tID < stride)
        {
            SumInt3(sdata[tID], sdata[tID + stride], temp);
            sData[tID] = temp;
        }
        __syncthreads();
    }

    // Sum the intermediate results accross a warp.
    // No need to write the answer to shared memory,
    // as only the contribution from tID == 0 will matter.
    if (tID < 32)
    {
        WarpReduceInt3(sdata[tID], tID, temp);
    }

    if (tID == 0)
    {
        output[blockIdx.x] = temp;
    }
}