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 }
177
178 destroyBufferStruct(pBuffers, pClState);
179
180 destroyClState(pClState);
181
182
183 // vlog_error("%s : implementation incomplete : FAIL\n", testName);
184 return 0; // -1; // fails on account of not being written.
185 }
186
187 static const char* patterns[] = {
188 ".EXTENSIONS.\n"
189 "__kernel void test_step_type(__global .TYPE..NUM. *source, __global int "
190 "*dest)\n"
191 "{\n"
192 " int tid = get_global_id(0);\n"
193 " dest[tid] = vec_step(.TYPE..NUM.);\n"
194 "\n"
195 "}\n",
196
197 ".EXTENSIONS.\n"
198 "__kernel void test_step_var(__global .TYPE..NUM. *source, __global int "
199 "*dest)\n"
200 "{\n"
201 " int tid = get_global_id(0);\n"
202 " dest[tid] = vec_step(source[tid]);\n"
203 "\n"
204 "}\n",
205
206 ".EXTENSIONS.\n"
207 " typedef .TYPE..NUM. TypeToTest;\n"
208 "__kernel void test_step_typedef_type(__global TypeToTest *source, "
209 "__global int *dest)\n"
210 "{\n"
211 " int tid = get_global_id(0);\n"
212 " dest[tid] = vec_step(TypeToTest);\n"
213 "\n"
214 "}\n",
215
216 ".EXTENSIONS.\n"
217 " typedef .TYPE..NUM. TypeToTest;\n"
218 "__kernel void test_step_typedef_var(__global TypeToTest *source, __global "
219 "int *dest)\n"
220 "{\n"
221 " int tid = get_global_id(0);\n"
222 " dest[tid] = vec_step(source[tid]);\n"
223 "\n"
224 "}\n",
225 };
226
227 /*
228 test_step_type,
229 test_step_var,
230 test_step_typedef_type,
231 test_step_typedef_var,
232 */
233
test_step_type(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)234 int test_step_type(cl_device_id deviceID, cl_context context,
235 cl_command_queue queue, int num_elements)
236 {
237 return test_step_internal(deviceID, context, queue, patterns[0],
238 "test_step_type");
239 }
240
test_step_var(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)241 int test_step_var(cl_device_id deviceID, cl_context context,
242 cl_command_queue queue, int num_elements)
243 {
244 return test_step_internal(deviceID, context, queue, patterns[1],
245 "test_step_var");
246 }
247
test_step_typedef_type(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)248 int test_step_typedef_type(cl_device_id deviceID, cl_context context,
249 cl_command_queue queue, int num_elements)
250 {
251 return test_step_internal(deviceID, context, queue, patterns[2],
252 "test_step_typedef_type");
253 }
254
test_step_typedef_var(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)255 int test_step_typedef_var(cl_device_id deviceID, cl_context context,
256 cl_command_queue queue, int num_elements)
257 {
258 return test_step_internal(deviceID, context, queue, patterns[3],
259 "test_step_typedef_var");
260 }
261