1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
|
// See https://github.com/pocl/pocl/issues/553
#include "pocl_opencl.h"
#define CL_HPP_ENABLE_EXCEPTIONS
#define CL_HPP_MINIMUM_OPENCL_VERSION 120
#define CL_HPP_TARGET_OPENCL_VERSION 120
#define CL_HPP_CL_1_2_DEFAULT_BUILD
#include <CL/opencl.hpp>
#include <iostream>
using namespace std;
const char *SOURCE = R"RAW(
// Expected output:
// outer=A inner=B
// + outer=A inner=B
// for each value of A and B.
// However I see three copies of the second line (starting with +).
// Commenting out any one line marked with YYYY bring it down to two copies,
// and commenting out any one line marked with XXXX gives the expected output.
__kernel void pocltest(int xarg1, int xarg2) {
int outerend = 1;
int innerend = 1;
outerend = 2; // YYYY
innerend = 2; // YYYY
int outer = 0;
int inner = 0;
int arg1 = 1;
int arg2 = 1;
arg1 = xarg1; // XXXX
arg2 = xarg2; // XXXX
for (outer = 0; outer < outerend; outer++) // XXXX
{
for (inner = 0; inner < innerend; inner++) // XXXX
{
//barrier(CLK_LOCAL_MEM_FENCE);
printf("outer=%d inner=%d lid=%d\n", outer, inner, get_local_id(0));
if (arg2 > arg1) // XXXX
{
barrier(CLK_LOCAL_MEM_FENCE); // XXXX
}
if (arg1 > 0) // XXXX
{
barrier(CLK_LOCAL_MEM_FENCE); // XXXX
}
printf("+ outer=%d inner=%d lid=%d\n", outer, inner, get_local_id(0));
//barrier(CLK_LOCAL_MEM_FENCE); /* This barrier also fixes it. */
}
}
}
)RAW";
int main(int argc, char *argv[])
{
cl::Platform platform = cl::Platform::getDefault();
cl::Device device = cl::Device::getDefault();
try {
cl::CommandQueue queue = cl::CommandQueue::getDefault();
cl::Program program(SOURCE, true);
#if (__GNUC__ > 5)
#pragma GCC diagnostic push
#pragma GCC diagnostic ignored "-Wignored-attributes"
#endif
auto kernel = cl::KernelFunctor<cl_int, cl_int>(program, "pocltest");
cl::Buffer buffer;
kernel(cl::EnqueueArgs(queue, cl::NDRange(2), cl::NDRange(2)), 1, 2);
#if (__GNUC__ > 5)
#pragma GCC diagnostic pop
#endif
queue.finish();
} catch (cl::Error &err) {
std::cout << "FAIL with OpenCL error = " << err.err() << std::endl;
return EXIT_FAILURE;
}
platform.unloadCompiler();
std::cout << "OK" << std::endl;
return EXIT_SUCCESS;
}
|
