warp shuffling для уменьшения массивов с любой длиной

Question

Я работаю над ядром Cuda, которое выполняет векторный точечный продукт (A x B). Я предположил, что длина каждого вектора кратно 32 (32,64, ...) и определяет размер блока равным длине массива. Каждый поток в блоке умножает один элемент A на соответствующий элемент B (поток i ==> psum = A [i] xB [i]). После умножения я использовал следующие функции, которые использовали метод перетаскивания основы для выполнения сокращения и вычисления суммы всех умножений.

__inline__ __device__ 
float warpReduceSum(float val) { 
    int warpSize =32; 
    for (int offset = warpSize/2; offset > 0; offset /= 2) 
     val += __shfl_down(val, offset); 
    return val; 
} 

__inline__ __device__ 
float blockReduceSum(float val) { 
    static __shared__ int shared[32]; // Shared mem for 32 partial sums 
    int lane = threadIdx.x % warpSize; 
    int wid = threadIdx.x/warpSize; 
    val = warpReduceSum(val);   // Each warp performs partial reduction 
    if (lane==0) 
     shared[wid]=val;    // Write reduced value to shared memory 
    __syncthreads();     // Wait for all partial reductions 
    //read from shared memory only if that warp existed 
    val = (threadIdx.x < blockDim.x/warpSize) ? shared[lane] : 0; 
    if (wid==0) 
     val = warpReduceSum(val);  // Final reduce within first warp 
    return val; 
} 

Я просто называю blockReduceSum (Psum), который Psum является умножение двух элементов на волоске.

Этот подход не работает, когда длина массива не кратная 32, поэтому мой вопрос заключается в том, можем ли мы изменить этот код, чтобы он также работал на любую длину? или это невозможно, потому что, если длина массива не кратная 32, некоторые перекосы имеют элементы, принадлежащие более чем одному массиву?

Answer 1

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

Чтобы ответить на ваш вопрос: вы можете использовать код, который вы написали, просто вызывая ядро ​​с числом нитей, являющегося ближайшим кратным 32 выше, чем N (длина массива) и вводя if заявление перед обращением к blockReduceSum что хотел бы это:

__global__ void kernel(float * A, float * B, int N) { 
    float psum = 0; 
    if(threadIdx.x < N) //threadIDx.x because your are using single block, you will need to change it to more general id once you move to multiple blocks 
     psum = A[threadIdx.x] * B[threadIdx.x]; 
    blockReduceSum(psum); 
    //The rest of computation 
} 

Таким образом, потоки, которые не имеют элемент массива, связанный с ними, но что нужно быть там из-за использования __shfl, будет способствовать 0 на сумму.