时间戳

首页 > 解决方案 > 为什么这个 printf 函数修复了我的 OpenCL 内核？

c++ - 为什么这个 printf 函数修复了我的 OpenCL 内核？

问题描述

我正在编写一个使用 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 到控制台？

将 OpenCL 2 与 RTX 2060 一起使用 - 如果需要任何其他信息，请告诉我，我会找到的！

标签： c++parallel-processingopenclopencl-c

解决方案

当你的代码：

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; 
  }
}

是馈送到这个网站它发现以下问题：

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, 1) in work group (4, 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, 1) in work group (4, 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, 0) in work group (0, 0), 6:16:
c_hist[id] = hist[id];

Write by work item (20, 1) in work group (4, 4), 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, 0), 6:18:
c_hist[id] = hist[id];

Write by work item (5, 1) in work group (4, 4), 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]:

推荐阅读