opencl:clEnqueueNDRangeKernel执行报错CL_OUT_OF_RESOURCES的一种情况

我的电脑上之前的显卡比较老并不支持opencl，所以我之前开发时opencl代码其实都是在CPU上跑的，现在所有的代码都调试通过了，决定装块新显卡用于程序的性能测试。

今天显卡到了，装上之后运行程序，

clEnqueueNDRangeKernel

在执行下面的kernel时报错：

CL_OUT_OF_RESOURCES

。

__kernel void prefix_sum_col_and_transpose( __constant SRC_TYPE *src, __global DST_TYPE * dst, uint width,uint height, uint src_width_step, uint dst_width_step){
.........// 代码实现部分略过
}

百撕不得其姐啊。。。。这代码在CPU上跑很正常，逻辑没问题呀。

最后发现只是kernel 指针参数的地址修饰符使用不当造成的。

上面这段代码，是用于图像积分图计算的，对给定的原图(src)数据计算积分图，输出到目标指针(dst)指向的全局内存中。因为src数据不允许被修改所以我想当然的把src指定为

__constant

。

而这里用

__constant

修饰是不对的。

__constant

和

__global

都是全局内存，

__constant

修饰的地址指向的是常量，不能被修改，但它们之间的区别却并不仅于此。

一个opencl设备的常量空间是有限制的，通过clGetDeviceInfo获取

CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE

可以知道一个opencl设备的最大常量缓冲区的尺寸，在我的显卡上，这个值是65536，简单通过命令行运行AMD APP SDK的

clinfo

就可以得到这个值，如下图：



因为图像的尺寸很容易就超过64kb，所以

clEnqueueNDRangeKernel

在执行kernel时无法将它放到opencl设备的constant buffer中，所以就会报错

CL_OUT_OF_RESOURCES

。

所以应该将

src

的地址修饰符从

__constant

改为

__global

,如果要禁止修改

src

指针的数据，前面用c语言标准的

const

关键字修饰这个指针就可以了，所以这个kernel函数正确的定义应该是这样：

__kernel void prefix_sum_col_and_transpose( const __global SRC_TYPE *src, __global DST_TYPE * dst, uint width,uint height, uint src_width_step, uint dst_width_step){
.........// 代码实现部分略过
}