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

Answer 1

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

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

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