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 <stdio.h>
17 #if defined(__APPLE__)
18 #include <OpenCL/opencl.h>
19 #include <OpenCL/cl_platform.h>
20 #else
21 #include <CL/opencl.h>
22 #include <CL/cl_platform.h>
23 #endif
24 #include "testBase.h"
25 #include "harness/typeWrappers.h"
26 #include "harness/testHarness.h"
27 #include "procs.h"
28
29
30 enum { SUCCESS, FAILURE };
31 typedef enum { NON_NULL_PATH, ADDROF_NULL_PATH, NULL_PATH } test_type;
32
33 #define NITEMS 4096
34
35 /* places the comparison result of value of the src ptr against 0 into each element of the output
36 * array, to allow testing that the kernel actually _gets_ the NULL value */
37 const char *kernel_string_long =
38 "kernel void test_kernel(global float *src, global long *dst)\n"
39 "{\n"
40 " uint tid = get_global_id(0);\n"
41 " dst[tid] = (long)(src != 0);\n"
42 "}\n";
43
44 // For gIsEmbedded
45 const char *kernel_string =
46 "kernel void test_kernel(global float *src, global int *dst)\n"
47 "{\n"
48 " uint tid = get_global_id(0);\n"
49 " dst[tid] = (int)(src != 0);\n"
50 "}\n";
51
52
53 /*
54 * The guts of the test:
55 * call setKernelArgs with a regular buffer, &NULL, or NULL depending on
56 * the value of 'test_type'
57 */
test_setargs_and_execution(cl_command_queue queue,cl_kernel kernel,cl_mem test_buf,cl_mem result_buf,test_type type)58 static int test_setargs_and_execution(cl_command_queue queue, cl_kernel kernel,
59 cl_mem test_buf, cl_mem result_buf, test_type type)
60 {
61 unsigned int test_success = 0;
62
63 unsigned int i;
64 cl_int status;
65 const char *typestr;
66
67 switch (type)
68 {
69 case NON_NULL_PATH:
70 status = clSetKernelArg(kernel, 0, sizeof(cl_mem), &test_buf);
71 typestr = "non-NULL";
72 break;
73 case ADDROF_NULL_PATH:
74 test_buf = NULL;
75 status = clSetKernelArg(kernel, 0, sizeof(cl_mem), &test_buf);
76 typestr = "&NULL";
77 break;
78 case NULL_PATH:
79 status = clSetKernelArg(kernel, 0, sizeof(cl_mem), NULL);
80 typestr = "NULL";
81 break;
82 }
83
84 log_info("Testing setKernelArgs with %s buffer.\n", typestr);
85
86 if (status != CL_SUCCESS) {
87 log_error("clSetKernelArg failed with status: %d\n", status);
88 return FAILURE; // no point in continuing *this* test
89 }
90
91 size_t global = NITEMS;
92 status = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global,
93 NULL, 0, NULL, NULL);
94 test_error(status, "NDRangeKernel failed.");
95
96 if (gIsEmbedded)
97 {
98 cl_int* host_result = (cl_int*)malloc(NITEMS*sizeof(cl_int));
99 status = clEnqueueReadBuffer(queue, result_buf, CL_TRUE, 0,
100 sizeof(cl_int)*NITEMS, host_result, 0, NULL, NULL);
101 test_error(status, "ReadBuffer failed.");
102 // in the non-null case, we expect NONZERO values:
103 if (type == NON_NULL_PATH) {
104 for (i=0; i<NITEMS; i++) {
105 if (host_result[i] == 0) {
106 log_error("failure: item %d in the result buffer was unexpectedly NULL.\n", i);
107 test_success = FAILURE; break;
108 }
109 }
110
111 } else if (type == ADDROF_NULL_PATH || type == NULL_PATH) {
112 for (i=0; i<NITEMS; i++) {
113 if (host_result[i] != 0) {
114 log_error("failure: item %d in the result buffer was unexpectedly non-NULL.\n", i);
115 test_success = FAILURE; break;
116 }
117 }
118 }
119 free(host_result);
120 }
121 else
122 {
123 cl_long* host_result = (cl_long*)malloc(NITEMS*sizeof(cl_long));
124 status = clEnqueueReadBuffer(queue, result_buf, CL_TRUE, 0,
125 sizeof(cl_long)*NITEMS, host_result, 0, NULL, NULL);
126 test_error(status, "ReadBuffer failed.");
127 // in the non-null case, we expect NONZERO values:
128 if (type == NON_NULL_PATH) {
129 for (i=0; i<NITEMS; i++) {
130 if (host_result[i] == 0) {
131 log_error("failure: item %d in the result buffer was unexpectedly NULL.\n", i);
132 test_success = FAILURE; break;
133 }
134 }
135 } else if (type == ADDROF_NULL_PATH || type == NULL_PATH) {
136 for (i=0; i<NITEMS; i++) {
137 if (host_result[i] != 0) {
138 log_error("failure: item %d in the result buffer was unexpectedly non-NULL.\n", i);
139 test_success = FAILURE; break;
140 }
141 }
142 }
143 free(host_result);
144 }
145
146 if (test_success == SUCCESS) {
147 log_info("\t%s ok.\n", typestr);
148 }
149
150 return test_success;
151 }
152
test_null_buffer_arg(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)153 int test_null_buffer_arg(cl_device_id device, cl_context context,
154 cl_command_queue queue, int num_elements)
155 {
156 unsigned int test_success = 0;
157 unsigned int buffer_size;
158 cl_int status;
159 cl_program program;
160 cl_kernel kernel;
161
162 // prep kernel:
163 if (gIsEmbedded)
164 status = create_single_kernel_helper(context, &program, &kernel, 1,
165 &kernel_string, "test_kernel");
166 else
167 status = create_single_kernel_helper(
168 context, &program, &kernel, 1, &kernel_string_long, "test_kernel");
169
170 test_error(status, "Unable to create kernel");
171
172 cl_mem dev_src = clCreateBuffer(context, CL_MEM_READ_ONLY, NITEMS*sizeof(cl_float),
173 NULL, NULL);
174
175 if (gIsEmbedded)
176 buffer_size = NITEMS*sizeof(cl_int);
177 else
178 buffer_size = NITEMS*sizeof(cl_long);
179
180 cl_mem dev_dst = clCreateBuffer(context, CL_MEM_WRITE_ONLY, buffer_size,
181 NULL, NULL);
182
183 // set the destination buffer normally:
184 status = clSetKernelArg(kernel, 1, sizeof(cl_mem), &dev_dst);
185 test_error(status, "SetKernelArg failed.");
186
187 //
188 // we test three cases:
189 //
190 // - typical case, used everyday: non-null buffer
191 // - the case of src as &NULL (the spec-compliance test)
192 // - the case of src as NULL (the backwards-compatibility test, Apple only)
193 //
194
195 test_success = test_setargs_and_execution(queue, kernel, dev_src, dev_dst, NON_NULL_PATH);
196 test_success |= test_setargs_and_execution(queue, kernel, dev_src, dev_dst, ADDROF_NULL_PATH);
197
198 #ifdef __APPLE__
199 test_success |= test_setargs_and_execution(queue, kernel, dev_src, dev_dst, NULL_PATH);
200 #endif
201
202 // clean up:
203 if (dev_src) clReleaseMemObject(dev_src);
204 clReleaseMemObject(dev_dst);
205 clReleaseKernel(kernel);
206 clReleaseProgram(program);
207
208 return test_success;
209 }
210