Получение неправильного вывода в параллельном коде в OpenCL

Question

Я новичок в параллельном программировании. Я пытаюсь наложить на ПрефиксSum проблема в OpenCL. Но я получаю неправильный результат. Поэтому во время отладки я изменил свое ядро ​​для выполнения простой операции. Я работаю над 64-разрядной машиной Windows8 с графическим процессором AMD.

Вот мой код ядра -

__kernel void add(__global float *input, __global float *output, __global float *temp) 
{ 
    int thid = get_global_id(0); 
    int pout = 0; 
    int pin = 1; 
    temp[pin*8 + thid] = input[thid]; 
    temp[pout*8 + thid] = input[thid]; 
    pout = 1-pout; 
    pin = 1-pout; 
    int offset = 1; 

    if(thid >= offset) { 
     temp[pout*8 + thid] =temp[pout*8 + thid] + temp[pin*8 + thid - offset]; 
    } else { 
     temp[pout*8 + thid] = temp[pin*8 + thid]; 
    } 

    barrier(CLK_GLOBAL_MEM_FENCE); 
    output[thid] = temp[pout*8 + thid]; 
} 

И это мой код хост -

int main(void) 
{ 
cl_context context; 
cl_context_properties properties[3]; 
cl_kernel kernel; 
cl_command_queue command_queue; 
cl_program program; 
cl_int err; 
cl_uint num_of_platforms=0; 
cl_platform_id platform_id; 
cl_device_id device_id; 
cl_uint num_of_devices=0; 
cl_mem inputA,inputB, output; 
outfile.open("shubham.txt"); 
size_t global=8; 

float inputDataA[DATA_SIZE]={1, 2, 3, 4, 5, 6, 7, 8}; 
float results[DATA_SIZE]={0}; 
float inputDataB[16] = {0}; 
float shubh[16] = {0}; 
int i;//,j; 

//cl_int infoSize = 10000; 
//size_t infoSize; 
//char *info; 
// retreive a list of platforms avaible 
//cl_int p = ; 


if(clGetPlatformIDs(1, &platform_id, &num_of_platforms) != CL_SUCCESS) 
{ 
    printf("Unable to get platform id\n"); 
    return 1; 
} 


// try to get a supported GPU device 
if (clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_CPU, 1, &device_id, &num_of_devices) != CL_SUCCESS) 
{ 
// printf("shbham"); 
printf("Unable to get device_id\n"); 
return 1; 
} 

// context properties list - must be terminated with 0 
properties[0]= CL_CONTEXT_PLATFORM; 
properties[1]= (cl_context_properties) platform_id; 
properties[2]= 0; 

// create a context with the GPU device 
context = clCreateContext(properties,1,&device_id,NULL,NULL,&err); 

// create command queue using the context and device 
command_queue = clCreateCommandQueue(context, device_id, 0, &err); 

// create a program from the kernel source code 
program = clCreateProgramWithSource(context,1,(const char **) &ProgramSource, NULL, &err); 

// compile the program 
if (clBuildProgram(program, 0, NULL, NULL, NULL, NULL) != CL_SUCCESS) 
{ 
printf("Error building program\n"); 
return 1; 
} 

// specify which kernel from the program to execute 
kernel = clCreateKernel(program, "add", &err); 

// create buffers for the input and ouput 

inputA = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(float) * DATA_SIZE, NULL, NULL); 
inputB = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * 16, NULL, NULL); 
output = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * DATA_SIZE, NULL, NULL); 

// load data into the input buffer 
clEnqueueWriteBuffer(command_queue, inputA, CL_TRUE, 0, sizeof(float) * DATA_SIZE, inputDataA, 0, NULL, NULL); 
clEnqueueWriteBuffer(command_queue, inputB, CL_TRUE, 0, sizeof(float) * 16, inputDataB, 0, NULL, NULL); 
clEnqueueWriteBuffer(command_queue, output, CL_TRUE, 0, sizeof(float) * DATA_SIZE, 0, 0, NULL, NULL); 

// set the argument list for the kernel command 
clSetKernelArg(kernel, 0, sizeof(cl_mem), &inputA); 
clSetKernelArg(kernel, 1, sizeof(cl_mem), &output); 
clSetKernelArg(kernel, 2, sizeof(cl_mem), &inputB); 

// enqueue the kernel command for execution 
clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global, NULL, 0, NULL, NULL); 
clFinish(command_queue); 

// copy the results from out of the output buffer 
clEnqueueReadBuffer(command_queue, output, CL_TRUE, 0, sizeof(float) *DATA_SIZE, results, 0, NULL, NULL); 
clEnqueueReadBuffer(command_queue, inputB, CL_TRUE, 0, sizeof(float) *16, shubh, 0, NULL, NULL); 

// print the results 
printf("output: "); 

for(i=0;i<DATA_SIZE; i++) 
{ 
printf("%f ",results[i]); 
outfile << results[i] << endl; 
} 
for(i=0;i<16;i++) 
{ 
outfile << shubh[i] <<" "; 
} 
// cleanup - release OpenCL resources 
clReleaseMemObject(inputA); 
clReleaseMemObject(inputB); 
clReleaseMemObject(output); 
clReleaseProgram(program); 
clReleaseKernel(kernel); 
clReleaseCommandQueue(command_queue); 
clReleaseContext(context); 
return 0; 
} 

Я бегу этот код с 8 рабочими элементами. Вход [1,2,3,4,5,6,7,8], а ожидаемый результат должен быть [1,3,5,7,9,11,13,15], но каждый раз, когда я запускаю свой код Я получаю разные результаты, такие как [1, 3, 5, 4, 5, 6, 7, 15]. Кажется, что thid не обновляют свой индекс в temp в состоянии «if».

Если проблема связана не с использованием функции atomic_add при добавлении в условие «если», то какой должен быть синтаксис, чтобы изменить его на атомный, я попытался сам, но получаю ошибку при компиляции.

Или, если есть какая-либо другая проблема, пожалуйста, помогите мне в ее исправлении.

PS. Я запускаю свой код с помощью DEVICE_TYPE_CPU, и он показывает ошибку при использовании DEVICE_TYPE_GPU. Надеюсь, это не причина проблемы.

Пожалуйста, помогите

Answer 1

EDIT: Если указать NULL для локального размера рабочей группы вы позволить решить реализации. Я предполагаю, что реализация AMD выбирает 1 как локальный размер рабочей группы, и у вас есть 8 рабочих групп размером 1. Таким образом, у вас есть 8 потоков, выполняющих расы данных в вашем массиве temp. Это глобальная память и поэтому ее разделяет между рабочими группами. барьеры здесь не помогают, потому что вы не можете синхронизировать между рабочими группами в OpenCL, но вам нужна такая синхронизация. Это также объяснит, почему ваш код работает корректно, если вы указали размер локальной группы ворг 8. Затем у вас есть 1 рабочая группа, и барьер может синхронизировать ваши потоки.

---

Хорошо давайте посмотрим на ядро:

__kernel void add(__global float *input, __global float *output, __global float *temp) 
{ 
    int thid = get_global_id(0); 
    int pout = 0; 
    int pin = 1; 
    temp[pin*8 + thid] = input[thid]; 
    temp[pout*8 + thid] = input[thid]; 
    pout = 1-pout; 
    pin = 1-pout; 
    int offset = 1; 

    if(thid >= offset) { 
     temp[pout*8 + thid] =temp[pout*8 + thid] + temp[pin*8 + thid - offset]; 
    } else { 
     temp[pout*8 + thid] = temp[pin*8 + thid]; 
    } 

    barrier(CLK_GLOBAL_MEM_FENCE); 
    output[thid] = temp[pout*8 + thid]; 
} 

Сначала я бы удалить дополнительную память, так как она просто копирует данные в два раза убийца производительность и может быть ваша проблема, а также. (Я не знаю аппаратное обеспечение, на котором запущено ядро, и если есть что-то вроде неявной синхронизации основы, как на графических процессорах Nvidia). Проблема здесь (с точки зрения параллельного программирования) - это простое условие гонки. Ваши потоки не записывались в temp перед другими потоками данных. Два решения: a) избавиться от temp, b) поставить барьер перед вашим оператором if. Однако в OpenCL барьер может синхронизировать потоки только в одной рабочей группе, поэтому это ядро ​​может вызвать ту же проблему, если вы используете более одной рабочей группы.

Поскольку вы только чтение из входных данных и записи на выход вам не нужна температура:

__kernel void add(__global float *input, __global float *output, __global float *temp) 
{ 
    int thid = get_global_id(0); 
    int offset = 1; 

    if(thid >= offset) { 
     output[thid] = input[thid] + input[thid - offset]; 
    } else { 
     output[thid] = input[thid]; 
    } 
} 

Это должно сделать это.