1 //
2 // Copyright (c) 2022 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 <algorithm>
18 #include <numeric>
19 #include <string>
20 #include <vector>
21
22 #include "procs.h"
23 #include "harness/integer_ops_test_info.h"
24 #include "harness/testHarness.h"
25
cpu_bit_reverse(T base)26 template <typename T> static T cpu_bit_reverse(T base)
27 {
28 T result = 0;
29
30 const size_t count = sizeof(T) * 8;
31 for (size_t i = 0; i < count; i++)
32 {
33 if (base & ((T)1 << i))
34 {
35 result |= ((T)1 << (count - i - 1));
36 }
37 }
38 return result;
39 }
40
41 template <typename T>
calculate_reference(std::vector<T> & ref,const std::vector<T> & base)42 static void calculate_reference(std::vector<T>& ref, const std::vector<T>& base)
43 {
44 ref.resize(base.size());
45 for (size_t i = 0; i < base.size(); i++)
46 {
47 ref[i] = cpu_bit_reverse(base[i]);
48 }
49 }
50
51 static constexpr const char* kernel_source = R"CLC(
52 __kernel void test_bit_reverse(__global TYPE* dst, __global TYPE* base)
53 {
54 int index = get_global_id(0);
55 dst[index] = bit_reverse(base[index]);
56 }
57 )CLC";
58
59 static constexpr const char* kernel_source_vec3 = R"CLC(
60 __kernel void test_bit_reverse(__global BASETYPE* dst, __global BASETYPE* base)
61 {
62 int index = get_global_id(0);
63 TYPE s = vload3(index, base);
64 TYPE d = bit_reverse(s);
65 vstore3(d, index, dst);
66 }
67 )CLC";
68
69 template <typename T, size_t N>
test_vectype(cl_device_id device,cl_context context,cl_command_queue queue)70 static int test_vectype(cl_device_id device, cl_context context,
71 cl_command_queue queue)
72 {
73 cl_int error = CL_SUCCESS;
74
75 clProgramWrapper program;
76 clKernelWrapper kernel;
77
78 std::string buildOptions{ "-DTYPE=" };
79 buildOptions += TestInfo<T>::deviceTypeName;
80 if (N > 1)
81 {
82 buildOptions += std::to_string(N);
83 }
84 buildOptions += " -DBASETYPE=";
85 buildOptions += TestInfo<T>::deviceTypeName;
86
87 const size_t ELEMENTS_TO_TEST = 65536;
88 std::vector<T> base(ELEMENTS_TO_TEST * N);
89 fill_vector_with_random_data(base);
90
91 std::vector<T> reference;
92 calculate_reference(reference, base);
93
94 const char* source = (N == 3) ? kernel_source_vec3 : kernel_source;
95 error =
96 create_single_kernel_helper(context, &program, &kernel, 1, &source,
97 "test_bit_reverse", buildOptions.c_str());
98 test_error(error, "Unable to create test_bit_reverse kernel");
99
100 clMemWrapper src;
101 clMemWrapper dst;
102
103 dst =
104 clCreateBuffer(context, 0, reference.size() * sizeof(T), NULL, &error);
105 test_error(error, "Unable to create output buffer");
106
107 src = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, base.size() * sizeof(T),
108 base.data(), &error);
109 test_error(error, "Unable to create base buffer");
110
111 error = clSetKernelArg(kernel, 0, sizeof(dst), &dst);
112 test_error(error, "Unable to set output buffer kernel arg");
113
114 error = clSetKernelArg(kernel, 1, sizeof(src), &src);
115 test_error(error, "Unable to set base buffer kernel arg");
116
117 size_t global_work_size[] = { reference.size() / N };
118 error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_work_size,
119 NULL, 0, NULL, NULL);
120 test_error(error, "Unable to enqueue test kernel");
121
122 error = clFinish(queue);
123 test_error(error, "clFinish failed after test kernel");
124
125 std::vector<T> results(reference.size(), 99);
126 error =
127 clEnqueueReadBuffer(queue, dst, CL_TRUE, 0, results.size() * sizeof(T),
128 results.data(), 0, NULL, NULL);
129 test_error(error, "Unable to read data after test kernel");
130
131 if (results != reference)
132 {
133 log_error("Result buffer did not match reference buffer!\n");
134 return TEST_FAIL;
135 }
136
137 return TEST_PASS;
138 }
139
140 template <typename T>
test_type(cl_device_id device,cl_context context,cl_command_queue queue)141 static int test_type(cl_device_id device, cl_context context,
142 cl_command_queue queue)
143 {
144 log_info(" testing type %s\n", TestInfo<T>::deviceTypeName);
145
146 return test_vectype<T, 1>(device, context, queue)
147 | test_vectype<T, 2>(device, context, queue)
148 | test_vectype<T, 3>(device, context, queue)
149 | test_vectype<T, 4>(device, context, queue)
150 | test_vectype<T, 8>(device, context, queue)
151 | test_vectype<T, 16>(device, context, queue);
152 }
153
test_extended_bit_ops_reverse(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)154 int test_extended_bit_ops_reverse(cl_device_id device, cl_context context,
155 cl_command_queue queue, int num_elements)
156 {
157 if (is_extension_available(device, "cl_khr_extended_bit_ops"))
158 {
159 int result = TEST_PASS;
160
161 result |= test_type<cl_char>(device, context, queue);
162 result |= test_type<cl_uchar>(device, context, queue);
163 result |= test_type<cl_short>(device, context, queue);
164 result |= test_type<cl_ushort>(device, context, queue);
165 result |= test_type<cl_int>(device, context, queue);
166 result |= test_type<cl_uint>(device, context, queue);
167 if (gHasLong)
168 {
169 result |= test_type<cl_long>(device, context, queue);
170 result |= test_type<cl_ulong>(device, context, queue);
171 }
172 return result;
173 }
174
175 log_info("cl_khr_extended_bit_ops is not supported\n");
176 return TEST_SKIPPED_ITSELF;
177 }
178