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){
	.........// 代码实现部分略过
}