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 #ifndef TEST_CONFORMANCE_CLCPP_ATOMICS_ATOMIC_FETCH_HPP
17 #define TEST_CONFORMANCE_CLCPP_ATOMICS_ATOMIC_FETCH_HPP
18
19 #include "../common.hpp"
20 #include "../funcs_test_utils.hpp"
21
22
23 const size_t atomic_bucket_size = 100;
24
25 // -----------------------------------------------------------------------------------
26 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
27 // -----------------------------------------------------------------------------------
28 #if defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
29 template <class func_type, class type>
generate_kernel_atomic_fetch(func_type func)30 std::string generate_kernel_atomic_fetch(func_type func)
31 {
32 std::string in1_value = "input[gid]";
33 std::string out1_value = "output[gid / " + std::to_string(atomic_bucket_size) + "]";
34 std::string function_call = "atomic_" + func.str() + "(&" + out1_value + ", " + in1_value + ")";
35 return
36 "" + func.defs() +
37 "__kernel void test_" + func.str() + "(global " + type_name<type>() + " *input, global atomic_" + type_name<type>() + " *output)\n"
38 "{\n"
39 " size_t gid = get_global_id(0);\n"
40 " " + function_call + ";\n"
41 "}\n";
42 }
43 #else
44 template <class func_type, class type>
generate_kernel_atomic_fetch(func_type func)45 std::string generate_kernel_atomic_fetch(func_type func)
46 {
47 std::string in1_value = "input[gid]";
48 std::string out1_value = "output[gid / " + std::to_string(atomic_bucket_size) + "]";
49 std::string function_call = func.str() + "(" + in1_value + ")";
50 return
51 "" + func.defs() +
52 "" + func.headers() +
53 "#include <opencl_memory>\n"
54 "#include <opencl_work_item>\n"
55 "using namespace cl;\n"
56 "__kernel void test_" + func.str() + "(global_ptr<" + type_name<type>() + "[]> input,"
57 "global_ptr<atomic<" + type_name<type>() + ">[]> output)\n"
58 "{\n"
59 " size_t gid = get_global_id(0);\n"
60 " " + out1_value + "." + function_call + ";\n"
61 "}\n";
62 }
63 #endif
64
65 template<class TYPE, class atomic_fetch>
verify_atomic_fetch(const std::vector<TYPE> & in,const std::vector<TYPE> & out,atomic_fetch op)66 bool verify_atomic_fetch(const std::vector<TYPE> &in, const std::vector<TYPE> &out, atomic_fetch op)
67 {
68 for (size_t i = 0; i < out.size(); i++)
69 {
70 TYPE expected = op.init_out();
71 for (size_t k = 0; k < atomic_bucket_size; k++)
72 {
73 const size_t in_i = i * atomic_bucket_size + k;
74 if (in_i >= in.size())
75 break;
76 expected = op(expected, in[in_i]);
77 }
78 if (expected != out[i])
79 {
80 print_error_msg(expected, out[i], i, op);
81 return false;
82 }
83 }
84 return true;
85 }
86
87 template <class atomic_fetch>
test_atomic_fetch_func(cl_device_id device,cl_context context,cl_command_queue queue,size_t count,atomic_fetch op)88 int test_atomic_fetch_func(cl_device_id device, cl_context context, cl_command_queue queue, size_t count, atomic_fetch op)
89 {
90 cl_mem buffers[2];
91 cl_program program;
92 cl_kernel kernel;
93 size_t work_size[1];
94 int err;
95
96 typedef typename atomic_fetch::in_type TYPE;
97
98 // Don't run test for unsupported types
99 if (!(type_supported<TYPE>(device)))
100 {
101 return CL_SUCCESS;
102 }
103 if (sizeof(TYPE) == 8 &&
104 (!is_extension_available(device, "cl_khr_int64_base_atomics") ||
105 !is_extension_available(device, "cl_khr_int64_extended_atomics")))
106 {
107 return CL_SUCCESS;
108 }
109
110 std::string code_str = generate_kernel_atomic_fetch<atomic_fetch, TYPE>(op);
111 std::string kernel_name("test_"); kernel_name += op.str();
112
113 // -----------------------------------------------------------------------------------
114 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
115 // -----------------------------------------------------------------------------------
116 // Only OpenCL C++ to SPIR-V compilation
117 #if defined(DEVELOPMENT) && defined(ONLY_SPIRV_COMPILATION)
118 err = create_opencl_kernel(context, &program, &kernel, code_str, kernel_name);
119 RETURN_ON_ERROR(err)
120 return err;
121 // Use OpenCL C kernels instead of OpenCL C++ kernels (test C++ host code)
122 #elif defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
123 err = create_opencl_kernel(context, &program, &kernel, code_str, kernel_name, "-cl-std=CL2.0", false);
124 RETURN_ON_ERROR(err)
125 #else
126 err = create_opencl_kernel(context, &program, &kernel, code_str, kernel_name);
127 RETURN_ON_ERROR(err)
128 #endif
129
130 std::vector<TYPE> input = generate_input<TYPE>(count, op.min1(), op.max1(), std::vector<TYPE>());
131 std::vector<TYPE> output = generate_output<TYPE>((count - 1) / atomic_bucket_size + 1);
132
133 buffers[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(TYPE) * input.size(), NULL, &err);
134 RETURN_ON_CL_ERROR(err, "clCreateBuffer")
135
136 buffers[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(TYPE) * output.size(), NULL, &err);
137 RETURN_ON_CL_ERROR(err, "clCreateBuffer")
138
139 err = clEnqueueWriteBuffer(
140 queue, buffers[0], CL_TRUE, 0, sizeof(TYPE) * input.size(),
141 static_cast<void *>(input.data()), 0, NULL, NULL
142 );
143 RETURN_ON_CL_ERROR(err, "clEnqueueWriteBuffer")
144
145 const TYPE pattern = op.init_out();
146 err = clEnqueueFillBuffer(queue, buffers[1], &pattern, sizeof(pattern), 0, sizeof(TYPE) * output.size(), 0, NULL, NULL);
147 RETURN_ON_CL_ERROR(err, "clEnqueueFillBuffer")
148
149 err = clSetKernelArg(kernel, 0, sizeof(buffers[0]), &buffers[0]);
150 RETURN_ON_CL_ERROR(err, "clSetKernelArg")
151 err = clSetKernelArg(kernel, 1, sizeof(buffers[1]), &buffers[1]);
152 RETURN_ON_CL_ERROR(err, "clSetKernelArg")
153
154 work_size[0] = count;
155 err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, work_size, NULL, 0, NULL, NULL);
156 RETURN_ON_CL_ERROR(err, "clEnqueueNDRangeKernel")
157
158 err = clEnqueueReadBuffer(
159 queue, buffers[1], CL_TRUE, 0, sizeof(TYPE) * output.size(),
160 static_cast<void *>(output.data()), 0, NULL, NULL
161 );
162 RETURN_ON_CL_ERROR(err, "clEnqueueReadBuffer")
163
164 if (!verify_atomic_fetch(input, output, op))
165 {
166 RETURN_ON_ERROR_MSG(-1, "test_%s %s failed", op.str().c_str(), type_name<TYPE>().c_str());
167 }
168 log_info("test_%s %s passed\n", op.str().c_str(), type_name<TYPE>().c_str());
169
170 clReleaseMemObject(buffers[0]);
171 clReleaseMemObject(buffers[1]);
172 clReleaseKernel(kernel);
173 clReleaseProgram(program);
174 return err;
175 }
176
177
178 template<class TYPE>
179 struct atomic_fetch
180 {
181 typedef TYPE in_type;
182
decl_stratomic_fetch183 std::string decl_str()
184 {
185 return type_name<TYPE>();
186 }
187
defsatomic_fetch188 std::string defs()
189 {
190 std::string defs;
191 if (sizeof(TYPE) == 8)
192 {
193 defs += "#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable\n";
194 defs += "#pragma OPENCL EXTENSION cl_khr_int64_extended_atomics : enable\n";
195 }
196 return defs;
197 }
198
headersatomic_fetch199 std::string headers()
200 {
201 return "#include <opencl_atomic>\n";
202 }
203
min1atomic_fetch204 TYPE min1()
205 {
206 return 0;
207 }
208
max1atomic_fetch209 TYPE max1()
210 {
211 return 1000;
212 }
213 };
214
215
216 #define DEF_ATOMIC_FETCH_FUNC(CLASS_NAME, FUNC_NAME, HOST_FUNC_EXPRESSION, INIT_OUT) \
217 template<class TYPE> \
218 struct CLASS_NAME : public atomic_fetch<TYPE> \
219 { \
220 std::string str() \
221 { \
222 return #FUNC_NAME; \
223 } \
224 \
225 TYPE init_out() \
226 { \
227 return INIT_OUT; \
228 } \
229 \
230 TYPE operator()(const TYPE& x, const TYPE& y) \
231 { \
232 return HOST_FUNC_EXPRESSION; \
233 } \
234 };
235
236 DEF_ATOMIC_FETCH_FUNC(atomic_fetch_add, fetch_add, x + y, 0)
237 DEF_ATOMIC_FETCH_FUNC(atomic_fetch_sub, fetch_sub, x - y, (std::numeric_limits<TYPE>::max)())
238
239 DEF_ATOMIC_FETCH_FUNC(atomic_fetch_and, fetch_and, x & y, (std::numeric_limits<TYPE>::max)())
240 DEF_ATOMIC_FETCH_FUNC(atomic_fetch_or, fetch_or, x | y, 0)
241 DEF_ATOMIC_FETCH_FUNC(atomic_fetch_xor, fetch_xor, x ^ y, 0)
242
243 DEF_ATOMIC_FETCH_FUNC(atomic_fetch_max, fetch_max, (std::max)(x, y), 0)
244 DEF_ATOMIC_FETCH_FUNC(atomic_fetch_min, fetch_min, (std::min)(x, y), (std::numeric_limits<TYPE>::max)())
245
246 #undef DEF_ATOMIC_FETCH_FUNC
247
248
AUTO_TEST_CASE(test_atomic_fetch)249 AUTO_TEST_CASE(test_atomic_fetch)
250 (cl_device_id device, cl_context context, cl_command_queue queue, int n_elems)
251 {
252 int error = CL_SUCCESS;
253 int last_error = CL_SUCCESS;
254
255 #define TEST_ATOMIC_MACRO(TEST_CLASS) \
256 last_error = test_atomic_fetch_func( \
257 device, context, queue, n_elems, TEST_CLASS \
258 ); \
259 CHECK_ERROR(last_error) \
260 error |= last_error;
261
262 TEST_ATOMIC_MACRO((atomic_fetch_add<cl_int>()))
263 TEST_ATOMIC_MACRO((atomic_fetch_add<cl_uint>()))
264 TEST_ATOMIC_MACRO((atomic_fetch_add<cl_long>()))
265 TEST_ATOMIC_MACRO((atomic_fetch_add<cl_ulong>()))
266
267 TEST_ATOMIC_MACRO((atomic_fetch_sub<cl_int>()))
268 TEST_ATOMIC_MACRO((atomic_fetch_sub<cl_uint>()))
269 TEST_ATOMIC_MACRO((atomic_fetch_sub<cl_long>()))
270 TEST_ATOMIC_MACRO((atomic_fetch_sub<cl_ulong>()))
271
272 TEST_ATOMIC_MACRO((atomic_fetch_and<cl_int>()))
273 TEST_ATOMIC_MACRO((atomic_fetch_and<cl_uint>()))
274 TEST_ATOMIC_MACRO((atomic_fetch_and<cl_long>()))
275 TEST_ATOMIC_MACRO((atomic_fetch_and<cl_ulong>()))
276
277 TEST_ATOMIC_MACRO((atomic_fetch_or<cl_int>()))
278 TEST_ATOMIC_MACRO((atomic_fetch_or<cl_uint>()))
279 TEST_ATOMIC_MACRO((atomic_fetch_or<cl_long>()))
280 TEST_ATOMIC_MACRO((atomic_fetch_or<cl_ulong>()))
281
282 TEST_ATOMIC_MACRO((atomic_fetch_xor<cl_int>()))
283 TEST_ATOMIC_MACRO((atomic_fetch_xor<cl_uint>()))
284 TEST_ATOMIC_MACRO((atomic_fetch_xor<cl_long>()))
285 TEST_ATOMIC_MACRO((atomic_fetch_xor<cl_ulong>()))
286
287 TEST_ATOMIC_MACRO((atomic_fetch_max<cl_int>()))
288 TEST_ATOMIC_MACRO((atomic_fetch_max<cl_uint>()))
289 TEST_ATOMIC_MACRO((atomic_fetch_max<cl_long>()))
290 TEST_ATOMIC_MACRO((atomic_fetch_max<cl_ulong>()))
291
292 TEST_ATOMIC_MACRO((atomic_fetch_min<cl_int>()))
293 TEST_ATOMIC_MACRO((atomic_fetch_min<cl_uint>()))
294 TEST_ATOMIC_MACRO((atomic_fetch_min<cl_long>()))
295 TEST_ATOMIC_MACRO((atomic_fetch_min<cl_ulong>()))
296
297 #undef TEST_ATOMIC_MACRO
298
299 if (error != CL_SUCCESS)
300 {
301 return -1;
302 }
303 return error;
304 }
305
306 #endif // TEST_CONFORMANCE_CLCPP_ATOMICS_ATOMIC_FETCH_HPP
307