1 //
2 // Copyright (c) 2017 The Khronos Group Inc.
3 //
4 // Licensed under the Apache License, Version 2.0 (the "License");
5 // you may not use this file except in compliance with the License.
6 // You may obtain a copy of the License at
7 //
8 // http://www.apache.org/licenses/LICENSE-2.0
9 //
10 // Unless required by applicable law or agreed to in writing, software
11 // distributed under the License is distributed on an "AS IS" BASIS,
12 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13 // See the License for the specific language governing permissions and
14 // limitations under the License.
15 //
16 #include "testBase.h"
17
18 #include <vector>
19
20 const char *pragma_unroll_kernels[] = {
21 "__kernel void pragma_unroll(__global uint *dst)\n"
22 "{\n"
23 " size_t tid = get_global_id(0);\n"
24 " __attribute__((opencl_unroll_hint))\n"
25 " for(size_t i = 0; i < 100; ++i)\n"
26 " dst[i] = i;\n"
27 "}\n",
28 "__kernel void pragma_unroll(__global uint *dst)\n"
29 "{\n"
30 " size_t tid = get_global_id(0);\n"
31 " __attribute__((opencl_unroll_hint(1)))\n"
32 " for(size_t i = 0; i < 100; ++i)\n"
33 " dst[i] = i;\n"
34 "}\n",
35 "__kernel void pragma_unroll(__global uint *dst)\n"
36 "{\n"
37 " size_t tid = get_global_id(0);\n"
38 " __attribute__((opencl_unroll_hint(10)))\n"
39 " for(size_t i = 0; i < 100; ++i)\n"
40 " dst[i] = i;\n"
41 "}\n",
42 "__kernel void pragma_unroll(__global uint *dst)\n"
43 "{\n"
44 " size_t tid = get_global_id(0);\n"
45 " __attribute__((opencl_unroll_hint(100)))\n"
46 " for(size_t i = 0; i < 100; ++i)\n"
47 " dst[i] = i;\n"
48 "}\n",
49 "__kernel void pragma_unroll(__global uint *dst)\n"
50 "{\n"
51 " size_t tid = get_global_id(0);\n"
52 " size_t n = (tid + 1) * 100;\n"
53 " __attribute__((opencl_unroll_hint))\n"
54 " for(size_t i = 0; i < n; ++i)\n"
55 " dst[i] = i;\n"
56 "}\n",
57 "__kernel void pragma_unroll(__global uint *dst)\n"
58 "{\n"
59 " size_t tid = get_global_id(0);\n"
60 " size_t n = (tid + 1) * 100;\n"
61 " __attribute__((opencl_unroll_hint(1)))\n"
62 " for(size_t i = 0; i < n; ++i)\n"
63 " dst[i] = i;\n"
64 "}\n",
65 "__kernel void pragma_unroll(__global uint *dst)\n"
66 "{\n"
67 " size_t tid = get_global_id(0);\n"
68 " size_t n = (tid + 1) * 100;\n"
69 " __attribute__((opencl_unroll_hint(10)))\n"
70 " for(size_t i = 0; i < n; ++i)\n"
71 " dst[i] = i;\n"
72 "}\n",
73 "__kernel void pragma_unroll(__global uint *dst)\n"
74 "{\n"
75 " size_t tid = get_global_id(0);\n"
76 " size_t n = (tid + 1) * 100;\n"
77 " __attribute__((opencl_unroll_hint(100)))\n"
78 " for(size_t i = 0; i < n; ++i)\n"
79 " dst[i] = i;\n"
80 "}\n",
81 "__kernel void pragma_unroll(__global uint *dst)\n"
82 "{\n"
83 " size_t tid = get_global_id(0);\n"
84 " size_t i = 0;\n"
85 " __attribute__((opencl_unroll_hint))\n"
86 " while(i < 100) {\n"
87 " dst[i] = i;\n"
88 " ++i;\n"
89 " }\n"
90 "}\n",
91 "__kernel void pragma_unroll(__global uint *dst)\n"
92 "{\n"
93 " size_t tid = get_global_id(0);\n"
94 " size_t i = 0;\n"
95 " __attribute__((opencl_unroll_hint(1)))\n"
96 " while(i < 100) {\n"
97 " dst[i] = i;\n"
98 " ++i;\n"
99 " }\n"
100 "}\n",
101 "__kernel void pragma_unroll(__global uint *dst)\n"
102 "{\n"
103 " size_t tid = get_global_id(0);\n"
104 " size_t i = 0;\n"
105 " __attribute__((opencl_unroll_hint(10)))\n"
106 " while(i < 100) {\n"
107 " dst[i] = i;\n"
108 " ++i;\n"
109 " }\n"
110 "}\n",
111 "__kernel void pragma_unroll(__global uint *dst)\n"
112 "{\n"
113 " size_t tid = get_global_id(0);\n"
114 " size_t i = 0;\n"
115 " __attribute__((opencl_unroll_hint(100)))\n"
116 " while(i < 100) {\n"
117 " dst[i] = i;\n"
118 " ++i;\n"
119 " }\n"
120 "}\n",
121 "__kernel void pragma_unroll(__global uint *dst)\n"
122 "{\n"
123 " size_t tid = get_global_id(0);\n"
124 " size_t n = (tid + 1) * 100;\n"
125 " size_t i = 0;\n"
126 " __attribute__((opencl_unroll_hint))\n"
127 " while(i < n) {\n"
128 " dst[i] = i;\n"
129 " ++i;\n"
130 " }\n"
131 "}\n",
132 "__kernel void pragma_unroll(__global uint *dst)\n"
133 "{\n"
134 " size_t tid = get_global_id(0);\n"
135 " size_t n = (tid + 1) * 100;\n"
136 " size_t i = 0;\n"
137 " __attribute__((opencl_unroll_hint(1)))\n"
138 " while(i < n) {\n"
139 " dst[i] = i;\n"
140 " ++i;\n"
141 " }\n"
142 "}\n",
143 "__kernel void pragma_unroll(__global uint *dst)\n"
144 "{\n"
145 " size_t tid = get_global_id(0);\n"
146 " size_t n = (tid + 1) * 100;\n"
147 " size_t i = 0;\n"
148 " __attribute__((opencl_unroll_hint(10)))\n"
149 " while(i < n) {\n"
150 " dst[i] = i;\n"
151 " ++i;\n"
152 " }\n"
153 "}\n",
154 "__kernel void pragma_unroll(__global uint *dst)\n"
155 "{\n"
156 " size_t tid = get_global_id(0);\n"
157 " size_t n = (tid + 1) * 100;\n"
158 " size_t i = 0;\n"
159 " __attribute__((opencl_unroll_hint(100)))\n"
160 " while(i < n) {\n"
161 " dst[i] = i;\n"
162 " ++i;\n"
163 " }\n"
164 "}\n",
165 "__kernel void pragma_unroll(__global uint *dst)\n"
166 "{\n"
167 " size_t tid = get_global_id(0);\n"
168 " size_t i = 0;\n"
169 " __attribute__((opencl_unroll_hint))\n"
170 " do {\n"
171 " dst[i] = i;\n"
172 " ++i;\n"
173 " } while(i < 100);\n"
174 "}\n",
175 "__kernel void pragma_unroll(__global uint *dst)\n"
176 "{\n"
177 " size_t tid = get_global_id(0);\n"
178 " size_t i = 0;\n"
179 " __attribute__((opencl_unroll_hint(1)))\n"
180 " do {\n"
181 " dst[i] = i;\n"
182 " ++i;\n"
183 " } while(i < 100);\n"
184 "}\n",
185 "__kernel void pragma_unroll(__global uint *dst)\n"
186 "{\n"
187 " size_t tid = get_global_id(0);\n"
188 " size_t i = 0;\n"
189 " __attribute__((opencl_unroll_hint(10)))\n"
190 " do {\n"
191 " dst[i] = i;\n"
192 " ++i;\n"
193 " } while(i < 100);\n"
194 "}\n",
195 "__kernel void pragma_unroll(__global uint *dst)\n"
196 "{\n"
197 " size_t tid = get_global_id(0);\n"
198 " size_t i = 0;\n"
199 " __attribute__((opencl_unroll_hint(100)))\n"
200 " do {\n"
201 " dst[i] = i;\n"
202 " ++i;\n"
203 " } while(i < 100);\n"
204 "}\n",
205 "__kernel void pragma_unroll(__global uint *dst)\n"
206 "{\n"
207 " size_t tid = get_global_id(0);\n"
208 " size_t n = (tid + 1) * 100;\n"
209 " size_t i = 0;\n"
210 " __attribute__((opencl_unroll_hint))\n"
211 " do {\n"
212 " dst[i] = i;\n"
213 " ++i;\n"
214 " } while(i < n);\n"
215 "}\n",
216 "__kernel void pragma_unroll(__global uint *dst)\n"
217 "{\n"
218 " size_t tid = get_global_id(0);\n"
219 " size_t n = (tid + 1) * 100;\n"
220 " size_t i = 0;\n"
221 " __attribute__((opencl_unroll_hint(1)))\n"
222 " do {\n"
223 " dst[i] = i;\n"
224 " ++i;\n"
225 " } while(i < n);\n"
226 "}\n",
227 "__kernel void pragma_unroll(__global uint *dst)\n"
228 "{\n"
229 " size_t tid = get_global_id(0);\n"
230 " size_t n = (tid + 1) * 100;\n"
231 " size_t i = 0;\n"
232 " __attribute__((opencl_unroll_hint(10)))\n"
233 " do {\n"
234 " dst[i] = i;\n"
235 " ++i;\n"
236 " } while(i < n);\n"
237 "}\n",
238 "__kernel void pragma_unroll(__global uint *dst)\n"
239 "{\n"
240 " size_t tid = get_global_id(0);\n"
241 " size_t n = (tid + 1) * 100;\n"
242 " size_t i = 0;\n"
243 " __attribute__((opencl_unroll_hint(100)))\n"
244 " do {\n"
245 " dst[i] = i;\n"
246 " ++i;\n"
247 " } while(i < n);\n"
248 "}\n",
249 };
250
test_pragma_unroll(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)251 int test_pragma_unroll(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
252 const size_t ELEMENT_NUM = 100;
253 const size_t KERNEL_NUM = 24;
254
255 cl_int error;
256
257 //execute all kernels and check if the results are as expected
258 for (size_t kernelIdx = 0; kernelIdx < KERNEL_NUM; ++kernelIdx) {
259 clProgramWrapper program;
260 clKernelWrapper kernel;
261 if (create_single_kernel_helper(
262 context, &program, &kernel, 1,
263 (const char **)&pragma_unroll_kernels[kernelIdx], "pragma_unroll"))
264 {
265 log_error("The program we attempted to compile was: \n%s\n",
266 pragma_unroll_kernels[kernelIdx]);
267 return -1;
268 }
269
270 clMemWrapper buffer = clCreateBuffer(context, CL_MEM_READ_WRITE, ELEMENT_NUM * sizeof(cl_uint), NULL, &error);
271 test_error(error, "clCreateBuffer failed");
272
273 error = clSetKernelArg(kernel, 0, sizeof(buffer), &buffer);
274 test_error(error, "clSetKernelArg failed");
275
276 //only one thread should be enough to verify if kernel is fully functional
277 size_t workSize = 1;
278 error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &workSize, NULL, 0, NULL, NULL);
279 test_error(error, "clEnqueueNDRangeKernel failed");
280
281 std::vector<cl_uint> results(ELEMENT_NUM, 0);
282 error = clEnqueueReadBuffer(queue, buffer, CL_TRUE, 0, ELEMENT_NUM * sizeof(cl_uint), &results[0], 0, NULL, NULL);
283 test_error(error, "clEnqueueReadBuffer failed");
284
285 for (size_t i = 0; i < ELEMENT_NUM; ++i) {
286 if (results[i] != i) {
287 log_error("Kernel %d returned invalid result. Test: %d, expected: %d\n", kernelIdx + 1, results[i], i);
288 return -1;
289 }
290 }
291 }
292
293 return 0;
294 }
295