openCL CL_OUT_OF_RESOURCES Ошибка

Question

Я пытаюсь преобразовать код, написанный в Cuda, в openCL и столкнуться с некоторыми проблемами. Моя конечная цель - реализовать код на плате Odroid XU3 с графическим процессором Mali T628.

Для того, чтобы упростить переход и сэкономить время, пытаясь отладки OpenCL ядер я сделал следующие шаги:

1. Реализовать код в Cuda и протестировать его на Nvidia GeForce 760
2. Реализовать код в openCL и протестируйте его на Nvidia GeForce 760
3. проверьте код openCL на плате Odroid XU3 с графическим процессором Mali T628.

Я знаю, что разные архитектуры могут иметь разные оптимизации, но на данный момент это не моя главная проблема. Мне удалось запустить код openCL на моем графическом процессоре Nvidia без каких-либо явных проблем, но все равно получаю странные ошибки при попытке запустить код на плате Odroid. Я знаю, что разные архитектуры имеют различную обработку исключений и т. Д., Но я не уверен, как их решить.

Поскольку код OpenCL работает на моем Nvidia я полагаю, что мне удалось сделать правильный переход между резьбовыми/блоками -> WorkItems/Workgroups и т.д. я уже исправлен ряд вопросов, которые относятся к вопросу cl_device_max_work_group_size так, что не может быть куазой.

При запуске кода я получаю ошибку «CL_OUT_OF_RESOURCES». Я сузил причину ошибки до 2 строк в коде, но не уверен, чтобы исправить эти проблемы.

ошибка вызвана следующими линиями:

1. lowestDist [pixelNum] = partialDiffSumTemp; обе переменные являются частными переменными ядра, и поэтому я не вижу никакой потенциальной проблемы.
2. d_disparityLeft [globalMemIdx + TILE_BOUNDARY_WIDTH - WINDOW_RADIUS + 0] = bestDisparity [0]; Здесь я предполагаю, что причина «OUT_OF_BOUND», но не уверен, как отлаживать ее, поскольку исходный код не имеет никаких проблем.

Мой Kernel код является:

#define ALIGN_IMAGE_WIDTH   64 
#define NUM_PIXEL_PER_THREAD  4 

#define MIN_DISPARITY    0 
#define MAX_DISPARITY    55 

#define WINDOW_SIZE    19 
#define WINDOW_RADIUS    (WINDOW_SIZE/2) 

#define TILE_SHARED_MEM_WIDTH  96      
#define TILE_SHARED_MEM_HEIGHT  32 
#define TILE_BOUNDARY_WIDTH  64 
#define TILE_BOUNDARY_HEIGHT  (2 * WINDOW_RADIUS) 

#define BLOCK_WIDTH    (TILE_SHARED_MEM_WIDTH - TILE_BOUNDARY_WIDTH) 
#define BLOCK_HEIGHT    (TILE_SHARED_MEM_HEIGHT - TILE_BOUNDARY_HEIGHT) 

#define THREAD_NUM_WIDTH   8 
#define THREADS_NUM_HEIGHT   TILE_SHARED_MEM_HEIGHT 

//TODO fix input arguments 
__kernel void hello_kernel(__global unsigned char* d_leftImage, 
          __global unsigned char* d_rightImage, 
          __global float* d_disparityLeft) { 

    int blockX  = get_group_id(0); 
    int blockY  = get_group_id(1); 
    int threadX  = get_local_id(0); 
    int threadY  = get_local_id(1); 

    __local unsigned char leftImage  [TILE_SHARED_MEM_WIDTH * TILE_SHARED_MEM_HEIGHT]; 
    __local unsigned char rightImage  [TILE_SHARED_MEM_WIDTH * TILE_SHARED_MEM_HEIGHT]; 
    __local unsigned int partialDiffSum [BLOCK_WIDTH   * TILE_SHARED_MEM_HEIGHT]; 

    int alignedImageWidth = 640; 
    int partialDiffSumTemp; 
    float bestDisparity[4] = {0,0,0,0}; 
    int lowestDist[4]; 
     lowestDist[0] = 214748364; 
     lowestDist[1] = 214748364; 
     lowestDist[2] = 214748364; 
     lowestDist[3] = 214748364; 

    // Read image blocks into shared memory. read is done at 32bit integers on a uchar array. each thread reads 3 integers(12byte) 96/12=8threads 
    int sharedMemIdx = threadY * TILE_SHARED_MEM_WIDTH + 4 * threadX; 
    int globalMemIdx = (blockY * BLOCK_HEIGHT + threadY) * alignedImageWidth + blockX * BLOCK_WIDTH + 4 * threadX; 

    for (int i = 0; i < 4; i++) { 
     leftImage [sharedMemIdx      + i ] = d_leftImage [globalMemIdx      + i]; 
     leftImage [sharedMemIdx + 4 * THREAD_NUM_WIDTH + i ] = d_leftImage [globalMemIdx + 4 * THREAD_NUM_WIDTH + i]; 
     leftImage [sharedMemIdx + 8 * THREAD_NUM_WIDTH + i ] = d_leftImage [globalMemIdx + 8 * THREAD_NUM_WIDTH + i]; 
     rightImage[sharedMemIdx      + i ] = d_rightImage[globalMemIdx      + i]; 
     rightImage[sharedMemIdx + 4 * THREAD_NUM_WIDTH + i ] = d_rightImage[globalMemIdx + 4 * THREAD_NUM_WIDTH + i]; 
     rightImage[sharedMemIdx + 8 * THREAD_NUM_WIDTH + i ] = d_rightImage[globalMemIdx + 8 * THREAD_NUM_WIDTH + i]; 
    } 

    barrier(CLK_LOCAL_MEM_FENCE); 

    int imageIdx = sharedMemIdx + TILE_BOUNDARY_WIDTH - WINDOW_RADIUS; 
    int partialSumIdx = threadY * BLOCK_WIDTH + 4 * threadX; 

    for(int dispLevel = MIN_DISPARITY; dispLevel <= MAX_DISPARITY; dispLevel++) { 

     // horizontal partial sum 
     partialDiffSumTemp = 0; 
     #pragma unroll 
     for(int i = imageIdx - WINDOW_RADIUS; i <= imageIdx + WINDOW_RADIUS; i++) { 
        //partialDiffSumTemp += calcDiff(leftImage [i], rightImage[i - dispLevel]); 
         partialDiffSumTemp += abs(leftImage[i] - rightImage[i - dispLevel]); 
     } 
     partialDiffSum[partialSumIdx] = partialDiffSumTemp; 

     barrier(CLK_LOCAL_MEM_FENCE); 

     for (int pixelNum = 1, i = imageIdx - WINDOW_RADIUS; pixelNum < NUM_PIXEL_PER_THREAD; pixelNum++, i++) { 
      partialDiffSum[partialSumIdx + pixelNum] = partialDiffSum[partialSumIdx + pixelNum - 1] + 
                 abs(leftImage[i + WINDOW_SIZE] - rightImage[i - dispLevel + WINDOW_SIZE]) - 
                 abs(leftImage[i]    - rightImage[i - dispLevel]); 
     } 

     barrier(CLK_LOCAL_MEM_FENCE); 

     // vertical sum 
     if(threadY >= WINDOW_RADIUS && threadY < TILE_SHARED_MEM_HEIGHT - WINDOW_RADIUS) { 

      for (int pixelNum = 0; pixelNum < NUM_PIXEL_PER_THREAD; pixelNum++) { 
       int rowIdx = partialSumIdx - WINDOW_RADIUS * BLOCK_WIDTH; 
       partialDiffSumTemp = 0; 

        for(int i = -WINDOW_RADIUS; i <= WINDOW_RADIUS; i++,rowIdx += BLOCK_WIDTH) { 
          partialDiffSumTemp += partialDiffSum[rowIdx + pixelNum]; 
        } 

        if (partialDiffSumTemp < lowestDist[pixelNum]) { 
         lowestDist[pixelNum] = partialDiffSumTemp; 
         bestDisparity[pixelNum] = dispLevel - 1; 
        } 


      } 
     } 

    } 

    if (threadY >= WINDOW_RADIUS && threadY < TILE_SHARED_MEM_HEIGHT - WINDOW_RADIUS && blockY < 32) { 

     d_disparityLeft[globalMemIdx + TILE_BOUNDARY_WIDTH - WINDOW_RADIUS + 0] = bestDisparity[0]; 
     d_disparityLeft[globalMemIdx + TILE_BOUNDARY_WIDTH - WINDOW_RADIUS + 1] = bestDisparity[1]; 
     d_disparityLeft[globalMemIdx + TILE_BOUNDARY_WIDTH - WINDOW_RADIUS + 2] = bestDisparity[2]; 
     d_disparityLeft[globalMemIdx + TILE_BOUNDARY_WIDTH - WINDOW_RADIUS + 3] = bestDisparity[3]; 
    } 

} 

Спасибо за помощь

Юваль

Answer 1

Из моего опыта NVidia графических процессоров не всегда врезаться на из связанного доступа и много раз ядра все еще возвращает ожидаемые результаты.

Используйте printf для проверки индексов. Если установлен драйвер Nvidia OpenCL 1.2, то printf должен быть доступен как основная функция. Насколько я проверял, Mali-T628 использует OpenCL 1.1, затем проверьте, доступно ли в качестве расширения поставщика printf. Также вы можете запустить свое ядро ​​на процессоре AMD/Intel, где доступно printf (OpenCL 1.2/2.0).

Альтернативный способ проверки индексов может проходить через массив __global int* debug, где вы будете хранить индексы, а затем проверять их на хосте. Не забудьте выделить его достаточно большим, чтобы записывать из привязанного индекса.