【发布时间】：2015-06-07 17:46:09
【问题描述】：
我正在尝试将用 Cuda 编写的代码转换为 openCL，但遇到了一些麻烦。我的最终目标是在带有 Mali T628 GPU 的 Odroid XU3 板上实现代码。
为了简化转换并节省尝试调试 openCL 内核的时间，我执行了以下步骤：
- 在 Cuda 中实现代码并在 Nvidia GeForce 760 上进行测试
- 在 openCL 中实现代码并在 Nvidia GeForce 760 上进行测试
- 在带有 Mali T628 GPU 的 Odroid XU3 板上测试 openCL 代码。
我知道不同的架构可能有不同的优化，但这不是我现在主要关心的问题。我设法在我的 Nvidia GPU 上运行 openCL 代码，没有明显问题，但是在尝试在 Odroid 板上运行代码时不断出现奇怪的错误。我知道不同的架构对异常等有不同的处理，但我不确定如何解决这些问题。
由于 openCL 代码在我的 Nvidia 上运行，我假设我设法在线程/块 -> workItems/workGroups 等之间进行了正确的转换。 我已经修复了几个与 cl_device_max_work_group_size 问题相关的问题，因此这不是原因。
运行代码时，我收到“CL_OUT_OF_RESOURCES”错误。我已将错误原因缩小到代码中的 2 行，但不确定是否能解决这些问题。
错误是由以下几行引起的：
- lowestDist[pixelNum] = partialDiffSumTemp;这两个变量都是内核的私有变量，因此我没有看到任何潜在问题。
- 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];
}
}
感谢大家的帮助
尤瓦尔
【问题讨论】：
-
GPU 代码很难调试，尤其是在涉及异常硬件时。很难想象这个“问题”的“答案”会是什么样子：人们只能尝试猜测可能有什么问题。但是，越界访问可能会导致
CL_OUT_OF_RESOURCES错误是正确的。所以printf调试的替代方案：你也可以用cuda-memcheck YourProgram.exe运行你的程序：它会打印是否有无效的内存访问（甚至可以获取行号信息，但我不确定) -
我知道这是旧的，但我遇到了类似的问题。我正在启动多个内核，但我不断收到“资源不足”错误。在我减少了内核中私有变量的使用后，大多数内核现在运行没有错误，所以它可能会用完寄存器......？这是一个非常奇怪的问题，我还没有修复最后一个内核。需要注意的另一件事是，Mali GPU 将其共享内存类型报告为“全局”，因此可能不会从中获得任何性能提升，并且在访问本地内存时会出现这些错误。因此，一种可能的解决方案是消除共享内存的使用。
-
用户在 ARM 社区论坛上发布了这个问题，看来问题出在本地工作大小上。解决这个问题也解决了我的问题。这很奇怪，因为我使用的是 8*32 本地工作大小，因此我预计会出现一个错误，抱怨工作大小无效（就像以前一样）。