0 
__global__ void sum(const float * __restrict__ indata, float * __restrict__ outdata) { 
    unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x; 
    // --- Specialize BlockReduce for type float. 
    typedef cub::BlockReduce<float, BLOCKSIZE> BlockReduceT; 
    // --- Allocate temporary storage in shared memory 
    __shared__ typename BlockReduceT::TempStorage temp_storage;    
    float result;
    if(tid < N) result = BlockReduceT(temp_storage).Sum(indata[tid]);   
    // --- Update block reduction value
    if(threadIdx.x == 0) outdata[blockIdx.x] = result;  
   return;  
}

我已经成功地用 cuda cub 测试了归约和（如上面的代码片段所示），我想根据这段代码执行两个向量的内积。但我对此有一些困惑：

  1. 我们需要两个用于 inner_product 的输入向量，我需要在对得到的新向量进行归约和之前对这两个输入向量进行逐个乘法运算。

  2. 在 cuda cub 的代码示例中，输入向量的维度等于 blocknumber*threadnumber。如果我们有一个非常大的向量怎么办。

4

1 回答 1

1

  1. 是的，使用 cub，并假设您的向量是单独存储的（即不交错存储），您需要先进行元素乘法。另一方面，thrust transform_reduce可以在单个函数调用中处理它。

  2. blocknumber*threadnumber 应该为您提供所需的所有范围。在 cc3.0 或更高版本的 GPU 上，blocknumber (ie gridDim.x) 的范围可达 2^31-1，threadnumber (ie blockDim.x) 的范围可达 1024。这使您可以处理 2^40 个元素。如果每个元素是 4 个字节，这将构成（即需要）2^42 个字节。这大约是 4TB（如果您考虑 2 个输入向量，则为两倍），这比目前任何 GPU 内存都要大得多。因此，在用完网格维度之前，您将用完 GPU 内存空间。

请注意，您显示的是cub::BlockReduce. 但是，如果您正在做两个大向量的向量点积，您可能想要cub::DeviceReduce改用。

于 2018-06-04T01:08:45.023 回答