У меня проблема с OpenCL, и я надеюсь, что кто-нибудь подскажет, в чем причина. Ниже приведена версия программы, сведенная к проблеме. У меня есть массив ввода int размером 4000. В моем ядре я делаю сканирование. Очевидно, есть хорошие способы сделать это параллельно, но чтобы воспроизвести проблему, только один поток выполняет все вычисления. Перед сканированием входные данные (result_mask) имеют только значения 0 или 1.
__kernel void
sel_a(__global db_tuple * input,
__global int * result_mask,
__global int * result_count,
const unsigned int max_id)
{
// update mask based on input in parallel
mem_fence(CLK_GLOBAL_MEM_FENCE);
if(gid == 0)
{
int i, c = 0;
for(i = 0; i < max_id; i++)
{
if(result_mask[i]!=0)
{
c++;
result_mask[i] = 5;
}
else
{
result_mask[i] = 5;
}
}
*result_count = c;
}
}
Ожидаемым результатом будет количество элементов, которые изначально имели значение, отличное от 0, и только 5 в маске результата. Однако это не так. Вывод выглядит так:
...
5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5
5 5 5 5 5 5 5 5 5 0 0 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 0 1 0 0 0 1 0 0 0 1 0 0 0 0 0 0 0 0 0 0
0 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 5 5 5 5 5 5 5 5 5 5 5
5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5
...
Я получаю этот блок из 80 элементов где-то после прибл. 3200 элементов. Это не всегда одинаковые позиции, но всегда одинаковое количество элементов - 80. И это становится еще более странным - если я изменю первую строку на
если (gid == 2000)
проблема ушла Однако, поиграв с идентификатором потока, я пришел к выводу, что проблема не исчезла, она просто переместилась. Используя поток 1425, я получаю проблему половину времени, и когда я получаю ее, глючный блок находится в конце массива. Следовательно, я предполагаю, что когда у меня нет 0 и 1, блок «переместился» еще дальше назад. Для еще большего волнения - когда я увеличиваю размер ввода до 5000, вывод полностью состоит из 0 с. Кроме того, следующий код не будет работать:
if(gid == 0)
{
int i, c = 0;
for(i = 0; i < max_id; i++)
{
if(result_mask[i]!=0)
{
c++;
result_mask[i] = 5;
}
else
{
result_mask[i] = 5;
}
}
*result_count = c;
}
if(gid == 3999)
{
int i, c = 0;
for(i = 0; i < max_id; i++)
{
if(result_mask[i]!=0)
{
c++;
result_mask[i] = 5;
}
else
{
result_mask[i] = 5;
}
}
*result_count = c;
}
, тогда как только
if(gid == 3999)
{
int i, c = 0;
for(i = 0; i < max_id; i++)
{
if(result_mask[i]!=0)
{
c++;
result_mask[i] = 5;
}
else
{
result_mask[i] = 5;
}
}
*result_count = c;
}
будет работать (опять же, вероятно, при большем входе, он может не работать). Ниже приведены некоторые сведения об устройстве:
Device name: GeForce 9600M GT
Device vendor: NVIDIA
Clock frequency: 1250 MHz
Max compute units: 4
Global memory size: 256 MB
Local memory size:. 16 KB
Max memory allocation size: 128 MB
Max work group size: 512
Очевидно, я упускаю что-то большое здесь. Сначала я подумал, что это какой-то конфликт памяти, когда блок из 80 элементов переопределяется другим «потоком». Но чем больше я об этом думаю, тем меньше в этом смысла.
Я буду очень благодарен за любые подсказки!
Благодаря.
EDIT:
Извините за поздний ответ. Поэтому я изменил код, сведя его к минимуму, чтобы воспроизвести проблему. Ниже приведен c-код программы:
#include <stdio.h>
#include <stdlib.h>
#include <OpenCL/openCL.h>
#define INPUTSIZE (200)
typedef struct tag_openCL
{
cl_device_id device;
cl_context ctx;
cl_command_queue queue;
cl_program program;
} openCL;
int main(void)
{
int err;
openCL* cl_ctx = malloc(sizeof(openCL));
if(!cl_ctx)
exit(1);
err = clGetDeviceIDs(NULL, CL_DEVICE_TYPE_GPU, 1, &cl_ctx->device, NULL);
cl_ctx->ctx = clCreateContext(0, 1, &cl_ctx->device, clLogMessagesToStdoutAPPLE, NULL, &err);
cl_ctx->queue = clCreateCommandQueue(cl_ctx->ctx, cl_ctx->device, CL_QUEUE_PROFILING_ENABLE, &err);
printf("Successfully created context and queue for openCL device. \n");
/* Build program */
char * kernel_source = "__kernel void \
sel(__global int * input, \
__global int * result_mask, \
const unsigned int max_id) \
{ \
int gid = get_global_id(0); \
\
result_mask[gid] = input[gid] % 2 == 0; \
result_mask[gid] &= (input[gid] + 1) % 3 == 0; \
\
if(gid == 0) { \
int i; \
for(i = 0; i < max_id; i++) { \
if(result_mask[i]) { \
result_mask[i] = 5; \
} \
else { \
result_mask[i] = 5; \
} \
} \
} \
}";
cl_program prog = clCreateProgramWithSource(cl_ctx->ctx, 1, (const char**)&kernel_source, NULL, &err);
cl_ctx->program = prog;
err = clBuildProgram(cl_ctx->program, 0, NULL, NULL, NULL, NULL);
cl_kernel kernel = clCreateKernel(cl_ctx->program, "sel", &err);
/* create dummy input data */
int * input = calloc(sizeof(int), INPUTSIZE);
int k;
for(k = 0; k < INPUTSIZE; k++)
{
input[k] = abs((k % 5) - (k % 3))+ k % 2;
}
cl_mem source, intermediate;
unsigned int problem_size = INPUTSIZE;
source = clCreateBuffer(cl_ctx->ctx, CL_MEM_READ_WRITE, problem_size * sizeof(int), NULL, NULL);
clEnqueueWriteBuffer(cl_ctx->queue, source, CL_TRUE, 0, problem_size * sizeof(int), (void*) input, 0, NULL, NULL);
intermediate = clCreateBuffer(cl_ctx->ctx, CL_MEM_READ_WRITE, problem_size * sizeof(int), NULL, NULL);
int arg = 0;
clSetKernelArg(kernel, arg++, sizeof(cl_mem), &source);
clSetKernelArg(kernel, arg++, sizeof(cl_mem), &intermediate);
clSetKernelArg(kernel, arg++, sizeof(unsigned int), &problem_size);
size_t global_work_size = problem_size;
size_t local_work_size = 1;
clEnqueueNDRangeKernel(cl_ctx->queue, kernel, 1, NULL, &global_work_size, &local_work_size, 0, NULL, NULL);
clFinish(cl_ctx->queue);
// read results
int * result = calloc(sizeof(int), problem_size );
clEnqueueReadBuffer(cl_ctx->queue, intermediate, CL_TRUE, 0, problem_size * sizeof(int), result, 0, NULL, NULL);
clFinish(cl_ctx->queue);
int j;
for(j=1; j<=problem_size; j++)
{
printf("%i \t", result[j-1]);
if(j%10 ==0 && j>0)
printf("\n");
}
return EXIT_SUCCESS;
}
Результат по-прежнему недетерминирован, я получаю 0 и 1 в случайных позициях на выходе. Для локального размера рабочей группы 1 они находятся в первой половине массива, для размера 2 - во второй половине, для размера 4 это выглядит нормально для 200 элементов, но снова есть 0 и 1 для размер проблемы 400. Кроме того, для глобальной рабочей группы размером 1 все работает нормально. То есть, если я использую два ядра - одно для параллельных вычислений с глобальным размером рабочей группы [размер задачи], а второе с глобальным размером рабочей группы 1, все прекрасно работает. Опять же, я прекрасно понимаю, что это не способ сделать это (ядро, выполняющее такой последовательный код), однако, я хотел бы знать, почему это не работает, так как, похоже, я что-то упустил.
Спасибо,
Васил