1 // See https://github.com/pocl/pocl/issues/757
2 // Caused an out of bounds read.
3
4 #include <algorithm>
5 #include <cassert>
6 #include <iostream>
7 #include <vector>
8
9 #define CL_HPP_MINIMUM_OPENCL_VERSION 120
10 #define CL_HPP_TARGET_OPENCL_VERSION 120
11 #include <CL/cl2.hpp>
12
13 const char *SOURCE = R"RAW(
14 #define lid(N) ((int) get_local_id(N))
15 #define gid(N) ((int) get_group_id(N))
16
17 __kernel void __attribute__ ((reqd_work_group_size(128, 1, 1))) grudge_assign_0(
18 int const grdg_n,
19 __global double *__restrict__ expr_8,
20 int const expr_8_offset,
21 __global double const *__restrict__ grdg_sub_discr_dx0_dr0,
22 int const grdg_sub_discr_dx0_dr0_offset)
23 {
24 if (-1 + -128 * gid(0) + -1 * lid(0) + grdg_n >= 0)
25 expr_8[expr_8_offset + 128 * gid(0) + lid(0)] = grdg_sub_discr_dx0_dr0[grdg_sub_discr_dx0_dr0_offset + 128 * gid(0) + lid(0)];
26 }
27 )RAW";
28
main(int argc,char * argv[])29 int main(int argc, char *argv[]) {
30 int n = 8;
31
32 cl::Device device = cl::Device::getDefault();
33 cl::CommandQueue queue = cl::CommandQueue::getDefault();
34 cl::Program program(SOURCE, true);
35
36 // Create buffers on the device.
37 cl::Buffer buffer_A(CL_MEM_READ_WRITE, sizeof(double) * n);
38 cl::Buffer buffer_B(CL_MEM_READ_WRITE, sizeof(double) * n);
39
40 std::vector<double> A(n);
41 std::vector<double> B(n);
42 std::fill(B.begin(), B.end(), 1);
43
44 // Write arrays to the device.
45 queue.enqueueWriteBuffer(buffer_A, CL_TRUE, 0, sizeof(double) * n, A.data());
46 queue.enqueueWriteBuffer(buffer_B, CL_TRUE, 0, sizeof(double) * n, B.data());
47
48 // Run the kernel.
49 cl::Kernel knl(program, "grudge_assign_0");
50 int sz = 1;
51 knl.setArg(0, sz);
52 knl.setArg(1, buffer_A);
53 knl.setArg(2, 0);
54 knl.setArg(3, buffer_B);
55 knl.setArg(4, 0);
56 queue.enqueueNDRangeKernel(
57 knl,
58 cl::NullRange,
59 cl::NDRange(((sz+127)/128)*128),
60 cl::NDRange(128));
61 queue.finish();
62
63 // Read result A from the device.
64 queue.enqueueReadBuffer(buffer_A, CL_TRUE, 0, sizeof(double) * n, A.data());
65
66 for (int i = 0; i < n; ++i) {
67 if (i < sz) {
68 assert(A[i] == 1);
69 } else {
70 assert(A[i] == 0);
71 }
72 }
73
74 return EXIT_SUCCESS;
75 }
76