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 "procs.h"
17
18 #define TEST_VALUE_POSITIVE( string_name, name, value ) \
19 { \
20 if (name < value) { \
21 log_error("FAILED: " string_name ": " #name " < " #value "\n"); \
22 errors++;\
23 } else { \
24 log_info("\t" string_name ": " #name " >= " #value "\n"); \
25 } \
26 }
27
28 #define TEST_VALUE_NEGATIVE( string_name, name, value ) \
29 { \
30 if (name > value) { \
31 log_error("FAILED: " string_name ": " #name " > " #value "\n"); \
32 errors++;\
33 } else { \
34 log_info("\t" string_name ": " #name " <= " #value "\n"); \
35 } \
36 }
37
38 #define TEST_VALUE_EQUAL_LITERAL( string_name, name, value ) \
39 { \
40 if (name != value) { \
41 log_error("FAILED: " string_name ": " #name " != " #value "\n"); \
42 errors++;\
43 } else { \
44 log_info("\t" string_name ": " #name " = " #value "\n"); \
45 } \
46 }
47
48 #define TEST_VALUE_EQUAL( string_name, name, value ) \
49 { \
50 if (name != value) { \
51 log_error("FAILED: " string_name ": " #name " != %a (%17.21g)\n", value, value); \
52 errors++;\
53 } else { \
54 log_info("\t" string_name ": " #name " = %a (%17.21g)\n", value, value); \
55 } \
56 }
57
test_host_numeric_constants(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)58 int test_host_numeric_constants(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
59 {
60 int errors = 0;
61 TEST_VALUE_EQUAL_LITERAL( "CL_CHAR_BIT", CL_CHAR_BIT, 8)
62 TEST_VALUE_EQUAL_LITERAL( "CL_SCHAR_MAX", CL_SCHAR_MAX, 127)
63 TEST_VALUE_EQUAL_LITERAL( "CL_SCHAR_MIN", CL_SCHAR_MIN, (-127-1))
64 TEST_VALUE_EQUAL_LITERAL( "CL_CHAR_MAX", CL_CHAR_MAX, CL_SCHAR_MAX)
65 TEST_VALUE_EQUAL_LITERAL( "CL_CHAR_MIN", CL_CHAR_MIN, CL_SCHAR_MIN)
66 TEST_VALUE_EQUAL_LITERAL( "CL_UCHAR_MAX", CL_UCHAR_MAX, 255)
67 TEST_VALUE_EQUAL_LITERAL( "CL_SHRT_MAX", CL_SHRT_MAX, 32767)
68 TEST_VALUE_EQUAL_LITERAL( "CL_SHRT_MIN", CL_SHRT_MIN, (-32767-1))
69 TEST_VALUE_EQUAL_LITERAL( "CL_USHRT_MAX", CL_USHRT_MAX, 65535)
70 TEST_VALUE_EQUAL_LITERAL( "CL_INT_MAX", CL_INT_MAX, 2147483647)
71 TEST_VALUE_EQUAL_LITERAL( "CL_INT_MIN", CL_INT_MIN, (-2147483647-1))
72 TEST_VALUE_EQUAL_LITERAL( "CL_UINT_MAX", CL_UINT_MAX, 0xffffffffU)
73 TEST_VALUE_EQUAL_LITERAL( "CL_LONG_MAX", CL_LONG_MAX, ((cl_long) 0x7FFFFFFFFFFFFFFFLL))
74 TEST_VALUE_EQUAL_LITERAL( "CL_LONG_MIN", CL_LONG_MIN, ((cl_long) -0x7FFFFFFFFFFFFFFFLL - 1LL))
75 TEST_VALUE_EQUAL_LITERAL( "CL_ULONG_MAX", CL_ULONG_MAX, ((cl_ulong) 0xFFFFFFFFFFFFFFFFULL))
76
77 TEST_VALUE_EQUAL_LITERAL( "CL_FLT_DIG", CL_FLT_DIG, 6)
78 TEST_VALUE_EQUAL_LITERAL( "CL_FLT_MANT_DIG", CL_FLT_MANT_DIG, 24)
79 TEST_VALUE_EQUAL_LITERAL( "CL_FLT_MAX_10_EXP", CL_FLT_MAX_10_EXP, +38)
80 TEST_VALUE_EQUAL_LITERAL( "CL_FLT_MAX_EXP", CL_FLT_MAX_EXP, +128)
81 TEST_VALUE_EQUAL_LITERAL( "CL_FLT_MIN_10_EXP", CL_FLT_MIN_10_EXP, -37)
82 TEST_VALUE_EQUAL_LITERAL( "CL_FLT_MIN_EXP", CL_FLT_MIN_EXP, -125)
83 TEST_VALUE_EQUAL_LITERAL( "CL_FLT_RADIX", CL_FLT_RADIX, 2)
84 TEST_VALUE_EQUAL_LITERAL( "CL_FLT_MAX", CL_FLT_MAX, MAKE_HEX_FLOAT( 0x1.fffffep127f, 0x1fffffeL, 103))
85 TEST_VALUE_EQUAL_LITERAL( "CL_FLT_MIN", CL_FLT_MIN, MAKE_HEX_FLOAT(0x1.0p-126f, 0x1L, -126))
86 TEST_VALUE_EQUAL_LITERAL( "CL_FLT_EPSILON", CL_FLT_EPSILON, MAKE_HEX_FLOAT(0x1.0p-23f, 0x1L, -23))
87
88 TEST_VALUE_EQUAL_LITERAL( "CL_DBL_DIG", CL_DBL_DIG, 15)
89 TEST_VALUE_EQUAL_LITERAL( "CL_DBL_MANT_DIG", CL_DBL_MANT_DIG, 53)
90 TEST_VALUE_EQUAL_LITERAL( "CL_DBL_MAX_10_EXP", CL_DBL_MAX_10_EXP, +308)
91 TEST_VALUE_EQUAL_LITERAL( "CL_DBL_MAX_EXP", CL_DBL_MAX_EXP, +1024)
92 TEST_VALUE_EQUAL_LITERAL( "CL_DBL_MIN_10_EXP", CL_DBL_MIN_10_EXP, -307)
93 TEST_VALUE_EQUAL_LITERAL( "CL_DBL_MIN_EXP", CL_DBL_MIN_EXP, -1021)
94 TEST_VALUE_EQUAL_LITERAL( "CL_DBL_RADIX", CL_DBL_RADIX, 2)
95 TEST_VALUE_EQUAL( "CL_DBL_MAX", CL_DBL_MAX, MAKE_HEX_DOUBLE(0x1.fffffffffffffp1023, 0x1fffffffffffffLL, 971))
96 TEST_VALUE_EQUAL( "CL_DBL_MIN", CL_DBL_MIN, MAKE_HEX_DOUBLE(0x1.0p-1022, 0x1LL, -1022))
97 TEST_VALUE_EQUAL( "CL_DBL_EPSILON", CL_DBL_EPSILON, MAKE_HEX_DOUBLE(0x1.0p-52, 0x1LL, -52))
98
99 TEST_VALUE_EQUAL( "CL_M_E", CL_M_E, MAKE_HEX_DOUBLE(0x1.5bf0a8b145769p+1, 0x15bf0a8b145769LL, -51) );
100 TEST_VALUE_EQUAL( "CL_M_LOG2E", CL_M_LOG2E, MAKE_HEX_DOUBLE(0x1.71547652b82fep+0, 0x171547652b82feLL, -52) );
101 TEST_VALUE_EQUAL( "CL_M_LOG10E", CL_M_LOG10E, MAKE_HEX_DOUBLE(0x1.bcb7b1526e50ep-2, 0x1bcb7b1526e50eLL, -54) );
102 TEST_VALUE_EQUAL( "CL_M_LN2", CL_M_LN2, MAKE_HEX_DOUBLE(0x1.62e42fefa39efp-1, 0x162e42fefa39efLL, -53) );
103 TEST_VALUE_EQUAL( "CL_M_LN10", CL_M_LN10, MAKE_HEX_DOUBLE(0x1.26bb1bbb55516p+1, 0x126bb1bbb55516LL, -51) );
104 TEST_VALUE_EQUAL( "CL_M_PI", CL_M_PI, MAKE_HEX_DOUBLE(0x1.921fb54442d18p+1, 0x1921fb54442d18LL, -51) );
105 TEST_VALUE_EQUAL( "CL_M_PI_2", CL_M_PI_2, MAKE_HEX_DOUBLE(0x1.921fb54442d18p+0, 0x1921fb54442d18LL, -52) );
106 TEST_VALUE_EQUAL( "CL_M_PI_4", CL_M_PI_4, MAKE_HEX_DOUBLE(0x1.921fb54442d18p-1, 0x1921fb54442d18LL, -53) );
107 TEST_VALUE_EQUAL( "CL_M_1_PI", CL_M_1_PI, MAKE_HEX_DOUBLE(0x1.45f306dc9c883p-2, 0x145f306dc9c883LL, -54) );
108 TEST_VALUE_EQUAL( "CL_M_2_PI", CL_M_2_PI, MAKE_HEX_DOUBLE(0x1.45f306dc9c883p-1, 0x145f306dc9c883LL, -53) );
109 TEST_VALUE_EQUAL( "CL_M_2_SQRTPI", CL_M_2_SQRTPI, MAKE_HEX_DOUBLE(0x1.20dd750429b6dp+0, 0x120dd750429b6dLL, -52) );
110 TEST_VALUE_EQUAL( "CL_M_SQRT2", CL_M_SQRT2, MAKE_HEX_DOUBLE(0x1.6a09e667f3bcdp+0, 0x16a09e667f3bcdLL, -52) );
111 TEST_VALUE_EQUAL( "CL_M_SQRT1_2", CL_M_SQRT1_2, MAKE_HEX_DOUBLE(0x1.6a09e667f3bcdp-1, 0x16a09e667f3bcdLL, -53) );
112
113 TEST_VALUE_EQUAL( "CL_M_E_F", CL_M_E_F, MAKE_HEX_FLOAT(0x1.5bf0a8p+1f, 0x15bf0a8L, -23));
114 TEST_VALUE_EQUAL( "CL_M_LOG2E_F", CL_M_LOG2E_F, MAKE_HEX_FLOAT(0x1.715476p+0f, 0x1715476L, -24));
115 TEST_VALUE_EQUAL( "CL_M_LOG10E_F", CL_M_LOG10E_F, MAKE_HEX_FLOAT(0x1.bcb7b2p-2f, 0x1bcb7b2L, -26));
116 TEST_VALUE_EQUAL( "CL_M_LN2_F", CL_M_LN2_F, MAKE_HEX_FLOAT(0x1.62e43p-1f, 0x162e43L, -21) );
117 TEST_VALUE_EQUAL( "CL_M_LN10_F", CL_M_LN10_F, MAKE_HEX_FLOAT(0x1.26bb1cp+1f, 0x126bb1cL, -23));
118 TEST_VALUE_EQUAL( "CL_M_PI_F", CL_M_PI_F, MAKE_HEX_FLOAT(0x1.921fb6p+1f, 0x1921fb6L, -23));
119 TEST_VALUE_EQUAL( "CL_M_PI_2_F", CL_M_PI_2_F, MAKE_HEX_FLOAT(0x1.921fb6p+0f, 0x1921fb6L, -24));
120 TEST_VALUE_EQUAL( "CL_M_PI_4_F", CL_M_PI_4_F, MAKE_HEX_FLOAT(0x1.921fb6p-1f, 0x1921fb6L, -25));
121 TEST_VALUE_EQUAL( "CL_M_1_PI_F", CL_M_1_PI_F, MAKE_HEX_FLOAT(0x1.45f306p-2f, 0x145f306L, -26));
122 TEST_VALUE_EQUAL( "CL_M_2_PI_F", CL_M_2_PI_F, MAKE_HEX_FLOAT(0x1.45f306p-1f, 0x145f306L, -25));
123 TEST_VALUE_EQUAL( "CL_M_2_SQRTPI_F", CL_M_2_SQRTPI_F,MAKE_HEX_FLOAT(0x1.20dd76p+0f, 0x120dd76L, -24));
124 TEST_VALUE_EQUAL( "CL_M_SQRT2_F", CL_M_SQRT2_F, MAKE_HEX_FLOAT(0x1.6a09e6p+0f, 0x16a09e6L, -24));
125 TEST_VALUE_EQUAL( "CL_M_SQRT1_2_F", CL_M_SQRT1_2_F, MAKE_HEX_FLOAT(0x1.6a09e6p-1f, 0x16a09e6L, -25));
126
127 return errors;
128 }
129
130
131 const char *kernel_int_float[] = {
132 "__kernel void test( __global float *float_out, __global int *int_out, __global uint *uint_out) \n"
133 "{\n"
134 " int_out[0] = CHAR_BIT;\n"
135 " int_out[1] = SCHAR_MAX;\n"
136 " int_out[2] = SCHAR_MIN;\n"
137 " int_out[3] = CHAR_MAX;\n"
138 " int_out[4] = CHAR_MIN;\n"
139 " int_out[5] = UCHAR_MAX;\n"
140 " int_out[6] = SHRT_MAX;\n"
141 " int_out[7] = SHRT_MIN;\n"
142 " int_out[8] = USHRT_MAX;\n"
143 " int_out[9] = INT_MAX;\n"
144 " int_out[10] = INT_MIN;\n"
145 " uint_out[0] = UINT_MAX;\n"
146
147 " int_out[11] = FLT_DIG;\n"
148 " int_out[12] = FLT_MANT_DIG;\n"
149 " int_out[13] = FLT_MAX_10_EXP;\n"
150 " int_out[14] = FLT_MAX_EXP;\n"
151 " int_out[15] = FLT_MIN_10_EXP;\n"
152 " int_out[16] = FLT_MIN_EXP;\n"
153 " int_out[17] = FLT_RADIX;\n"
154 "#ifdef __IMAGE_SUPPORT__\n"
155 " int_out[18] = __IMAGE_SUPPORT__;\n"
156 "#else\n"
157 " int_out[18] = 0xf00baa;\n"
158 "#endif\n"
159 " float_out[0] = FLT_MAX;\n"
160 " float_out[1] = FLT_MIN;\n"
161 " float_out[2] = FLT_EPSILON;\n"
162 " float_out[3] = M_E_F;\n"
163 " float_out[4] = M_LOG2E_F;\n"
164 " float_out[5] = M_LOG10E_F;\n"
165 " float_out[6] = M_LN2_F;\n"
166 " float_out[7] = M_LN10_F;\n"
167 " float_out[8] = M_PI_F;\n"
168 " float_out[9] = M_PI_2_F;\n"
169 " float_out[10] = M_PI_4_F;\n"
170 " float_out[11] = M_1_PI_F;\n"
171 " float_out[12] = M_2_PI_F;\n"
172 " float_out[13] = M_2_SQRTPI_F;\n"
173 " float_out[14] = M_SQRT2_F;\n"
174 " float_out[15] = M_SQRT1_2_F;\n"
175 "}\n"
176 };
177
178 const char *kernel_long[] = {
179 "__kernel void test(__global long *long_out, __global ulong *ulong_out) \n"
180 "{\n"
181 " long_out[0] = LONG_MAX;\n"
182 " long_out[1] = LONG_MIN;\n"
183 " ulong_out[0] = ULONG_MAX;\n"
184 "}\n"
185 };
186
187 const char *kernel_double[] = {
188 "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n"
189 "__kernel void test( __global double *double_out, __global long *long_out ) \n "
190 "{\n"
191 " long_out[0] = DBL_DIG;\n"
192 " long_out[1] = DBL_MANT_DIG;\n"
193 " long_out[2] = DBL_MAX_10_EXP;\n"
194 " long_out[3] = DBL_MAX_EXP;\n"
195 " long_out[4] = DBL_MIN_10_EXP;\n"
196 " long_out[5] = DBL_MIN_EXP;\n"
197 " long_out[6] = DBL_RADIX;\n"
198 " double_out[0] = DBL_MAX;\n"
199 " double_out[1] = DBL_MIN;\n"
200 " double_out[2] = DBL_EPSILON;\n"
201 " double_out[3] = M_E;\n"
202 " double_out[4] = M_LOG2E;\n"
203 " double_out[5] = M_LOG10E;\n"
204 " double_out[6] = M_LN2;\n"
205 " double_out[7] = M_LN10;\n"
206 " double_out[8] = M_PI;\n"
207 " double_out[9] = M_PI_2;\n"
208 " double_out[10] = M_PI_4;\n"
209 " double_out[11] = M_1_PI;\n"
210 " double_out[12] = M_2_PI;\n"
211 " double_out[13] = M_2_SQRTPI;\n"
212 " double_out[14] = M_SQRT2;\n"
213 " double_out[15] = M_SQRT1_2;\n"
214 "}\n"
215 };
216
217
test_kernel_numeric_constants(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)218 int test_kernel_numeric_constants(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
219 {
220 int error, errors = 0;
221 // clProgramWrapper program;
222 // clKernelWrapper kernel;
223 // clMemWrapper streams[3];
224 cl_program program;
225 cl_kernel kernel;
226 cl_mem streams[3];
227
228 size_t threads[] = {1,1,1};
229 cl_float float_out[16];
230 cl_int int_out[19];
231 cl_uint uint_out[1];
232 cl_long long_out[7];
233 cl_ulong ulong_out[1];
234 cl_double double_out[16];
235
236 /** INTs and FLOATs **/
237
238 // Create the kernel
239 if( create_single_kernel_helper( context, &program, &kernel, 1, kernel_int_float, "test" ) != 0 )
240 {
241 return -1;
242 }
243
244 /* Create some I/O streams */
245 streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float_out),
246 NULL, &error);
247 test_error( error, "Creating test array failed" );
248 streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(int_out),
249 NULL, &error);
250 test_error( error, "Creating test array failed" );
251 streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(uint_out),
252 NULL, &error);
253 test_error( error, "Creating test array failed" );
254
255 error = clSetKernelArg(kernel, 1, sizeof( streams[1] ), &streams[1]);
256 test_error( error, "Unable to set indexed kernel arguments" );
257 error = clSetKernelArg(kernel, 0, sizeof( streams[0] ), &streams[0]);
258 test_error( error, "Unable to set indexed kernel arguments" );
259 error = clSetKernelArg(kernel, 2, sizeof( streams[2] ), &streams[2]);
260 test_error( error, "Unable to set indexed kernel arguments" );
261
262 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL );
263 test_error( error, "Kernel execution failed" );
264
265 error = clEnqueueReadBuffer( queue, streams[0], CL_TRUE, 0, sizeof(float_out), (void*)float_out, 0, NULL, NULL );
266 test_error( error, "Unable to get result data" );
267 error = clEnqueueReadBuffer( queue, streams[1], CL_TRUE, 0, sizeof(int_out), (void*)int_out, 0, NULL, NULL );
268 test_error( error, "Unable to get result data" );
269 error = clEnqueueReadBuffer( queue, streams[2], CL_TRUE, 0, sizeof(uint_out), (void*)uint_out, 0, NULL, NULL );
270 test_error( error, "Unable to get result data" );
271
272 TEST_VALUE_EQUAL_LITERAL( "CHAR_BIT", int_out[0], 8)
273 TEST_VALUE_EQUAL_LITERAL( "SCHAR_MAX", int_out[1], 127)
274 TEST_VALUE_EQUAL_LITERAL( "SCHAR_MIN", int_out[2], (-127-1))
275 TEST_VALUE_EQUAL_LITERAL( "CHAR_MAX", int_out[3], CL_SCHAR_MAX)
276 TEST_VALUE_EQUAL_LITERAL( "CHAR_MIN", int_out[4], CL_SCHAR_MIN)
277 TEST_VALUE_EQUAL_LITERAL( "UCHAR_MAX", int_out[5], 255)
278 TEST_VALUE_EQUAL_LITERAL( "SHRT_MAX", int_out[6], 32767)
279 TEST_VALUE_EQUAL_LITERAL( "SHRT_MIN",int_out[7], (-32767-1))
280 TEST_VALUE_EQUAL_LITERAL( "USHRT_MAX", int_out[8], 65535)
281 TEST_VALUE_EQUAL_LITERAL( "INT_MAX", int_out[9], 2147483647)
282 TEST_VALUE_EQUAL_LITERAL( "INT_MIN", int_out[10], (-2147483647-1))
283 TEST_VALUE_EQUAL_LITERAL( "UINT_MAX", uint_out[0], 0xffffffffU)
284
285 TEST_VALUE_EQUAL_LITERAL( "FLT_DIG", int_out[11], 6)
286 TEST_VALUE_EQUAL_LITERAL( "FLT_MANT_DIG", int_out[12], 24)
287 TEST_VALUE_EQUAL_LITERAL( "FLT_MAX_10_EXP", int_out[13], +38)
288 TEST_VALUE_EQUAL_LITERAL( "FLT_MAX_EXP", int_out[14], +128)
289 TEST_VALUE_EQUAL_LITERAL( "FLT_MIN_10_EXP", int_out[15], -37)
290 TEST_VALUE_EQUAL_LITERAL( "FLT_MIN_EXP", int_out[16], -125)
291 TEST_VALUE_EQUAL_LITERAL( "FLT_RADIX", int_out[17], 2)
292 TEST_VALUE_EQUAL( "FLT_MAX", float_out[0], MAKE_HEX_FLOAT(0x1.fffffep127f, 0x1fffffeL, 103))
293 TEST_VALUE_EQUAL( "FLT_MIN", float_out[1], MAKE_HEX_FLOAT(0x1.0p-126f, 0x1L, -126))
294 TEST_VALUE_EQUAL( "FLT_EPSILON", float_out[2], MAKE_HEX_FLOAT(0x1.0p-23f, 0x1L, -23))
295 TEST_VALUE_EQUAL( "M_E_F", float_out[3], CL_M_E_F )
296 TEST_VALUE_EQUAL( "M_LOG2E_F", float_out[4], CL_M_LOG2E_F )
297 TEST_VALUE_EQUAL( "M_LOG10E_F", float_out[5], CL_M_LOG10E_F )
298 TEST_VALUE_EQUAL( "M_LN2_F", float_out[6], CL_M_LN2_F )
299 TEST_VALUE_EQUAL( "M_LN10_F", float_out[7], CL_M_LN10_F )
300 TEST_VALUE_EQUAL( "M_PI_F", float_out[8], CL_M_PI_F )
301 TEST_VALUE_EQUAL( "M_PI_2_F", float_out[9], CL_M_PI_2_F )
302 TEST_VALUE_EQUAL( "M_PI_4_F", float_out[10], CL_M_PI_4_F )
303 TEST_VALUE_EQUAL( "M_1_PI_F", float_out[11], CL_M_1_PI_F )
304 TEST_VALUE_EQUAL( "M_2_PI_F", float_out[12], CL_M_2_PI_F )
305 TEST_VALUE_EQUAL( "M_2_SQRTPI_F", float_out[13], CL_M_2_SQRTPI_F )
306 TEST_VALUE_EQUAL( "M_SQRT2_F", float_out[14], CL_M_SQRT2_F )
307 TEST_VALUE_EQUAL( "M_SQRT1_2_F", float_out[15], CL_M_SQRT1_2_F )
308
309 // We need to check these values against what we know is supported on the device
310 if( checkForImageSupport( deviceID ) == 0 )
311 { // has images
312 // If images are supported, the constant should have been defined to the value 1
313 if( int_out[18] == 0xf00baa )
314 {
315 log_error( "FAILURE: __IMAGE_SUPPORT__ undefined even though images are supported\n" );
316 return -1;
317 }
318 else if( int_out[18] != 1 )
319 {
320 log_error( "FAILURE: __IMAGE_SUPPORT__ defined, but to the wrong value (defined as %d, spec states it should be 1)\n", int_out[18] );
321 return -1;
322 }
323 }
324 else
325 { // no images
326 // If images aren't supported, the constant should be undefined
327 if( int_out[18] != 0xf00baa )
328 {
329 log_error( "FAILURE: __IMAGE_SUPPORT__ defined to value %d even though images aren't supported", int_out[18] );
330 return -1;
331 }
332 }
333 log_info( "\t__IMAGE_SUPPORT__: %d\n", int_out[18]);
334
335 clReleaseMemObject(streams[0]); streams[0] = NULL;
336 clReleaseMemObject(streams[1]); streams[1] = NULL;
337 clReleaseMemObject(streams[2]); streams[2] = NULL;
338 clReleaseKernel(kernel); kernel = NULL;
339 clReleaseProgram(program); program = NULL;
340
341 /** LONGs **/
342
343 if(!gHasLong) {
344 log_info("Longs not supported; skipping long tests.\n");
345 }
346 else
347 {
348 // Create the kernel
349 if( create_single_kernel_helper( context, &program, &kernel, 1, kernel_long, "test" ) != 0 )
350 {
351 return -1;
352 }
353
354 streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
355 sizeof(long_out), NULL, &error);
356 test_error( error, "Creating test array failed" );
357 streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
358 sizeof(ulong_out), NULL, &error);
359 test_error( error, "Creating test array failed" );
360
361 error = clSetKernelArg(kernel, 1, sizeof( streams[1] ), &streams[1]);
362 test_error( error, "Unable to set indexed kernel arguments" );
363 error = clSetKernelArg(kernel, 0, sizeof( streams[0] ), &streams[0]);
364 test_error( error, "Unable to set indexed kernel arguments" );
365
366 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL );
367 test_error( error, "Kernel execution failed" );
368
369 error = clEnqueueReadBuffer( queue, streams[0], CL_TRUE, 0, sizeof(long_out), &long_out, 0, NULL, NULL );
370 test_error( error, "Unable to get result data" );
371 error = clEnqueueReadBuffer( queue, streams[1], CL_TRUE, 0, sizeof(ulong_out), &ulong_out, 0, NULL, NULL );
372 test_error( error, "Unable to get result data" );
373
374 TEST_VALUE_EQUAL_LITERAL( "LONG_MAX", long_out[0], ((cl_long) 0x7FFFFFFFFFFFFFFFLL))
375 TEST_VALUE_EQUAL_LITERAL( "LONG_MIN", long_out[1], ((cl_long) -0x7FFFFFFFFFFFFFFFLL - 1LL))
376 TEST_VALUE_EQUAL_LITERAL( "ULONG_MAX", ulong_out[0], ((cl_ulong) 0xFFFFFFFFFFFFFFFFULL))
377
378 clReleaseMemObject(streams[0]); streams[0] = NULL;
379 clReleaseMemObject(streams[1]); streams[1] = NULL;
380 clReleaseKernel(kernel); kernel = NULL;
381 clReleaseProgram(program); program = NULL;
382 }
383
384 /** DOUBLEs **/
385
386 if(!is_extension_available(deviceID, "cl_khr_fp64")) {
387 log_info("Extension cl_khr_fp64 not supported; skipping double tests.\n");
388 }
389 else
390 {
391 // Create the kernel
392 if( create_single_kernel_helper( context, &program, &kernel, 1, kernel_double, "test" ) != 0 )
393 {
394 return -1;
395 }
396
397 streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
398 sizeof(double_out), NULL, &error);
399 test_error( error, "Creating test array failed" );
400 streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
401 sizeof(long_out), NULL, &error);
402 test_error( error, "Creating test array failed" );
403
404 error = clSetKernelArg(kernel, 1, sizeof( streams[1] ), &streams[1]);
405 test_error( error, "Unable to set indexed kernel arguments" );
406 error = clSetKernelArg(kernel, 0, sizeof( streams[0] ), &streams[0]);
407 test_error( error, "Unable to set indexed kernel arguments" );
408
409 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL );
410 test_error( error, "Kernel execution failed" );
411
412 error = clEnqueueReadBuffer( queue, streams[0], CL_TRUE, 0, sizeof(double_out), &double_out, 0, NULL, NULL );
413 test_error( error, "Unable to get result data" );
414 error = clEnqueueReadBuffer( queue, streams[1], CL_TRUE, 0, sizeof(long_out), &long_out, 0, NULL, NULL );
415 test_error( error, "Unable to get result data" );
416
417 TEST_VALUE_EQUAL_LITERAL( "DBL_DIG", long_out[0], 15)
418 TEST_VALUE_EQUAL_LITERAL( "DBL_MANT_DIG", long_out[1], 53)
419 TEST_VALUE_EQUAL_LITERAL( "DBL_MAX_10_EXP", long_out[2], +308)
420 TEST_VALUE_EQUAL_LITERAL( "DBL_MAX_EXP", long_out[3], +1024)
421 TEST_VALUE_EQUAL_LITERAL( "DBL_MIN_10_EXP", long_out[4], -307)
422 TEST_VALUE_EQUAL_LITERAL( "DBL_MIN_EXP", long_out[5], -1021)
423 TEST_VALUE_EQUAL_LITERAL( "DBL_RADIX", long_out[6], 2)
424 TEST_VALUE_EQUAL( "DBL_MAX", double_out[0], MAKE_HEX_DOUBLE(0x1.fffffffffffffp1023, 0x1fffffffffffffLL, 971))
425 TEST_VALUE_EQUAL( "DBL_MIN", double_out[1], MAKE_HEX_DOUBLE(0x1.0p-1022, 0x1LL, -1022))
426 TEST_VALUE_EQUAL( "DBL_EPSILON", double_out[2], MAKE_HEX_DOUBLE(0x1.0p-52, 0x1LL, -52))
427 //TEST_VALUE_EQUAL( "M_E", double_out[3], CL_M_E )
428 TEST_VALUE_EQUAL( "M_LOG2E", double_out[4], CL_M_LOG2E )
429 TEST_VALUE_EQUAL( "M_LOG10E", double_out[5], CL_M_LOG10E )
430 TEST_VALUE_EQUAL( "M_LN2", double_out[6], CL_M_LN2 )
431 TEST_VALUE_EQUAL( "M_LN10", double_out[7], CL_M_LN10 )
432 TEST_VALUE_EQUAL( "M_PI", double_out[8], CL_M_PI )
433 TEST_VALUE_EQUAL( "M_PI_2", double_out[9], CL_M_PI_2 )
434 TEST_VALUE_EQUAL( "M_PI_4", double_out[10], CL_M_PI_4 )
435 TEST_VALUE_EQUAL( "M_1_PI", double_out[11], CL_M_1_PI )
436 TEST_VALUE_EQUAL( "M_2_PI", double_out[12], CL_M_2_PI )
437 TEST_VALUE_EQUAL( "M_2_SQRTPI", double_out[13], CL_M_2_SQRTPI )
438 TEST_VALUE_EQUAL( "M_SQRT2", double_out[14], CL_M_SQRT2 )
439 TEST_VALUE_EQUAL( "M_SQRT1_2", double_out[15], CL_M_SQRT1_2 )
440
441 clReleaseMemObject(streams[0]); streams[0] = NULL;
442 clReleaseMemObject(streams[1]); streams[1] = NULL;
443 clReleaseKernel(kernel); kernel = NULL;
444 clReleaseProgram(program); program = NULL;
445 }
446
447 error = clFinish(queue);
448 test_error(error, "clFinish failed");
449
450 return errors;
451 }
452
453
454 const char *kernel_constant_limits[] = {
455 "__kernel void test( __global int *intOut, __global float *floatOut ) \n"
456 "{\n"
457 " intOut[0] = isinf( MAXFLOAT ) ? 1 : 0;\n"
458 " intOut[1] = isnormal( MAXFLOAT ) ? 1 : 0;\n"
459 " intOut[2] = isnan( MAXFLOAT ) ? 1 : 0;\n"
460 " intOut[3] = sizeof( MAXFLOAT );\n"
461 " intOut[4] = ( MAXFLOAT == FLT_MAX ) ? 1 : 0;\n"
462 // " intOut[5] = ( MAXFLOAT == CL_FLT_MAX ) ? 1 : 0;\n"
463 " intOut[6] = ( MAXFLOAT == MAXFLOAT ) ? 1 : 0;\n"
464 " intOut[7] = ( MAXFLOAT == 0x1.fffffep127f ) ? 1 : 0;\n"
465 " floatOut[0] = MAXFLOAT;\n"
466 "}\n"
467 };
468
469 const char *kernel_constant_extended_limits[] = {
470 "__kernel void test( __global int *intOut, __global float *floatOut ) \n"
471 "{\n"
472 " intOut[0] = ( INFINITY == HUGE_VALF ) ? 1 : 0;\n"
473 " intOut[1] = sizeof( INFINITY );\n"
474 " intOut[2] = isinf( INFINITY ) ? 1 : 0;\n"
475 " intOut[3] = isnormal( INFINITY ) ? 1 : 0;\n"
476 " intOut[4] = isnan( INFINITY ) ? 1 : 0;\n"
477 " intOut[5] = ( INFINITY > MAXFLOAT ) ? 1 : 0;\n"
478 " intOut[6] = ( -INFINITY < -MAXFLOAT ) ? 1 : 0;\n"
479 " intOut[7] = ( ( MAXFLOAT + MAXFLOAT ) == INFINITY ) ? 1 : 0;\n"
480 " intOut[8] = ( nextafter( MAXFLOAT, INFINITY ) == INFINITY ) ? 1 : 0;\n"
481 " intOut[9] = ( nextafter( -MAXFLOAT, -INFINITY ) == -INFINITY ) ? 1 : 0;\n"
482 " intOut[10] = ( INFINITY == INFINITY ) ? 1 : 0;\n"
483 " intOut[11] = ( as_uint( INFINITY ) == 0x7f800000 ) ? 1 : 0;\n"
484 " floatOut[0] = INFINITY;\n"
485 "\n"
486 " intOut[12] = sizeof( HUGE_VALF );\n"
487 " intOut[13] = ( HUGE_VALF == INFINITY ) ? 1 : 0;\n"
488 " floatOut[1] = HUGE_VALF;\n"
489 "\n"
490 " intOut[14] = ( NAN == NAN ) ? 1 : 0;\n"
491 " intOut[15] = ( NAN != NAN ) ? 1 : 0;\n"
492 " intOut[16] = isnan( NAN ) ? 1 : 0;\n"
493 " intOut[17] = isinf( NAN ) ? 1 : 0;\n"
494 " intOut[18] = isnormal( NAN ) ? 1 : 0;\n"
495 " intOut[19] = ( ( as_uint( NAN ) & 0x7fffffff ) > 0x7f800000 ) ? 1 : 0;\n"
496 " intOut[20] = sizeof( NAN );\n"
497 " floatOut[2] = NAN;\n"
498 "\n"
499 " intOut[21] = isnan( INFINITY / INFINITY ) ? 1 : 0;\n"
500 " intOut[22] = isnan( INFINITY - INFINITY ) ? 1 : 0;\n"
501 " intOut[23] = isnan( 0.f / 0.f ) ? 1 : 0;\n"
502 " intOut[24] = isnan( INFINITY * 0.f ) ? 1 : 0;\n"
503 " intOut[25] = ( INFINITY == NAN ); \n"
504 " intOut[26] = ( -INFINITY == NAN ); \n"
505 " intOut[27] = ( INFINITY > NAN ); \n"
506 " intOut[28] = ( -INFINITY < NAN ); \n"
507 " intOut[29] = ( INFINITY != NAN ); \n"
508 " intOut[30] = ( NAN > INFINITY ); \n"
509 " intOut[31] = ( NAN < -INFINITY ); \n"
510
511 "}\n"
512 };
513
514 const char *kernel_constant_double_limits[] = {
515 "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n"
516 "__kernel void test( __global int *intOut, __global double *doubleOut ) \n"
517 "{\n"
518 " intOut[0] = sizeof( HUGE_VAL );\n"
519 " intOut[1] = ( HUGE_VAL == INFINITY ) ? 1 : 0;\n"
520 " intOut[2] = isinf( HUGE_VAL ) ? 1 : 0;\n"
521 " intOut[3] = isnormal( HUGE_VAL ) ? 1 : 0;\n"
522 " intOut[4] = isnan( HUGE_VAL ) ? 1 : 0;\n"
523 " intOut[5] = ( HUGE_VAL == HUGE_VALF ) ? 1 : 0;\n"
524 " intOut[6] = ( as_ulong( HUGE_VAL ) == 0x7ff0000000000000UL ) ? 1 : 0;\n"
525 " doubleOut[0] = HUGE_VAL;\n"
526 "}\n"
527 };
528
529 #define TEST_FLOAT_ASSERTION( a, msg, f ) if( !( a ) ) { log_error( "ERROR: Float constant failed requirement: %s (bitwise value is 0x%8.8x)\n", msg, *( (uint32_t *)&f ) ); return -1; }
530 #define TEST_DOUBLE_ASSERTION( a, msg, f ) if( !( a ) ) { log_error( "ERROR: Double constant failed requirement: %s (bitwise value is 0x%16.16llx)\n", msg, *( (uint64_t *)&f ) ); return -1; }
531
test_kernel_limit_constants(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)532 int test_kernel_limit_constants(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
533 {
534 int error;
535 size_t threads[] = {1,1,1};
536 clMemWrapper intStream, floatStream, doubleStream;
537 cl_int intOut[ 32 ];
538 cl_float floatOut[ 3 ];
539 cl_double doubleOut[ 1 ];
540
541
542 /* Create some I/O streams */
543 intStream = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(intOut), NULL,
544 &error);
545 test_error( error, "Creating test array failed" );
546 floatStream = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(floatOut),
547 NULL, &error);
548 test_error( error, "Creating test array failed" );
549
550 // Stage 1: basic limits on MAXFLOAT
551 {
552 clProgramWrapper program;
553 clKernelWrapper kernel;
554
555 if( create_single_kernel_helper( context, &program, &kernel, 1, kernel_constant_limits, "test" ) != 0 )
556 {
557 return -1;
558 }
559
560 error = clSetKernelArg( kernel, 0, sizeof( intStream ), &intStream );
561 test_error( error, "Unable to set indexed kernel arguments" );
562 error = clSetKernelArg( kernel, 1, sizeof( floatStream ), &floatStream );
563 test_error( error, "Unable to set indexed kernel arguments" );
564
565 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL );
566 test_error( error, "Kernel execution failed" );
567
568 error = clEnqueueReadBuffer( queue, intStream, CL_TRUE, 0, sizeof(intOut), intOut, 0, NULL, NULL );
569 test_error( error, "Unable to get result data" );
570 error = clEnqueueReadBuffer( queue, floatStream, CL_TRUE, 0, sizeof(floatOut), floatOut, 0, NULL, NULL );
571 test_error( error, "Unable to get result data" );
572
573 // Test MAXFLOAT properties
574 TEST_FLOAT_ASSERTION( intOut[0] == 0, "isinf( MAXFLOAT ) = false", floatOut[0] )
575 TEST_FLOAT_ASSERTION( intOut[1] == 1, "isnormal( MAXFLOAT ) = true", floatOut[0] )
576 TEST_FLOAT_ASSERTION( intOut[2] == 0, "isnan( MAXFLOAT ) = false", floatOut[0] )
577 TEST_FLOAT_ASSERTION( intOut[3] == 4, "sizeof( MAXFLOAT ) = 4", floatOut[0] )
578 TEST_FLOAT_ASSERTION( intOut[4] == 1, "MAXFLOAT = FLT_MAX", floatOut[0] )
579 TEST_FLOAT_ASSERTION( floatOut[0] == CL_FLT_MAX, "MAXFLOAT = CL_FLT_MAX", floatOut[0] )
580 TEST_FLOAT_ASSERTION( intOut[6] == 1, "MAXFLOAT = MAXFLOAT", floatOut[0] )
581 TEST_FLOAT_ASSERTION( floatOut[0] == MAKE_HEX_FLOAT( 0x1.fffffep127f, 0x1fffffeL, 103), "MAXFLOAT = 0x1.fffffep127f", floatOut[0] )
582 }
583
584 // Stage 2: INFINITY and NAN
585 char profileStr[128] = "";
586 error = clGetDeviceInfo( deviceID, CL_DEVICE_PROFILE, sizeof( profileStr ), &profileStr, NULL );
587 test_error( error, "Unable to run INFINITY/NAN tests (unable to get CL_DEVICE_PROFILE" );
588
589 bool testInfNan = true;
590 if( strcmp( profileStr, "EMBEDDED_PROFILE" ) == 0 )
591 {
592 // We test if we're not an embedded profile, OR if the inf/nan flag in the config is set
593 cl_device_fp_config single = 0;
594 error = clGetDeviceInfo( deviceID, CL_DEVICE_SINGLE_FP_CONFIG, sizeof( single ), &single, NULL );
595 test_error( error, "Unable to run INFINITY/NAN tests (unable to get FP_CONFIG bits)" );
596
597 if( ( single & CL_FP_INF_NAN ) == 0 )
598 {
599 log_info( "Skipping INFINITY and NAN tests on embedded device (INF/NAN not supported on this device)" );
600 testInfNan = false;
601 }
602 }
603
604 if( testInfNan )
605 {
606 clProgramWrapper program;
607 clKernelWrapper kernel;
608
609 if( create_single_kernel_helper( context, &program, &kernel, 1, kernel_constant_extended_limits, "test" ) != 0 )
610 {
611 return -1;
612 }
613
614 error = clSetKernelArg( kernel, 0, sizeof( intStream ), &intStream );
615 test_error( error, "Unable to set indexed kernel arguments" );
616 error = clSetKernelArg( kernel, 1, sizeof( floatStream ), &floatStream );
617 test_error( error, "Unable to set indexed kernel arguments" );
618
619 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL );
620 test_error( error, "Kernel execution failed" );
621
622 error = clEnqueueReadBuffer( queue, intStream, CL_TRUE, 0, sizeof(intOut), intOut, 0, NULL, NULL );
623 test_error( error, "Unable to get result data" );
624 error = clEnqueueReadBuffer( queue, floatStream, CL_TRUE, 0, sizeof(floatOut), floatOut, 0, NULL, NULL );
625 test_error( error, "Unable to get result data" );
626
627 TEST_FLOAT_ASSERTION( intOut[0] == 1, "INFINITY == HUGE_VALF", intOut[0] )
628 TEST_FLOAT_ASSERTION( intOut[1] == 4, "sizeof( INFINITY ) == 4", intOut[1] )
629 TEST_FLOAT_ASSERTION( intOut[2] == 1, "isinf( INFINITY ) == true", intOut[2] )
630 TEST_FLOAT_ASSERTION( intOut[3] == 0, "isnormal( INFINITY ) == false", intOut[3] )
631 TEST_FLOAT_ASSERTION( intOut[4] == 0, "isnan( INFINITY ) == false", intOut[4] )
632 TEST_FLOAT_ASSERTION( intOut[5] == 1, "INFINITY > MAXFLOAT", intOut[5] )
633 TEST_FLOAT_ASSERTION( intOut[6] == 1, "-INFINITY < -MAXFLOAT", intOut[6] )
634 TEST_FLOAT_ASSERTION( intOut[7] == 1, "( MAXFLOAT + MAXFLOAT ) == INFINITY", intOut[7] )
635 TEST_FLOAT_ASSERTION( intOut[8] == 1, "nextafter( MAXFLOAT, INFINITY ) == INFINITY", intOut[8] )
636 TEST_FLOAT_ASSERTION( intOut[9] == 1, "nextafter( -MAXFLOAT, -INFINITY ) == -INFINITY", intOut[9] )
637 TEST_FLOAT_ASSERTION( intOut[10] == 1, "INFINITY = INFINITY", intOut[10] )
638 TEST_FLOAT_ASSERTION( intOut[11] == 1, "asuint( INFINITY ) == 0x7f800000", intOut[11] )
639 TEST_FLOAT_ASSERTION( *( (uint32_t *)&floatOut[0] ) == 0x7f800000, "asuint( INFINITY ) == 0x7f800000", floatOut[0] )
640 TEST_FLOAT_ASSERTION( floatOut[1] == INFINITY, "INFINITY == INFINITY", floatOut[1] )
641
642 TEST_FLOAT_ASSERTION( intOut[12] == 4, "sizeof( HUGE_VALF ) == 4", intOut[12] )
643 TEST_FLOAT_ASSERTION( intOut[13] == 1, "HUGE_VALF == INFINITY", intOut[13] )
644 TEST_FLOAT_ASSERTION( floatOut[1] == HUGE_VALF, "HUGE_VALF == HUGE_VALF", floatOut[1] )
645
646 TEST_FLOAT_ASSERTION( intOut[14] == 0, "(NAN == NAN) = false", intOut[14] )
647 TEST_FLOAT_ASSERTION( intOut[15] == 1, "(NAN != NAN) = true", intOut[15] )
648 TEST_FLOAT_ASSERTION( intOut[16] == 1, "isnan( NAN ) = true", intOut[16] )
649 TEST_FLOAT_ASSERTION( intOut[17] == 0, "isinf( NAN ) = false", intOut[17] )
650 TEST_FLOAT_ASSERTION( intOut[18] == 0, "isnormal( NAN ) = false", intOut[18] )
651 TEST_FLOAT_ASSERTION( intOut[19] == 1, "( as_uint( NAN ) & 0x7fffffff ) > 0x7f800000", intOut[19] )
652 TEST_FLOAT_ASSERTION( intOut[20] == 4, "sizeof( NAN ) = 4", intOut[20] )
653 TEST_FLOAT_ASSERTION( ( *( (uint32_t *)&floatOut[2] ) & 0x7fffffff ) > 0x7f800000, "( as_uint( NAN ) & 0x7fffffff ) > 0x7f800000", floatOut[2] )
654
655 TEST_FLOAT_ASSERTION( intOut[ 21 ] == 1, "isnan( INFINITY / INFINITY ) = true", intOut[ 21 ] )
656 TEST_FLOAT_ASSERTION( intOut[ 22 ] == 1, "isnan( INFINITY - INFINITY ) = true", intOut[ 22 ] )
657 TEST_FLOAT_ASSERTION( intOut[ 23 ] == 1, "isnan( 0.f / 0.f ) = true", intOut[ 23 ] )
658 TEST_FLOAT_ASSERTION( intOut[ 24 ] == 1, "isnan( INFINITY * 0.f ) = true", intOut[ 24 ] )
659 TEST_FLOAT_ASSERTION( intOut[ 25 ] == 0, "( INFINITY == NAN ) = false", intOut[ 25 ] )
660 TEST_FLOAT_ASSERTION( intOut[ 26 ] == 0, "(-INFINITY == NAN ) = false", intOut[ 26 ] )
661 TEST_FLOAT_ASSERTION( intOut[ 27 ] == 0, "( INFINITY > NAN ) = false", intOut[ 27 ] )
662 TEST_FLOAT_ASSERTION( intOut[ 28 ] == 0, "(-INFINITY < NAN ) = false", intOut[ 28 ] )
663 TEST_FLOAT_ASSERTION( intOut[ 29 ] == 1, "( INFINITY != NAN ) = true", intOut[ 29 ] )
664 TEST_FLOAT_ASSERTION( intOut[ 30 ] == 0, "( NAN < INFINITY ) = false", intOut[ 30 ] )
665 TEST_FLOAT_ASSERTION( intOut[ 31 ] == 0, "( NAN > -INFINITY ) = false", intOut[ 31 ] )
666 }
667
668 // Stage 3: limits on HUGE_VAL (double)
669 if( !is_extension_available( deviceID, "cl_khr_fp64" ) )
670 log_info( "Note: Skipping double HUGE_VAL tests (doubles unsupported on device)\n" );
671 else
672 {
673 cl_device_fp_config config = 0;
674 error = clGetDeviceInfo( deviceID, CL_DEVICE_DOUBLE_FP_CONFIG, sizeof( config ), &config, NULL );
675 test_error( error, "Unable to run INFINITY/NAN tests (unable to get double FP_CONFIG bits)" );
676
677 if( ( config & CL_FP_INF_NAN ) == 0 )
678 log_info( "Skipping HUGE_VAL tests (INF/NAN not supported on this device)" );
679 else
680 {
681 clProgramWrapper program;
682 clKernelWrapper kernel;
683
684 if( create_single_kernel_helper( context, &program, &kernel, 1, kernel_constant_double_limits, "test" ) != 0 )
685 {
686 return -1;
687 }
688
689 doubleStream = clCreateBuffer(context, CL_MEM_READ_WRITE,
690 sizeof(doubleOut), NULL, &error);
691 test_error( error, "Creating test array failed" );
692
693 error = clSetKernelArg( kernel, 0, sizeof( intStream ), &intStream );
694 test_error( error, "Unable to set indexed kernel arguments" );
695 error = clSetKernelArg( kernel, 1, sizeof( doubleStream ), &doubleStream );
696 test_error( error, "Unable to set indexed kernel arguments" );
697
698 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL );
699 test_error( error, "Kernel execution failed" );
700
701 error = clEnqueueReadBuffer( queue, intStream, CL_TRUE, 0, sizeof(intOut), intOut, 0, NULL, NULL );
702 test_error( error, "Unable to get result data" );
703 error = clEnqueueReadBuffer( queue, doubleStream, CL_TRUE, 0, sizeof(doubleOut), doubleOut, 0, NULL, NULL );
704 test_error( error, "Unable to get result data" );
705
706 TEST_DOUBLE_ASSERTION( intOut[0] == 8, "sizeof( HUGE_VAL ) = 8", intOut[0] )
707 TEST_DOUBLE_ASSERTION( intOut[1] == 1, "HUGE_VAL = INFINITY", intOut[1] )
708 TEST_DOUBLE_ASSERTION( intOut[2] == 1, "isinf( HUGE_VAL ) = true", intOut[2] )
709 TEST_DOUBLE_ASSERTION( intOut[3] == 0, "isnormal( HUGE_VAL ) = false", intOut[3] )
710 TEST_DOUBLE_ASSERTION( intOut[4] == 0, "isnan( HUGE_VAL ) = false", intOut[4] )
711 TEST_DOUBLE_ASSERTION( intOut[5] == 1, "HUGE_VAL = HUGE_VAL", intOut[5] )
712 TEST_DOUBLE_ASSERTION( intOut[6] == 1, "as_ulong( HUGE_VAL ) = 0x7ff0000000000000UL", intOut[6] )
713 TEST_DOUBLE_ASSERTION( *( (uint64_t *)&doubleOut[0] ) == 0x7ff0000000000000ULL, "as_ulong( HUGE_VAL ) = 0x7ff0000000000000UL", doubleOut[0] )
714 }
715 }
716
717 return 0;
718 }
719
720
721