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 #include "harness/os_helpers.h"
18
19 const char *define_kernel_code[] = {
20 " #define VALUE\n"
21 "__kernel void define_test(__global int *src, __global int *dstA, __global int *dstB)\n"
22 "{\n"
23 " int tid = get_global_id(0);\n"
24 "#ifdef VALUE\n"
25 " dstA[tid] = src[tid] * 2;\n"
26 "#else\n"
27 " dstA[tid] = src[tid] * 4;\n"
28 "#endif\n"
29 "\n"
30 "#undef VALUE\n"
31 "#ifdef VALUE\n"
32 " dstB[tid] = src[tid] * 2;\n"
33 "#else\n"
34 " dstB[tid] = src[tid] * 4;\n"
35 "#endif\n"
36 "\n"
37 "}\n"};
38
39
40
41
test_preprocessor_define_udef(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)42 int test_preprocessor_define_udef(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
43
44 cl_int error;
45 clKernelWrapper kernel;
46 clProgramWrapper program;
47 clMemWrapper buffer[3];
48 cl_int *srcData, *resultData;
49 int i;
50 MTdata d;
51
52 error = create_single_kernel_helper(context, &program, &kernel, 1, define_kernel_code, "define_test");
53 if (error)
54 return -1;
55
56 buffer[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, num_elements*sizeof(cl_int), NULL, &error);
57 test_error( error, "clCreateBuffer failed");
58 buffer[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, num_elements*sizeof(cl_int), NULL, &error);
59 test_error( error, "clCreateBuffer failed");
60 buffer[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, num_elements*sizeof(cl_int), NULL, &error);
61 test_error( error, "clCreateBuffer failed");
62
63 srcData = (cl_int*)malloc(sizeof(cl_int)*num_elements);
64 if (srcData == NULL) {
65 log_error("Failed to allocate storage for source data (%d cl_ints).\n", num_elements);
66 return -1;
67 }
68
69 d = init_genrand( gRandomSeed );
70 for (i=0; i<num_elements; i++)
71 srcData[i] = (int)get_random_float(-1024, 1024,d);
72 free_mtdata(d); d = NULL;
73
74 resultData = (cl_int*)malloc(sizeof(cl_int)*num_elements);
75 if (resultData == NULL) {
76 free(srcData);
77 log_error("Failed to allocate storage for result data (%d cl_ints).\n", num_elements);
78 return -1;
79 }
80
81 error = clSetKernelArg(kernel, 0, sizeof(buffer[0]), &buffer[0]);
82 test_error(error, "clSetKernelArg failed");
83 error = clSetKernelArg(kernel, 1, sizeof(buffer[1]), &buffer[1]);
84 test_error(error, "clSetKernelArg failed");
85 error = clSetKernelArg(kernel, 2, sizeof(buffer[2]), &buffer[2]);
86 test_error(error, "clSetKernelArg failed");
87
88
89 error = clEnqueueWriteBuffer(queue, buffer[0], CL_TRUE, 0, num_elements*sizeof(cl_int), srcData, 0, NULL, NULL);
90 test_error(error, "clEnqueueWriteBuffer failed");
91
92 size_t threads[3] = {num_elements, 0, 0};
93 error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL);
94 test_error(error, "clEnqueueNDRangeKernel failed");
95
96 error = clEnqueueReadBuffer(queue, buffer[1], CL_TRUE, 0, num_elements*sizeof(cl_int), resultData, 0, NULL, NULL);
97 test_error(error, "clEnqueueReadBuffer failed");
98
99 for (i=0; i<num_elements; i++)
100 if (resultData[i] != srcData[i]*2) {
101 free(srcData);
102 free(resultData);
103 return -1;
104 }
105
106 error = clEnqueueReadBuffer(queue, buffer[2], CL_TRUE, 0, num_elements*sizeof(cl_int), resultData, 0, NULL, NULL);
107 test_error(error, "clEnqueueReadBuffer failed");
108
109 for (i=0; i<num_elements; i++)
110 if (resultData[i] != srcData[i]*4) {
111 free(srcData);
112 free(resultData);
113 return -1;
114 }
115
116 free(srcData);
117 free(resultData);
118 return 0;
119 }
120
121
122 const char *include_kernel_code =
123 "#include \"%s\"\n"
124 "__kernel void include_test(__global int *src, __global int *dstA)\n"
125 "{\n"
126 " int tid = get_global_id(0);\n"
127 "#ifdef HEADER_FOUND\n"
128 " dstA[tid] = HEADER_FOUND;\n"
129 "#else\n"
130 " dstA[tid] = 0;\n"
131 "#endif\n"
132 "\n"
133 "}\n";
134
135
test_preprocessor_include(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)136 int test_preprocessor_include(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
137
138 cl_int error;
139 clKernelWrapper kernel;
140 clProgramWrapper program;
141 clMemWrapper buffer[2];
142 cl_int *resultData;
143 int i;
144
145 char include_dir[4096] = {0};
146 char include_kernel[4096] = {0};
147
148 char const * sep = get_dir_sep();
149 char const * path = get_exe_dir();
150
151 /* Build with the include directory defined */
152 sprintf(include_dir,"%s%sincludeTestDirectory%stestIncludeFile.h", path, sep, sep);
153 sprintf(include_kernel, include_kernel_code, include_dir);
154 free( (void *) sep );
155 free( (void *) path );
156
157 const char* test_kernel[] = { include_kernel, 0 };
158 error = create_single_kernel_helper(context, &program, &kernel, 1, test_kernel, "include_test");
159 if (error)
160 return -1;
161
162 buffer[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, num_elements*sizeof(cl_int), NULL, &error);
163 test_error( error, "clCreateBuffer failed");
164 buffer[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, num_elements*sizeof(cl_int), NULL, &error);
165 test_error( error, "clCreateBuffer failed");
166
167 resultData = (cl_int*)malloc(sizeof(cl_int)*num_elements);
168 if (resultData == NULL) {
169 log_error("Failed to allocate storage for result data (%d cl_ints).\n", num_elements);
170 return -1;
171 }
172
173 error = clSetKernelArg(kernel, 0, sizeof(buffer[0]), &buffer[0]);
174 test_error(error, "clSetKernelArg failed");
175 error = clSetKernelArg(kernel, 1, sizeof(buffer[1]), &buffer[1]);
176 test_error(error, "clSetKernelArg failed");
177
178 size_t threads[3] = {num_elements, 0, 0};
179 error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL);
180 test_error(error, "clEnqueueNDRangeKernel failed");
181
182 error = clEnqueueReadBuffer(queue, buffer[1], CL_TRUE, 0, num_elements*sizeof(cl_int), resultData, 0, NULL, NULL);
183 test_error(error, "clEnqueueReadBuffer failed");
184
185 for (i=0; i<num_elements; i++)
186 if (resultData[i] != 12) {
187 free(resultData);
188 return -1;
189 }
190
191 free(resultData);
192 return 0;
193 }
194
195
196
197
198 const char *line_error_kernel_code[] = {
199 "__kernel void line_error_test(__global int *dstA)\n"
200 "{\n"
201 " int tid = get_global_id(0);\n"
202 "#line 124 \"fictitious/file/name.c\" \n"
203 "#error some error\n"
204 " dstA[tid] = tid;\n"
205 "\n"
206 "}\n"};
207
208
test_preprocessor_line_error(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)209 int test_preprocessor_line_error(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
210
211 cl_int error, error2;
212 clKernelWrapper kernel;
213 clProgramWrapper program;
214 clMemWrapper buffer[2];
215
216 char buildLog[ 1024 * 128 ];
217
218 log_info("test_preprocessor_line_error may report spurious ERRORS in the conformance log.\n");
219
220 /* Create the program object from source */
221 program = clCreateProgramWithSource( context, 1, line_error_kernel_code, NULL, &error );
222 test_error(error, "clCreateProgramWithSource failed");
223
224 /* Compile the program */
225 error2 = clBuildProgram( program, 0, NULL, NULL, NULL, NULL );
226 if (error2) {
227 log_info("Build error detected at clBuildProgram.");
228 } else {
229 log_info("Error not reported by clBuildProgram.\n");
230 }
231
232 cl_build_status status;
233 error = clGetProgramBuildInfo(program, deviceID, CL_PROGRAM_BUILD_STATUS, sizeof(status), &status, NULL);
234 test_error(error, "clGetProgramBuildInfo failed for CL_PROGRAM_BUILD_STATUS");
235 if (status != CL_BUILD_ERROR) {
236 log_error("Build status did not return CL_BUILD_ERROR for a program with #error defined.\n");
237 return -1;
238 } else if (status == CL_BUILD_ERROR || error2) {
239 error2 = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_LOG, sizeof( buildLog ), buildLog, NULL );
240 test_error( error2, "Unable to get program build log" );
241
242 log_info("Build failed as expected with #error in source:\n");
243 log_info( "Build log is: ------------\n" );
244 log_info( "%s\n", buildLog );
245 log_info( "Original source is: ------------\n" );
246 log_info( "%s", line_error_kernel_code[0] );
247 log_info( "\n----------\n" );
248
249 if (strstr(buildLog, "fictitious/file/name.c")) {
250 log_info("Found file name from #line param in log output.\n");
251 } else {
252 log_info("WARNING: Did not find file name from #line param in log output.\n");
253 }
254
255 if (strstr(buildLog, "124")) {
256 log_info("Found line number from #line param in log output.\n");
257 } else {
258 log_info("WARNING: Did not find line number from #line param in log output.\n");
259 }
260
261 log_info("test_preprocessor_line_error PASSED.\n");
262 return 0;
263 }
264
265 /* And create a kernel from it */
266 kernel = clCreateKernel( program, "line_error_test", &error );
267 test_error(error, "clCreateKernel failed");
268
269 buffer[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, num_elements*sizeof(cl_int), NULL, &error);
270 test_error( error, "clCreateBuffer failed");
271
272 error = clSetKernelArg(kernel, 0, sizeof(buffer[0]), &buffer[0]);
273 test_error(error, "clSetKernelArg failed");
274
275 size_t threads[3] = {num_elements, 0, 0};
276 error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL);
277 test_error(error, "clEnqueueNDRangeKernel failed");
278
279 log_error("Program built and ran with #error defined.");
280 return -1;
281 }
282
283
284
285 const char *pragma_kernel_code[] = {
286 "__kernel void pragma_test(__global int *dstA)\n"
287 "{\n"
288 "#pragma A fool thinks himself to be wise, but a wise man knows himself to be a fool.\n"
289 " int tid = get_global_id(0);\n"
290 "#pragma\n"
291 " dstA[tid] = tid;\n"
292 "#pragma mark Though I am not naturally honest, I am so sometimes by chance.\n"
293 "\n"
294 "}\n"};
295
296
test_preprocessor_pragma(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)297 int test_preprocessor_pragma(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
298
299 cl_int error;
300 clKernelWrapper kernel;
301 clProgramWrapper program;
302 clMemWrapper buffer[2];
303 cl_int *resultData;
304 int i;
305
306 error = create_single_kernel_helper(context, &program, &kernel, 1, pragma_kernel_code, "pragma_test");
307 if (error)
308 return -1;
309
310 buffer[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, num_elements*sizeof(cl_int), NULL, &error);
311 test_error( error, "clCreateBuffer failed");
312
313 error = clSetKernelArg(kernel, 0, sizeof(buffer[0]), &buffer[0]);
314 test_error(error, "clSetKernelArg failed");
315
316 size_t threads[3] = {num_elements, 0, 0};
317 error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL);
318 test_error(error, "clEnqueueNDRangeKernel failed");
319
320 resultData = (cl_int*)malloc(sizeof(cl_int)*num_elements);
321 if (resultData == NULL) {
322 log_error("Failed to allocate storage for result data (%d cl_ints).\n", num_elements);
323 return -1;
324 }
325
326 error = clEnqueueReadBuffer(queue, buffer[0], CL_TRUE, 0, num_elements*sizeof(cl_int), resultData, 0, NULL, NULL);
327 test_error(error, "clEnqueueReadBuffer failed");
328
329 for (i=0; i<num_elements; i++)
330 if (resultData[i] != i) {
331 free(resultData);
332 return -1;
333 }
334
335 free(resultData);
336 return 0;
337 }
338
339
340
341