У меня есть код для вычисления простых чисел, которые я распараллелил с использованием OpenMP:
#pragma omp parallel for private(i,j) reduction(+:pcount) schedule(dynamic)
for (i = sqrt_limit+1; i < limit; i++)
{
check = 1;
for (j = 2; j <= sqrt_limit; j++)
{
if ( !(j&1) && (i&(j-1)) == 0 )
{
check = 0;
break;
}
if ( j&1 && i%j == 0 )
{
check = 0;
break;
}
}
if (check)
pcount++;
}
Я пытаюсь перенести его на графический процессор, и я хотел бы уменьшить количество, как я сделал для примера OpenMPвыше.Ниже приведен мой код, который кроме предоставления неверных результатов также медленнее:
__global__ void sieve ( int *flags, int *o_flags, long int sqrootN, long int N)
{
long int gid = blockIdx.x*blockDim.x+threadIdx.x, tid = threadIdx.x, j;
__shared__ int s_flags[NTHREADS];
if (gid > sqrootN && gid < N)
s_flags[tid] = flags[gid];
else
return;
__syncthreads();
s_flags[tid] = 1;
for (j = 2; j <= sqrootN; j++)
{
if ( gid%j == 0 )
{
s_flags[tid] = 0;
break;
}
}
//reduce
for(unsigned int s=1; s < blockDim.x; s*=2)
{
if( tid % (2*s) == 0 )
{
s_flags[tid] += s_flags[tid + s];
}
__syncthreads();
}
//write results of this block to the global memory
if (tid == 0)
o_flags[blockIdx.x] = s_flags[0];
}
Прежде всего, как мне сделать это ядро быстрым, я думаю, что узким местом является цикл for, и я не уверенкак это заменить.И затем, мои подсчеты не верны.Я изменил оператор '%' и заметил некоторое преимущество.
В массиве flags я пометил простые числа от 2 до sqroot (N), в этом ядре я вычисляю простые числа из sqroot (N) к N, но мне нужно проверить, делится ли каждое число в {sqroot (N), N} на простые числа в {2, sqroot (N)}.В массиве o_flags хранятся частичные суммы для каждого блока.
РЕДАКТИРОВАТЬ: Следуя предложению, я изменил свой код (теперь я лучше понимаю, что такое комментарий к syncthreads);Я понял, что мне не нужен массив флагов, и в моем случае работают только глобальные индексы.На данный момент меня беспокоит медлительность кода (больше, чем корректность), которую можно отнести к циклу for.Кроме того, после определенного размера данных (100000) ядро выдавало неверные результаты для последующих размеров данных.Даже для данных размером менее 100000 результаты сокращения графического процессора неверны (участник форума NVidia отметил, что это может быть из-за того, что мой размер данных не равен степени 2).Итак, есть еще три (может быть связанных) вопроса -
Как я могу сделать это ядро быстрее?Является ли хорошей идеей использовать разделяемую память в моем случае, когда мне приходится циклически повторять каждый тид?
Почему он дает правильные результаты только для определенных размеров данных?
Как я могу изменить сокращение?
__global__ void sieve ( int *o_flags, long int sqrootN, long int N )
{
unsigned int gid = blockIdx.x*blockDim.x+threadIdx.x, tid = threadIdx.x;
volatile __shared__ int s_flags[NTHREADS];
s_flags[tid] = 1;
for (unsigned int j=2; j<=sqrootN; j++)
{
if ( gid % j == 0 )
s_flags[tid] = 0;
}
__syncthreads();
//reduce
reduce(s_flags, tid, o_flags);
}