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 "harness/compat.h"
17
18 #include <stdio.h>
19 #include <stdlib.h>
20 #include <string.h>
21 #include <sys/types.h>
22 #include <sys/stat.h>
23
24 #include <algorithm>
25 #include <vector>
26
27 #include "procs.h"
28
29 namespace {
30 const char* constant_kernel_code = R"(
31 __kernel void constant_kernel(__global float *out, __constant float *tmpF, __constant int *tmpI)
32 {
33 int tid = get_global_id(0);
34
35 float ftmp = tmpF[tid];
36 float Itmp = tmpI[tid];
37 out[tid] = ftmp * Itmp;
38 }
39 )";
40
41 const char* loop_constant_kernel_code = R"(
42 kernel void loop_constant_kernel(global float *out, constant float *i_pos, int num)
43 {
44 int tid = get_global_id(0);
45 float sum = 0;
46 for (int i = 0; i < num; i++) {
47 float pos = i_pos[i*3];
48 sum += pos;
49 }
50 out[tid] = sum;
51 }
52 )";
53
54
verify(std::vector<cl_float> & tmpF,std::vector<cl_int> & tmpI,std::vector<cl_float> & out)55 int verify(std::vector<cl_float>& tmpF, std::vector<cl_int>& tmpI,
56 std::vector<cl_float>& out)
57 {
58 for (int i = 0; i < out.size(); i++)
59 {
60 float f = tmpF[i] * tmpI[i];
61 if (out[i] != f)
62 {
63 log_error("CONSTANT test failed\n");
64 return -1;
65 }
66 }
67
68 log_info("CONSTANT test passed\n");
69 return 0;
70 }
71
verify_loop_constant(const std::vector<cl_float> & tmp,std::vector<cl_float> & out,cl_int l)72 int verify_loop_constant(const std::vector<cl_float>& tmp,
73 std::vector<cl_float>& out, cl_int l)
74 {
75 float sum = 0;
76 for (int j = 0; j < l; ++j) sum += tmp[j * 3];
77
78 auto predicate = [&sum](cl_float elem) { return sum != elem; };
79
80 if (std::any_of(out.cbegin(), out.cend(), predicate))
81 {
82 log_error("loop CONSTANT test failed\n");
83 return -1;
84 }
85
86 log_info("loop CONSTANT test passed\n");
87 return 0;
88 }
89
generate_random_inputs(std::vector<T> & v)90 template <typename T> void generate_random_inputs(std::vector<T>& v)
91 {
92 RandomSeed seed(gRandomSeed);
93
94 auto random_generator = [&seed]() {
95 return static_cast<T>(get_random_float(-0x02000000, 0x02000000, seed));
96 };
97
98 std::generate(v.begin(), v.end(), random_generator);
99 }
100 }
101
test_constant(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)102 int test_constant(cl_device_id device, cl_context context,
103 cl_command_queue queue, int num_elements)
104 {
105 clMemWrapper streams[3];
106 clProgramWrapper program;
107 clKernelWrapper kernel;
108
109 size_t global_threads[3];
110 int err;
111 cl_ulong maxSize, maxGlobalSize, maxAllocSize;
112 size_t num_floats, num_ints, constant_values;
113 RoundingMode oldRoundMode;
114 int isRTZ = 0;
115
116 /* Verify our test buffer won't be bigger than allowed */
117 err = clGetDeviceInfo(device, CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE,
118 sizeof(maxSize), &maxSize, 0);
119 test_error(err, "Unable to get max constant buffer size");
120 log_info("Device reports CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE %llu bytes.\n",
121 maxSize);
122
123 // Limit test buffer size to 1/4 of CL_DEVICE_GLOBAL_MEM_SIZE
124 err = clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE,
125 sizeof(maxGlobalSize), &maxGlobalSize, 0);
126 test_error(err, "Unable to get CL_DEVICE_GLOBAL_MEM_SIZE");
127
128 maxSize = std::min(maxSize, maxGlobalSize / 4);
129
130 err = clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE,
131 sizeof(maxAllocSize), &maxAllocSize, 0);
132 test_error(err, "Unable to get CL_DEVICE_MAX_MEM_ALLOC_SIZE");
133
134 maxSize = std::min(maxSize, maxAllocSize);
135
136 maxSize /= 4;
137 num_ints = static_cast<size_t>(maxSize / sizeof(cl_int));
138 num_floats = static_cast<size_t>(maxSize / sizeof(cl_float));
139 constant_values = std::min(num_floats, num_ints);
140
141
142 log_info(
143 "Test will attempt to use %lu bytes with one %lu byte constant int "
144 "buffer and one %lu byte constant float buffer.\n",
145 constant_values * sizeof(cl_int) + constant_values * sizeof(cl_float),
146 constant_values * sizeof(cl_int), constant_values * sizeof(cl_float));
147
148 std::vector<cl_int> tmpI(constant_values);
149 std::vector<cl_float> tmpF(constant_values);
150 std::vector<cl_float> out(constant_values);
151
152
153 streams[0] =
154 clCreateBuffer(context, CL_MEM_READ_WRITE,
155 sizeof(cl_float) * constant_values, nullptr, &err);
156 test_error(err, "clCreateBuffer failed");
157
158 streams[1] =
159 clCreateBuffer(context, CL_MEM_READ_WRITE,
160 sizeof(cl_float) * constant_values, nullptr, &err);
161 test_error(err, "clCreateBuffer failed");
162
163 streams[2] =
164 clCreateBuffer(context, CL_MEM_READ_WRITE,
165 sizeof(cl_int) * constant_values, nullptr, &err);
166 test_error(err, "clCreateBuffer failed");
167
168 generate_random_inputs(tmpI);
169 generate_random_inputs(tmpF);
170
171 err = clEnqueueWriteBuffer(queue, streams[1], CL_TRUE, 0,
172 sizeof(cl_float) * constant_values, tmpF.data(),
173 0, nullptr, nullptr);
174 test_error(err, "clEnqueueWriteBuffer failed");
175 err = clEnqueueWriteBuffer(queue, streams[2], CL_TRUE, 0,
176 sizeof(cl_int) * constant_values, tmpI.data(), 0,
177 nullptr, nullptr);
178 test_error(err, "clEnqueueWriteBuffer faile.");
179
180 err = create_single_kernel_helper(context, &program, &kernel, 1,
181 &constant_kernel_code, "constant_kernel");
182 test_error(err, "Failed to create kernel and program");
183
184
185 err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]);
186 err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]);
187 err |= clSetKernelArg(kernel, 2, sizeof streams[2], &streams[2]);
188 test_error(err, "clSetKernelArgs failed");
189
190 global_threads[0] = constant_values;
191 err = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, global_threads,
192 nullptr, 0, nullptr, nullptr);
193 test_error(err, "clEnqueueNDRangeKernel failed");
194
195 err = clEnqueueReadBuffer(queue, streams[0], CL_TRUE, 0,
196 sizeof(cl_float) * constant_values, out.data(), 0,
197 nullptr, nullptr);
198 test_error(err, "clEnqueueReadBuffer failed");
199
200 // If we only support rtz mode
201 if (CL_FP_ROUND_TO_ZERO == get_default_rounding_mode(device) && gIsEmbedded)
202 {
203 oldRoundMode = set_round(kRoundTowardZero, kfloat);
204 isRTZ = 1;
205 }
206
207 err = verify(tmpF, tmpI, out);
208
209 if (isRTZ) (void)set_round(oldRoundMode, kfloat);
210
211 // Loop constant buffer test
212 clProgramWrapper loop_program;
213 clKernelWrapper loop_kernel;
214 cl_int limit = 2;
215
216 memset(out.data(), 0, sizeof(cl_float) * constant_values);
217 err = create_single_kernel_helper(context, &loop_program, &loop_kernel, 1,
218 &loop_constant_kernel_code,
219 "loop_constant_kernel");
220 test_error(err, "Failed to create kernel and program");
221
222 err = clSetKernelArg(loop_kernel, 0, sizeof streams[0], &streams[0]);
223 err |= clSetKernelArg(loop_kernel, 1, sizeof streams[1], &streams[1]);
224 err |= clSetKernelArg(loop_kernel, 2, sizeof(limit), &limit);
225 test_error(err, "clSetKernelArgs failed");
226
227 err = clEnqueueNDRangeKernel(queue, loop_kernel, 1, nullptr, global_threads,
228 nullptr, 0, nullptr, nullptr);
229 test_error(err, "clEnqueueNDRangeKernel failed");
230
231 err = clEnqueueReadBuffer(queue, streams[0], CL_TRUE, 0,
232 sizeof(cl_float) * constant_values, out.data(), 0,
233 nullptr, nullptr);
234 test_error(err, "clEnqueueReadBuffer failed");
235
236 err = verify_loop_constant(tmpF, out, limit);
237
238
239 return err;
240 }
241