Блок сокращения в CUDA
Я пытаюсь сделать сокращение в CUDA, и я действительно новичок. В настоящее время я изучаю пример кода от NVIDIA.
Я думаю, я действительно не уверен, как установить размер блока и размер сетки, особенно когда мой входной массив больше (512 X 512
), чем один размер блока.
Вот код
template <unsigned int blockSize>
__global__ void reduce6(int *g_idata, int *g_odata, unsigned int n)
{
extern __shared__ int sdata[];
unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*(blockSize*2) + tid;
unsigned int gridSize = blockSize*2*gridDim.x;
sdata[tid] = 0;
while (i < n)
{
sdata[tid] += g_idata[i] + g_idata[i+blockSize];
i += gridSize;
}
__syncthreads();
if (blockSize >= 512) { if (tid < 256) { sdata[tid] += sdata[tid + 256]; } __syncthreads(); }
if (blockSize >= 256) { if (tid < 128) { sdata[tid] += sdata[tid + 128]; } __syncthreads(); }
if (blockSize >= 128) { if (tid < 64) { sdata[tid] += sdata[tid + 64]; } __syncthreads(); }
if (tid < 32)
{
if (blockSize >= 64) sdata[tid] += sdata[tid + 32];
if (blockSize >= 32) sdata[tid] += sdata[tid + 16];
if (blockSize >= 16) sdata[tid] += sdata[tid + 8];
if (blockSize >= 8) sdata[tid] += sdata[tid + 4];
if (blockSize >= 4) sdata[tid] += sdata[tid + 2];
if (blockSize >= 2) sdata[tid] += sdata[tid + 1];
}
if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}
Тем не менее, мне кажется, g_odata[blockIdx.x]
сохраняет частичные суммы из всех блоков, и, если я хочу получить окончательный результат, мне нужно сложить все условия в пределах g_odata[blockIdx.x]
массив.
Мне интересно: есть ли ядро для полного суммирования? или я здесь неправильно понимаю вещи? Буду очень признателен, если кто-нибудь сможет научить меня этому. Спасибо большое.
3 ответа
Чтобы иметь лучшее представление об этой теме, вы можете взглянуть на этот PDF-файл NVIDIA, который графически объясняет все стратегии, которые вы использовали в своем коде.
Ваше понимание верно. Показанные здесь сокращения в итоге приводят к последовательности блочных сумм, хранящихся в глобальной памяти.
Для суммирования всех этих блочных сумм требуется некоторая форма глобальной синхронизации. Вы должны дождаться завершения всех блоков, прежде чем сложить их суммы. На данный момент у вас есть несколько вариантов, некоторые из которых:
- запустить новое ядро после основного ядра для суммирования блок-сумм
- добавить суммы блоков на хосте
- использовать атомикс для сложения сумм блоков в конце основного ядра
- используйте такой метод, как уменьшение потока, чтобы сложить суммы блоков в основном ядре.
Если вы будете искать по тегу CUDA, вы сможете найти примеры всего этого и обсуждения их плюсов и минусов. Чтобы увидеть, как размещенное вами основное ядро используется для полного сокращения, посмотрите пример кода параллельного сокращения.
Роберт Кровелла уже ответил на этот вопрос, который в основном о понимании, а не производительности.
Тем не менее, для всех тех, кто сталкивается с этим вопросом, я просто хочу подчеркнуть, что CUB делает доступными функции сокращения блоков. Ниже я приведу простой проработанный пример использования CUB BlockReduce
,
#include <cub/cub.cuh>
#include <cuda.h>
#include "Utilities.cuh"
#include <iostream>
#define BLOCKSIZE 32
const int N = 1024;
/**************************/
/* BLOCK REDUCTION KERNEL */
/**************************/
__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;
}
/********/
/* MAIN */
/********/
int main() {
// --- Allocate host side space for
float *h_data = (float *)malloc(N * sizeof(float));
float *h_result = (float *)malloc((N / BLOCKSIZE) * sizeof(float));
float *d_data; gpuErrchk(cudaMalloc(&d_data, N * sizeof(float)));
float *d_result; gpuErrchk(cudaMalloc(&d_result, (N / BLOCKSIZE) * sizeof(float)));
for (int i = 0; i < N; i++) h_data[i] = (float)i;
gpuErrchk(cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice));
sum<<<iDivUp(N, BLOCKSIZE), BLOCKSIZE>>>(d_data, d_result);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
gpuErrchk(cudaMemcpy(h_result, d_result, (N / BLOCKSIZE) * sizeof(float), cudaMemcpyDeviceToHost));
std::cout << "output: ";
for(int i = 0; i < (N / BLOCKSIZE); i++) std::cout << h_result[i] << " ";
std::cout << std::endl;
gpuErrchk(cudaFree(d_data));
gpuErrchk(cudaFree(d_result));
return 0;
}
В этом примере массив длины N
создан, и результатом является сумма 32
последовательные элементы. Так
result[0] = data[0] + ... + data[31];
result[1] = data[32] + ... + data[63];
....