问题描述
我正在编写一个使用 Hillis-Steele 扫描模式计算累积直方图的内核。它似乎工作不正常 - 输出的数字太高且顺序错误。
在调试过程中,我添加了一个简单的 printf() 函数,它只是打印出每个运行的工作项的全局大小:
kernel void cumulative(global const int *hist,global int *c_hist) {
int id = get_global_id(0);
int size = get_global_size(0);
printf("%d\n",size); //MAKES THE CUMULATIVE HISTOGRAM CORRECT
for (int step = 1; step < size; step *= 2) {
c_hist[id] = hist[id];
if (id >= step) c_hist[id] += hist[id - step];
barrier(CLK_GLOBAL_MEM_FENCE);
global int* tmp = hist; hist = c_hist; c_hist = tmp;
}
}
有人能告诉我这里发生了什么吗?有什么方法可以在不一遍又一遍地向控制台打印 1024 的情况下修复此代码?
在 RTX 2060 上使用 OpenCL 2 - 如果需要任何其他信息,请告诉我,我会找到的!
解决方法
当您的代码:
kernel void cumulative(global const int *hist,global int *c_hist) {
int id = get_global_id(0);
int size = get_global_size(0);
for (int step = 1; step < size; step *= 2) {
c_hist[id] = hist[id];
if (id >= step) c_hist[id] += hist[id - step];
barrier(CLK_GLOBAL_MEM_FENCE);
global int* tmp = hist; hist = c_hist; c_hist = tmp;
}
}
馈送至 this site 它发现以下问题:
10:17: warning: initializing '__global int *' with an expression of type 'const __global int *' discards qualifiers [-Wincompatible-pointer-types-discards-qualifiers]
global int* tmp = hist; hist = c_hist; c_hist = tmp;
^ ~~~~
1 warning generated.
00d8ad302736fae9192e5f1d4027e53e.opt.bc: warning: Assuming the arguments 'hist','c_hist' of 'cumulative' on line 1 of to be non-aliased; please consider adding a restrict qualifier to these arguments
6:18: error: possible null pointer access for work item (1,1) in work group (4,4)
c_hist[id] = hist[id];
6:16: error: possible null pointer access for work item (5,4)
c_hist[id] = hist[id];
error: possible read-write race on c_hist[1008]:
Write by work item (0,0) in work group (0,0),6:16:
c_hist[id] = hist[id];
Read by work item (20,4),possible sources are:
6:18:
c_hist[id] = hist[id];
7:32:
if (id >= step) c_hist[id] += hist[id - step];
7:35:
if (id >= step) c_hist[id] += hist[id - step];
error: possible write-write race on hist[1008]:
Write by work item (0,6:16:
c_hist[id] = hist[id];
Write by work item (20,possible sources are:
6:16:
c_hist[id] = hist[id];
7:32:
if (id >= step) c_hist[id] += hist[id - step];
error: possible write-read race on hist[128]:
Read by work item (0,0) in work group (4,6:18:
c_hist[id] = hist[id];
Write by work item (5,possible sources are:
6:16:
c_hist[id] = hist[id];
7:32:
if (id >= step) c_hist[id] += hist[id - step];
GPUVerify kernel analyser finished with 0 verified,5 errors
注意它报告:error: possible write-read race on hist[128]: