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 "testBase.h"
17
18
19 #include "harness/conversions.h"
20 #include "harness/typeWrappers.h"
21 #include "harness/testHarness.h"
22
23 #include "structs.h"
24
25 #include "defines.h"
26
27 #include "type_replacer.h"
28
29
30 /*
31 test_step_type,
32 test_step_var,
33 test_step_typedef_type,
34 test_step_typedef_var,
35 */
36
37
test_step_internal(cl_device_id deviceID,cl_context context,cl_command_queue queue,const char * pattern,const char * testName)38 int test_step_internal(cl_device_id deviceID, cl_context context,
39 cl_command_queue queue, const char* pattern,
40 const char* testName)
41 {
42 int err;
43 int typeIdx, vecSizeIdx;
44
45 char tempBuffer[2048];
46
47 clState* pClState = newClState(deviceID, context, queue);
48 bufferStruct* pBuffers =
49 newBufferStruct(BUFFER_SIZE, BUFFER_SIZE, pClState);
50
51 if (pBuffers == NULL)
52 {
53 destroyClState(pClState);
54 vlog_error("%s : Could not create buffer\n", testName);
55 return -1;
56 }
57
58 // detect whether profile of the device is embedded
59 char profile[1024] = "";
60 err = clGetDeviceInfo(deviceID, CL_DEVICE_PROFILE, sizeof(profile), profile,
61 NULL);
62 if (err)
63 {
64 print_error(err, "clGetDeviceInfo for CL_DEVICE_PROFILE failed\n");
65 return -1;
66 }
67 gIsEmbedded = NULL != strstr(profile, "EMBEDDED_PROFILE");
68
69 for (typeIdx = 0; types[typeIdx] != kNumExplicitTypes; ++typeIdx)
70 {
71 if (types[typeIdx] == kDouble)
72 {
73 // If we're testing doubles, we need to check for support first
74 if (!is_extension_available(deviceID, "cl_khr_fp64"))
75 {
76 log_info("Not testing doubles (unsupported on this device)\n");
77 continue;
78 }
79 }
80
81 if (types[typeIdx] == kLong || types[typeIdx] == kULong)
82 {
83 // If we're testing long/ulong, we need to check for embedded
84 // support
85 if (gIsEmbedded
86 && !is_extension_available(deviceID, "cles_khr_int64"))
87 {
88 log_info("Not testing longs (unsupported on this embedded "
89 "device)\n");
90 continue;
91 }
92 }
93
94 char srcBuffer[2048];
95
96 doSingleReplace(tempBuffer, 2048, pattern, ".EXTENSIONS.",
97 types[typeIdx] == kDouble
98 ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable"
99 : "");
100
101 for (vecSizeIdx = 0; vecSizeIdx < NUM_VECTOR_SIZES; ++vecSizeIdx)
102 {
103 doReplace(srcBuffer, 2048, tempBuffer, ".TYPE.",
104 g_arrTypeNames[typeIdx], ".NUM.",
105 g_arrVecSizeNames[vecSizeIdx]);
106
107 if (srcBuffer[0] == '\0')
108 {
109 vlog_error("%s: failed to fill source buf for type %s%s\n",
110 testName, g_arrTypeNames[typeIdx],
111 g_arrVecSizeNames[vecSizeIdx]);
112 destroyBufferStruct(pBuffers, pClState);
113 destroyClState(pClState);
114 return -1;
115 }
116
117 err = clStateMakeProgram(pClState, srcBuffer, testName);
118 if (err)
119 {
120 vlog_error("%s: Error compiling \"\n%s\n\"", testName,
121 srcBuffer);
122 destroyBufferStruct(pBuffers, pClState);
123 destroyClState(pClState);
124 return -1;
125 }
126
127 err = pushArgs(pBuffers, pClState);
128 if (err != 0)
129 {
130 vlog_error("%s: failed to push args %s%s\n", testName,
131 g_arrTypeNames[typeIdx],
132 g_arrVecSizeNames[vecSizeIdx]);
133 destroyBufferStruct(pBuffers, pClState);
134 destroyClState(pClState);
135 return -1;
136 }
137
138 // now we run the kernel
139 err = runKernel(pClState, 1024);
140 if (err != 0)
141 {
142 vlog_error("%s: runKernel fail (%ld threads) %s%s\n", testName,
143 pClState->m_numThreads, g_arrTypeNames[typeIdx],
144 g_arrVecSizeNames[vecSizeIdx]);
145 destroyBufferStruct(pBuffers, pClState);
146 destroyClState(pClState);
147 return -1;
148 }
149
150 err = retrieveResults(pBuffers, pClState);
151 if (err != 0)
152 {
153 vlog_error("%s: failed to retrieve results %s%s\n", testName,
154 g_arrTypeNames[typeIdx],
155 g_arrVecSizeNames[vecSizeIdx]);
156 destroyBufferStruct(pBuffers, pClState);
157 destroyClState(pClState);
158 return -1;
159 }
160
161 err = checkCorrectnessStep(pBuffers, pClState,
162 g_arrTypeSizes[typeIdx],
163 g_arrVecSizes[vecSizeIdx]);
164
165 if (err != 0)
166 {
167 vlog_error("%s: incorrect results %s%s\n", testName,
168 g_arrTypeNames[typeIdx],
169 g_arrVecSizeNames[vecSizeIdx]);
170 vlog_error("%s: Source was \"\n%s\n\"", testName, srcBuffer);
171 destroyBufferStruct(pBuffers, pClState);
172 destroyClState(pClState);
173 return -1;
174 }
175
176 clStateDestroyProgramAndKernel(pClState);
177 }
178 }
179
180 destroyBufferStruct(pBuffers, pClState);
181
182 destroyClState(pClState);
183
184
185 // vlog_error("%s : implementation incomplete : FAIL\n", testName);
186 return 0; // -1; // fails on account of not being written.
187 }
188
189 static const char* patterns[] = {
190 ".EXTENSIONS.\n"
191 "__kernel void test_step_type(__global .TYPE..NUM. *source, __global int "
192 "*dest)\n"
193 "{\n"
194 " int tid = get_global_id(0);\n"
195 " dest[tid] = vec_step(.TYPE..NUM.);\n"
196 "\n"
197 "}\n",
198
199 ".EXTENSIONS.\n"
200 "__kernel void test_step_var(__global .TYPE..NUM. *source, __global int "
201 "*dest)\n"
202 "{\n"
203 " int tid = get_global_id(0);\n"
204 " dest[tid] = vec_step(source[tid]);\n"
205 "\n"
206 "}\n",
207
208 ".EXTENSIONS.\n"
209 " typedef .TYPE..NUM. TypeToTest;\n"
210 "__kernel void test_step_typedef_type(__global TypeToTest *source, "
211 "__global int *dest)\n"
212 "{\n"
213 " int tid = get_global_id(0);\n"
214 " dest[tid] = vec_step(TypeToTest);\n"
215 "\n"
216 "}\n",
217
218 ".EXTENSIONS.\n"
219 " typedef .TYPE..NUM. TypeToTest;\n"
220 "__kernel void test_step_typedef_var(__global TypeToTest *source, __global "
221 "int *dest)\n"
222 "{\n"
223 " int tid = get_global_id(0);\n"
224 " dest[tid] = vec_step(source[tid]);\n"
225 "\n"
226 "}\n",
227 };
228
229 /*
230 test_step_type,
231 test_step_var,
232 test_step_typedef_type,
233 test_step_typedef_var,
234 */
235
test_step_type(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)236 int test_step_type(cl_device_id deviceID, cl_context context,
237 cl_command_queue queue, int num_elements)
238 {
239 return test_step_internal(deviceID, context, queue, patterns[0],
240 "test_step_type");
241 }
242
test_step_var(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)243 int test_step_var(cl_device_id deviceID, cl_context context,
244 cl_command_queue queue, int num_elements)
245 {
246 return test_step_internal(deviceID, context, queue, patterns[1],
247 "test_step_var");
248 }
249
test_step_typedef_type(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)250 int test_step_typedef_type(cl_device_id deviceID, cl_context context,
251 cl_command_queue queue, int num_elements)
252 {
253 return test_step_internal(deviceID, context, queue, patterns[2],
254 "test_step_typedef_type");
255 }
256
test_step_typedef_var(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)257 int test_step_typedef_var(cl_device_id deviceID, cl_context context,
258 cl_command_queue queue, int num_elements)
259 {
260 return test_step_internal(deviceID, context, queue, patterns[3],
261 "test_step_typedef_var");
262 }
263