Эта простая функция сокращения содержится в одной из онлайн-презентаций CUDA.
__device__ void reducedSum(double* d_idata, double* d_odata, long LENGTH)
{
extern __shared__ double sdata[];
unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < LENGTH) {
sdata[tid] = d_idata[i];
__syncthreads();
printf("Kernel sdata : %d \n", sdata[tid]);
for (unsigned int s = 1; s < blockDim.x; s *= 2)
{
if (tid % (2 * s) == 0)
{
sdata[tid] += sdata[tid + s];
}
__syncthreads();
}
if (tid == 0) {
d_odata[blockIdx.x] = sdata[0];
}
}
}
Но printf здесь всегда печатает следующий вывод. Ожидается, что на самом деле он скопирует значения из массива d_idata и назначит его частично каждому блоку общей памяти. Однако этого не происходит.
Вызов в ядре выглядит следующим образом:
long LENGTH = 10;
long N = 5;
int threadsPerBlock = N;
int numBlocks = (threadsPerBlock + LENGTH - 1) / threadsPerBlock;
cudaCalc<<<numBlocks, threadsPerBlock, N*sizeof(double)>>> (d_vec1, d_vec2, d_dotProduct, ....)
Теперь внутри ядра я вызываю эту функцию extendedSum __device__ следующим образом .
__global__ void cudaCalc(int* d_vec1, int* d_vec2, double* d_dotProduct, ... )
{
int tid_0 = threadIdx.x;
int index = blockDim.x * blockIdx.x + threadIdx.x;
if (index < LENGTH) {
d_dotProduct[index] = (double) d_vec1[index] * d_vec2[index];
d_squared1[index] = (double)d_vec1[index] * d_vec1[index];
d_squared2[index] = (double)d_vec2[index] * d_vec2[index];
__syncthreads();
}
reducedSum(d_squared1, d_squaredSum1, LENGTH);
reducedSum(d_squared2, d_squaredSum2, LENGTH);
reducedSum(d_dotProduct, d_dotSum, LENGTH);
}
Может, какой-нибудь хороший сэр / мадам, пожалуйста, покажите мне, где я ошибаюсь? Я был в этом часами. Если вы хотите увидеть остальную часть кода, пожалуйста, запрос. Заранее спасибо.