Блог

Я хочу распараллелить распределение температур, используя технологию OpenCL. Я столкнулся с проблемой с моим графическим процессором — идентификатор рабочего элемента для всех остальных функций ядра одинаковый. Вместо результата, например, от 0 до 1024, я получил этот результат. Что я сделал неправильно?

введите описание изображения здесь

Source.cpp

 include lt;iostreamgt; #include lt;stringgt; #include lt;fstreamgt;  #include lt;omp.hgt; #include lt;CL/cl.hppgt;  float*** distributeOpenCL(float*** cuboid, int k, int m, int n) {  // OpenCL init  int size = k * m * n;  float*** hResult = initCuboid(k, m, n);  cl_platform_id platform;  cl_device_id device;  cl_int error = 0;  std::ifstream file("program.cl");  std::string fileText = std::string(std::istreambuf_iteratorlt;chargt;(file), std::istreambuf_iteratorlt;chargt;());  const char* srcText = fileText.data();  size_t srcLength = fileText.size();  cl_context context;  cl_program program;  cl_kernel kernel;  cl_command_queue queue;  cl_mem dCuboid, dRes;  size_t localSize[2] = { k,m };  size_t globalSize[2] = { ceil(size / (float)localSize[0]) * localSize[0], ceil(size / (float)localSize[1]) * localSize[1] };  // Get GPU  error |= clGetPlatformIDs(1, amp;platform, NULL);  error |= clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, amp;device, NULL);  // Compile and build  context = clCreateContext(NULL, 1, amp;device, NULL, NULL, amp;error);  program = clCreateProgramWithSource(context, 1, amp;srcText, amp;srcLength, amp;error);  error |= clBuildProgram(program, 1, amp;device, NULL, NULL, NULL);  // What funtion from file we have to run  kernel = clCreateKernel(program, "distributeKernel", amp;error);  // Add to Queue  queue = clCreateCommandQueueWithProperties(context, device, NULL, amp;error);  // Create buffer  dCuboid = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(float) * size, NULL, NULL);  dRes = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(float) * size, NULL, NULL);  // Write data to buffer  error |= clEnqueueWriteBuffer(queue, dCuboid, CL_TRUE, 0, sizeof(float) * size, cuboid, 0, NULL, NULL);  // Kernel args  error |= clSetKernelArg(kernel, 0, sizeof(cl_mem), amp;dCuboid);  error |= clSetKernelArg(kernel, 1, sizeof(int), amp;k);  error |= clSetKernelArg(kernel, 2, sizeof(int), amp;m);  error |= clSetKernelArg(kernel, 3, sizeof(int), amp;n);  error |= clSetKernelArg(kernel, 4, sizeof(cl_mem), amp;dRes);  // Start task  error |= clEnqueueNDRangeKernel(queue, kernel, 2, NULL, globalSize, localSize, 0, NULL, NULL);  // Wait execution  clFinish(queue);  // Read Result  error |= clEnqueueReadBuffer(queue, dRes, CL_TRUE, 0, sizeof(float) * size, hResult, 0, NULL, NULL);  //printCuboid(resP, k, m, n, resPFile);  // Deallocation  clReleaseKernel(kernel);  clReleaseMemObject(dCuboid);  clReleaseMemObject(dRes);  clReleaseCommandQueue(queue);  clReleaseProgram(program);  clReleaseContext(context);  return hResult; }  int main(int argc, char* argv[]) {  std::ofstream filledFile("filled.txt");  std::ofstream resLFile("resL.txt");  std::ofstream resPFile("resP.txt");  double durationL, durationP, time1, time2;  int k = 5, m = 5, n = 5, temp1 = 10, temp2 = 15;  float*** cuboid, *** resL, *** resP;  if (argc gt; 1) {  k = atoi(argv[1]), m = atoi(argv[2]), n = atoi(argv[3]),  temp1 = atoi(argv[4]), temp2 = atoi(argv[5]);  }  // Linear  cuboid = initCuboid(k, m, n);  fillCuboid(cuboid, k, m, n, temp1, temp2);  printCuboidToFile(cuboid, k, m, n, filledFile);  time1 = omp_get_wtime();  resL = distribute(cuboid, k, m, n);  time2 = omp_get_wtime();  durationL = time2 - time1;  printCuboidToFile(resL, k, m, n, resLFile);  // Parallel  time1 = omp_get_wtime();  resP = distributeOpenCL(cuboid, k, m, n);  time2 = omp_get_wtime();  durationP = time2 - time1;  //printCuboidToFile(resP, k, m, n, resPFile);  std::cout lt;lt; "Linear time: " lt;lt; durationL lt;lt; std::endl;  std::cout lt;lt; "Parallel time: " lt;lt; durationP lt;lt; std::endl;  std::cout lt;lt; "Parallel faster than linear on: " lt;lt; durationL - durationP lt;lt; std::endl;  // Delete 3d arrays, closing files  deleteCuboid(cuboid, k, m, n);  deleteCuboid(resL, k, m, n);  deleteCuboid(resP, k, m, n);  filledFile.close();  resLFile.close();   resPFile.close();  return 0; }  

program.cl

 __kernel void distributeKernel(__global float*** cuboid, int k, int m, int n, __global float*** result) {  int gz = get_global_id(0);  int gy = get_global_id(1);  printf("gy - %d n", amp;gy);  printf("gz - %d n", amp;gz);  bool isDissipated = false;  int size = k * m * n;  // Ends if temperatures in cube becomes balanced  while (!isDissipated) {  int dissipatedCount = 0;  for (int x = 0; x lt; n; x  ) {  // Calc average temperature  float sum = 0;  int count = 0;  float average;  for (int zSum = gz - 1; zSum lt;= gz   1; zSum  ) {  for (int ySum = gy - 1; ySum lt;= gy   1; ySum  ) {  for (int xSum = x - 1; xSum lt;= x   1; xSum  ) {  if (zSum gt;= 0 amp;amp; ySum gt;= 0 amp;amp; xSum gt;= 0  amp;amp; zSum lt; k amp;amp; ySum lt; m amp;amp; xSum lt; n) {  count  ;  sum  = result[gz][gy][xSum];  }  }  }  }  average = round(sum / count * 100) / 100;  if (average == result[gz][gy][x]) {  dissipatedCount  ;  }  else {  result[gz][gy][x] = average;  }  }  if (dissipatedCount == size) {  isDissipated = true;  }  } }  

Чтобы устранить проблему с якобы неправильным get_global_id() решением, начните с простой, минимальной программы добавления векторов в стиле «Привет, мир», а затем шаг за шагом переходите к приложению распределения температуры.

С вашим кодом я вижу несколько проблем:

1. В OpenCL могут быть только 1D указатели (с одним * ).

    __kernel void distributeKernel(__global float* cuboid, __global float* result)  

   Введите линейный индекс для доступа к более чем 1 измерению: например, для 2D int n = x y*get_global_size(0);

2. Из того, что я вижу, k , m , n являются размерами решетки. Полностью исключите их из ядра. Дозвонись size get_global_size(...) .
3. Ядро выглядит довольно сложным с большим количеством циклов и ветвлений. Это может свести на нет любое преимущество в производительности, которое вы надеетесь получить от распараллеливания GPU. Избавьтесь от петель и разветвлений, насколько это возможно. Кроме того, не должно быть никакого цикла над одним из размеров решетки, так как положение решетки-это то, что вы распараллеливаете.

Я бы также посоветовал использовать только 1D-распараллеливание в OpenCL и выполнять линейную индексацию самостоятельно. Это дает вам большую гибкость в отношении размера рабочей группы.