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 if (type == NON_NULL_PATH) {
68 status = clSetKernelArg(kernel, 0, sizeof(cl_mem), &test_buf);
69 typestr = "non-NULL";
70 } else if (type == ADDROF_NULL_PATH) {
71 test_buf = NULL;
72 status = clSetKernelArg(kernel, 0, sizeof(cl_mem), &test_buf);
73 typestr = "&NULL";
74 } else if (type == NULL_PATH) {
75 status = clSetKernelArg(kernel, 0, sizeof(cl_mem), NULL);
76 typestr = "NULL";
77 }
78
79 log_info("Testing setKernelArgs with %s buffer.\n", typestr);
80
81 if (status != CL_SUCCESS) {
82 log_error("clSetKernelArg failed with status: %d\n", status);
83 return FAILURE; // no point in continuing *this* test
84 }
85
86 size_t global = NITEMS;
87 status = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global,
88 NULL, 0, NULL, NULL);
89 test_error(status, "NDRangeKernel failed.");
90
91 if (gIsEmbedded)
92 {
93 cl_int* host_result = (cl_int*)malloc(NITEMS*sizeof(cl_int));
94 status = clEnqueueReadBuffer(queue, result_buf, CL_TRUE, 0,
95 sizeof(cl_int)*NITEMS, host_result, 0, NULL, NULL);
96 test_error(status, "ReadBuffer failed.");
97 // in the non-null case, we expect NONZERO values:
98 if (type == NON_NULL_PATH) {
99 for (i=0; i<NITEMS; i++) {
100 if (host_result[i] == 0) {
101 log_error("failure: item %d in the result buffer was unexpectedly NULL.\n", i);
102 test_success = FAILURE; break;
103 }
104 }
105
106 } else if (type == ADDROF_NULL_PATH || type == NULL_PATH) {
107 for (i=0; i<NITEMS; i++) {
108 if (host_result[i] != 0) {
109 log_error("failure: item %d in the result buffer was unexpectedly non-NULL.\n", i);
110 test_success = FAILURE; break;
111 }
112 }
113 }
114 free(host_result);
115 }
116 else
117 {
118 cl_long* host_result = (cl_long*)malloc(NITEMS*sizeof(cl_long));
119 status = clEnqueueReadBuffer(queue, result_buf, CL_TRUE, 0,
120 sizeof(cl_long)*NITEMS, host_result, 0, NULL, NULL);
121 test_error(status, "ReadBuffer failed.");
122 // in the non-null case, we expect NONZERO values:
123 if (type == NON_NULL_PATH) {
124 for (i=0; i<NITEMS; i++) {
125 if (host_result[i] == 0) {
126 log_error("failure: item %d in the result buffer was unexpectedly NULL.\n", i);
127 test_success = FAILURE; break;
128 }
129 }
130 } else if (type == ADDROF_NULL_PATH || type == NULL_PATH) {
131 for (i=0; i<NITEMS; i++) {
132 if (host_result[i] != 0) {
133 log_error("failure: item %d in the result buffer was unexpectedly non-NULL.\n", i);
134 test_success = FAILURE; break;
135 }
136 }
137 }
138 free(host_result);
139 }
140
141 if (test_success == SUCCESS) {
142 log_info("\t%s ok.\n", typestr);
143 }
144
145 return test_success;
146 }
147
test_null_buffer_arg(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)148 int test_null_buffer_arg(cl_device_id device, cl_context context,
149 cl_command_queue queue, int num_elements)
150 {
151 unsigned int test_success = 0;
152 unsigned int i;
153 unsigned int buffer_size;
154 cl_int status;
155 cl_program program;
156 cl_kernel kernel;
157
158 // prep kernel:
159 if (gIsEmbedded)
160 status = create_single_kernel_helper(context, &program, NULL, 1, &kernel_string, NULL);
161 else
162 status = create_single_kernel_helper(context, &program, NULL, 1, &kernel_string_long, NULL);
163
164 test_error(status, "Unable to build test program");
165
166 kernel = clCreateKernel(program, "test_kernel", &status);
167 test_error(status, "CreateKernel failed.");
168
169 cl_mem dev_src = clCreateBuffer(context, CL_MEM_READ_ONLY, NITEMS*sizeof(cl_float),
170 NULL, NULL);
171
172 if (gIsEmbedded)
173 buffer_size = NITEMS*sizeof(cl_int);
174 else
175 buffer_size = NITEMS*sizeof(cl_long);
176
177 cl_mem dev_dst = clCreateBuffer(context, CL_MEM_WRITE_ONLY, buffer_size,
178 NULL, NULL);
179
180 // set the destination buffer normally:
181 status = clSetKernelArg(kernel, 1, sizeof(cl_mem), &dev_dst);
182 test_error(status, "SetKernelArg failed.");
183
184 //
185 // we test three cases:
186 //
187 // - typical case, used everyday: non-null buffer
188 // - the case of src as &NULL (the spec-compliance test)
189 // - the case of src as NULL (the backwards-compatibility test, Apple only)
190 //
191
192 test_success = test_setargs_and_execution(queue, kernel, dev_src, dev_dst, NON_NULL_PATH);
193 test_success |= test_setargs_and_execution(queue, kernel, dev_src, dev_dst, ADDROF_NULL_PATH);
194
195 #ifdef __APPLE__
196 test_success |= test_setargs_and_execution(queue, kernel, dev_src, dev_dst, NULL_PATH);
197 #endif
198
199 // clean up:
200 if (dev_src) clReleaseMemObject(dev_src);
201 clReleaseMemObject(dev_dst);
202 clReleaseKernel(kernel);
203 clReleaseProgram(program);
204
205 return test_success;
206 }
207