dot_product с CUDA_CUB

Question

__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. В примерах кода куба размерность входных векторов равна blocknumber * threadnumber. Что делать, если у нас есть очень большой вектор.

Answer 1

1. Да, с cub, и предполагая, что ваши векторы были сохранены отдельно (т.е. не чередуются), вам нужно было бы сначала выполнить поэлементное умножение. С другой стороны, thrust transform_reduce может обработать его за один вызов функции.

2. blocknumber * threadnumber должен дать вам весь необходимый диапазон. на GPU cc3.0 или выше, номер блока (то есть gridDim.x) может варьироваться до 2^31-1 и номер резьбы (т.е. blockDim.x) может варьироваться до 1024. Это дает вам возможность обрабатывать 2^40 элементов. Если каждый элемент составляет 4 байта, это будет составлять (т.е. потребовать) 2^42 байта. Это около 4 ТБ (или вдвое больше, если вы рассматриваете 2 входных вектора), что намного больше, чем любая память GPU в настоящее время. Таким образом, вы исчерпаете пространство памяти GPU, прежде чем исчерпаете размерность сетки.

Обратите внимание, что вы показываете cub::BlockReduce, Однако, если вы делаете произведение векторной точки из двух больших векторов, вы можете использовать cub::DeviceReduce вместо.