1 //
2 // Copyright (c) 2020 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
17 #include "testBase.h"
18 #include "harness/featureHelpers.h"
19
20 #include <vector>
21
22 static const char* test_kernel = R"CLC(
23 __kernel void test(__global int* dst) {
24 dst[0] = 0;
25 }
26 )CLC";
27
28 // This sub-test checks that CL_DEVICE_OPENCL_C_VERSION meets any API
29 // requirements and that programs can be built for the reported OpenCL C version
30 // and all previous versions.
test_CL_DEVICE_OPENCL_C_VERSION(cl_device_id device,cl_context context)31 static int test_CL_DEVICE_OPENCL_C_VERSION(cl_device_id device,
32 cl_context context)
33 {
34 const Version latest_version = Version(3, 0);
35
36 const Version api_version = get_device_cl_version(device);
37 const Version clc_version = get_device_cl_c_version(device);
38
39 if (api_version > latest_version)
40 {
41 log_info("CL_DEVICE_VERSION is %s, which is bigger than %s.\n"
42 "Need to update the opencl_c_versions test!\n",
43 api_version.to_string().c_str(),
44 latest_version.to_string().c_str());
45 }
46
47 if (clc_version > latest_version)
48 {
49 log_info("CL_DEVICE_OPENCL_C_VERSION is %s, which is bigger than %s.\n"
50 "Need to update the opencl_c_versions test!\n",
51 clc_version.to_string().c_str(),
52 latest_version.to_string().c_str());
53 }
54
55 // For OpenCL 3.0, the minimum required OpenCL C version is OpenCL C 1.2.
56 // For OpenCL 2.x, the minimum required OpenCL C version is OpenCL C 2.0.
57 // For other OpenCL versions, the minimum required OpenCL C version is
58 // the same as the API version.
59 const Version min_clc_version = api_version == Version(3, 0)
60 ? Version(1, 2)
61 : api_version >= Version(2, 0) ? Version(2, 0) : api_version;
62 if (clc_version < min_clc_version)
63 {
64 log_error("The minimum required OpenCL C version for API version %s is "
65 "%s (got %s)!\n",
66 api_version.to_string().c_str(),
67 min_clc_version.to_string().c_str(),
68 clc_version.to_string().c_str());
69 return TEST_FAIL;
70 }
71
72 log_info(" testing compilation based on CL_DEVICE_OPENCL_C_VERSION\n");
73
74 struct TestCase
75 {
76 Version version;
77 const char* buildOptions;
78 };
79
80 std::vector<TestCase> tests;
81 tests.push_back({ Version(1, 1), "-cl-std=CL1.1" });
82 tests.push_back({ Version(1, 2), "-cl-std=CL1.2" });
83 tests.push_back({ Version(2, 0), "-cl-std=CL2.0" });
84 tests.push_back({ Version(3, 0), "-cl-std=CL3.0" });
85
86 for (const auto& testcase : tests)
87 {
88 if (clc_version >= testcase.version)
89 {
90 clProgramWrapper program;
91 cl_int error =
92 create_single_kernel_helper_create_program_for_device(
93 context, device, &program, 1, &test_kernel,
94 testcase.buildOptions);
95 test_error(error, "Unable to build program!");
96
97 log_info(" successfully built program with build options '%s'\n",
98 testcase.buildOptions);
99 }
100 }
101
102 return TEST_PASS;
103 }
104
105 // This sub-test checks that CL_DEVICE_OPENCL_C_ALL_VERSIONS includes any
106 // requirements for the API version, and that programs can be built for all
107 // reported versions.
test_CL_DEVICE_OPENCL_C_ALL_VERSIONS(cl_device_id device,cl_context context)108 static int test_CL_DEVICE_OPENCL_C_ALL_VERSIONS(cl_device_id device,
109 cl_context context)
110 {
111 // For now, the required OpenCL C version is the same as the API version.
112 const Version api_version = get_device_cl_version(device);
113 bool found_api_version = false;
114
115 log_info(
116 " testing compilation based on CL_DEVICE_OPENCL_C_ALL_VERSIONS\n");
117
118 cl_int error = CL_SUCCESS;
119
120 size_t sz = 0;
121 error =
122 clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_ALL_VERSIONS, 0, NULL, &sz);
123 test_error(error, "Unable to query CL_DEVICE_OPENCL_C_ALL_VERSIONS size");
124
125 std::vector<cl_name_version> clc_versions(sz / sizeof(cl_name_version));
126 error = clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_ALL_VERSIONS, sz,
127 clc_versions.data(), NULL);
128 test_error(error, "Unable to query CL_DEVICE_OPENCL_C_FEATURES");
129
130 for (const auto& clc_version : clc_versions)
131 {
132 const unsigned major = CL_VERSION_MAJOR(clc_version.version);
133 const unsigned minor = CL_VERSION_MINOR(clc_version.version);
134
135 if (strcmp(clc_version.name, "OpenCL C") == 0)
136 {
137 if (api_version == Version(major, minor))
138 {
139 found_api_version = true;
140 }
141
142 if (major == 1 && minor == 0)
143 {
144 log_info(
145 " skipping OpenCL C 1.0, there is no -cl-std=CL1.0.\n");
146 continue;
147 }
148
149 std::string buildOptions = "-cl-std=CL";
150 buildOptions += std::to_string(major);
151 buildOptions += ".";
152 buildOptions += std::to_string(minor);
153
154 clProgramWrapper program;
155 error = create_single_kernel_helper_create_program_for_device(
156 context, device, &program, 1, &test_kernel,
157 buildOptions.c_str());
158 test_error(error, "Unable to build program!");
159
160 log_info(" successfully built program with build options '%s'\n",
161 buildOptions.c_str());
162 }
163 else
164 {
165 log_error(" unknown OpenCL C name '%s'.\n", clc_version.name);
166 return TEST_FAIL;
167 }
168 }
169
170 if (!found_api_version)
171 {
172 log_error(" didn't find required OpenCL C version '%s'!\n",
173 api_version.to_string().c_str());
174 return TEST_FAIL;
175 }
176
177 return TEST_PASS;
178 }
179
180 // This sub-test checks that any required features are present for a specific
181 // CL_DEVICE_OPENCL_C_VERSION.
test_CL_DEVICE_OPENCL_C_VERSION_features(cl_device_id device,cl_context context)182 static int test_CL_DEVICE_OPENCL_C_VERSION_features(cl_device_id device,
183 cl_context context)
184 {
185 log_info(" testing for OPENCL_C_VERSION required features\n");
186
187 OpenCLCFeatures features;
188 int error = get_device_cl_c_features(device, features);
189 if (error)
190 {
191 log_error("Couldn't query OpenCL C features for the device!\n");
192 return TEST_FAIL;
193 }
194
195 const Version clc_version = get_device_cl_c_version(device);
196 if (clc_version >= Version(2, 0))
197 {
198 bool has_all_OpenCL_C_20_features =
199 features.supports__opencl_c_atomic_order_acq_rel
200 && features.supports__opencl_c_atomic_order_seq_cst
201 && features.supports__opencl_c_atomic_scope_device
202 && features.supports__opencl_c_atomic_scope_all_devices
203 && features.supports__opencl_c_device_enqueue
204 && features.supports__opencl_c_generic_address_space
205 && features.supports__opencl_c_pipes
206 && features.supports__opencl_c_program_scope_global_variables
207 && features.supports__opencl_c_work_group_collective_functions;
208
209 if (features.supports__opencl_c_images)
210 {
211 has_all_OpenCL_C_20_features = has_all_OpenCL_C_20_features
212 && features.supports__opencl_c_3d_image_writes
213 && features.supports__opencl_c_read_write_images;
214 }
215
216 test_assert_error(
217 has_all_OpenCL_C_20_features,
218 "At least one required OpenCL C 2.0 feature is missing!");
219 }
220
221 return TEST_PASS;
222 }
223
224 // This sub-test checks that all required OpenCL C versions are present for a
225 // specific CL_DEVICE_OPENCL_C_VERSION.
test_CL_DEVICE_OPENCL_C_VERSION_versions(cl_device_id device,cl_context context)226 static int test_CL_DEVICE_OPENCL_C_VERSION_versions(cl_device_id device,
227 cl_context context)
228 {
229 log_info(" testing for OPENCL_C_VERSION required versions\n");
230
231 const Version device_clc_version = get_device_cl_c_version(device);
232
233 std::vector<Version> test_clc_versions;
234 test_clc_versions.push_back(Version(1, 0));
235 test_clc_versions.push_back(Version(1, 1));
236 test_clc_versions.push_back(Version(1, 2));
237 test_clc_versions.push_back(Version(2, 0));
238 test_clc_versions.push_back(Version(3, 0));
239
240 cl_int error = CL_SUCCESS;
241
242 size_t sz = 0;
243 error =
244 clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_ALL_VERSIONS, 0, NULL, &sz);
245 test_error(error, "Unable to query CL_DEVICE_OPENCL_C_ALL_VERSIONS size");
246
247 std::vector<cl_name_version> device_clc_versions(sz
248 / sizeof(cl_name_version));
249 error = clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_ALL_VERSIONS, sz,
250 device_clc_versions.data(), NULL);
251 test_error(error, "Unable to query CL_DEVICE_OPENCL_C_FEATURES");
252
253 for (const auto& test_clc_version : test_clc_versions)
254 {
255 if (device_clc_version >= test_clc_version)
256 {
257 bool found = false;
258 for (const auto& check : device_clc_versions)
259 {
260 const unsigned major = CL_VERSION_MAJOR(check.version);
261 const unsigned minor = CL_VERSION_MINOR(check.version);
262
263 if (strcmp(check.name, "OpenCL C") == 0
264 && test_clc_version == Version(major, minor))
265 {
266 found = true;
267 break;
268 }
269 }
270
271 if (found)
272 {
273 log_info(" found OpenCL C version '%s'\n",
274 test_clc_version.to_string().c_str());
275 }
276 else
277 {
278 log_error("Didn't find OpenCL C version '%s'!\n",
279 test_clc_version.to_string().c_str());
280 return TEST_FAIL;
281 }
282 }
283 }
284
285
286 return TEST_PASS;
287 }
288
test_opencl_c_versions(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)289 int test_opencl_c_versions(cl_device_id device, cl_context context,
290 cl_command_queue queue, int num_elements)
291 {
292 check_compiler_available(device);
293
294 const Version version = get_device_cl_version(device);
295
296 int result = TEST_PASS;
297
298 result |= test_CL_DEVICE_OPENCL_C_VERSION(device, context);
299
300 if (version >= Version(3, 0))
301 {
302 result |= test_CL_DEVICE_OPENCL_C_ALL_VERSIONS(device, context);
303 result |= test_CL_DEVICE_OPENCL_C_VERSION_features(device, context);
304 result |= test_CL_DEVICE_OPENCL_C_VERSION_versions(device, context);
305 }
306
307 return result;
308 }
309