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