openCL CL_OUT_OF_RESOURCES錯誤

Question

我試圖將Cuda中寫入的代碼轉換爲openCL並遇到一些麻煩。我的最終目標是使用Mali T628 GPU在Odroid XU3開發板上實現代碼。

爲了簡化過渡，並節省時間試圖調試OpenCL內核我已經做了以下步驟：

1. 實現Cuda的代碼和測試它了NVIDIA GeForce 760
2. 落實在openCL中測試代碼並在Nvidia GeForce 760上測試它使用Mali T628 GPU測試Odroid XU3開發板上的openCL代碼。

我知道不同的架構可能有不同的優化，但這不是我現在主要關心的問題。我試圖在我的Nvidia GPU上運行openCL代碼，但沒有明顯的問題，但在嘗試在Odroid主板上運行代碼時仍然出現奇怪的錯誤。我知道不同的架構有不同的異常處理等，但我不知道如何解決這些問題。

由於OpenCL的代碼工作在我的Nvidia我認爲我能夠做到線程/塊之間的正確轉換 - >工作項/工作組等 我已經解決了，涉及到cl_device_max_work_group_size問題的幾個問題，使不能是cuase。

當運行代碼時，我得到一個「CL_OUT_OF_RESOURCES」錯誤。我已經將代碼中的錯誤原因縮小爲2行，但不確定解決這些問題。

錯誤是由以下幾行引起的：

1. lowestDist [pixelNum] = partialDiffSumTemp;這兩個變量都是內核的私有變量，因此我沒有看到任何潛在的問題。
2. d_disparityLeft [globalMemIdx + TILE_BOUNDARY_WIDTH - WINDOW_RADIUS + 0] = bestDisparity [0]; 在這裏我猜原因是「OUT_OF_BOUND」，但不知道如何調試它，因爲原來的代碼沒有任何問題。

我的內核代碼是：

#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 GPU並不總是會崩潰越界訪問並多次內核仍然返回預期的結果。

使用printf來檢查索引。如果您安裝了Nvidia OpenCL 1.2驅動程序，則應該將printf作爲核心功能。據我查看Mali-T628使用OpenCL 1.1，然後檢查printf是否可用作供應商擴展。你也可以在AMD/Intel CPU上運行內核，其中printf可用（OpenCL 1.2/2.0）。

檢查索引的替代方法可以傳遞__global int* debug數組，您可以在其中存儲索引，然後在主機上檢查它們。確保分配足夠大的空間，以便記錄出界索引。