| // |
| // Copyright (c) 2017 The Khronos Group Inc. |
| // |
| // Licensed under the Apache License, Version 2.0 (the "License"); |
| // you may not use this file except in compliance with the License. |
| // You may obtain a copy of the License at |
| // |
| // http://www.apache.org/licenses/LICENSE-2.0 |
| // |
| // Unless required by applicable law or agreed to in writing, software |
| // distributed under the License is distributed on an "AS IS" BASIS, |
| // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. |
| // See the License for the specific language governing permissions and |
| // limitations under the License. |
| // |
| #include "procs.h" |
| |
| #define TEST_VALUE_POSITIVE( string_name, name, value ) \ |
| { \ |
| if (name < value) { \ |
| log_error("FAILED: " string_name ": " #name " < " #value "\n"); \ |
| errors++;\ |
| } else { \ |
| log_info("\t" string_name ": " #name " >= " #value "\n"); \ |
| } \ |
| } |
| |
| #define TEST_VALUE_NEGATIVE( string_name, name, value ) \ |
| { \ |
| if (name > value) { \ |
| log_error("FAILED: " string_name ": " #name " > " #value "\n"); \ |
| errors++;\ |
| } else { \ |
| log_info("\t" string_name ": " #name " <= " #value "\n"); \ |
| } \ |
| } |
| |
| #define TEST_VALUE_EQUAL_LITERAL( string_name, name, value ) \ |
| { \ |
| if (name != value) { \ |
| log_error("FAILED: " string_name ": " #name " != " #value "\n"); \ |
| errors++;\ |
| } else { \ |
| log_info("\t" string_name ": " #name " = " #value "\n"); \ |
| } \ |
| } |
| |
| #define TEST_VALUE_EQUAL( string_name, name, value ) \ |
| { \ |
| if (name != value) { \ |
| log_error("FAILED: " string_name ": " #name " != %a (%17.21g)\n", value, value); \ |
| errors++;\ |
| } else { \ |
| log_info("\t" string_name ": " #name " = %a (%17.21g)\n", value, value); \ |
| } \ |
| } |
| |
| int test_host_numeric_constants(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| int errors = 0; |
| TEST_VALUE_EQUAL_LITERAL( "CL_CHAR_BIT", CL_CHAR_BIT, 8) |
| TEST_VALUE_EQUAL_LITERAL( "CL_SCHAR_MAX", CL_SCHAR_MAX, 127) |
| TEST_VALUE_EQUAL_LITERAL( "CL_SCHAR_MIN", CL_SCHAR_MIN, (-127-1)) |
| TEST_VALUE_EQUAL_LITERAL( "CL_CHAR_MAX", CL_CHAR_MAX, CL_SCHAR_MAX) |
| TEST_VALUE_EQUAL_LITERAL( "CL_CHAR_MIN", CL_CHAR_MIN, CL_SCHAR_MIN) |
| TEST_VALUE_EQUAL_LITERAL( "CL_UCHAR_MAX", CL_UCHAR_MAX, 255) |
| TEST_VALUE_EQUAL_LITERAL( "CL_SHRT_MAX", CL_SHRT_MAX, 32767) |
| TEST_VALUE_EQUAL_LITERAL( "CL_SHRT_MIN", CL_SHRT_MIN, (-32767-1)) |
| TEST_VALUE_EQUAL_LITERAL( "CL_USHRT_MAX", CL_USHRT_MAX, 65535) |
| TEST_VALUE_EQUAL_LITERAL( "CL_INT_MAX", CL_INT_MAX, 2147483647) |
| TEST_VALUE_EQUAL_LITERAL( "CL_INT_MIN", CL_INT_MIN, (-2147483647-1)) |
| TEST_VALUE_EQUAL_LITERAL( "CL_UINT_MAX", CL_UINT_MAX, 0xffffffffU) |
| TEST_VALUE_EQUAL_LITERAL( "CL_LONG_MAX", CL_LONG_MAX, ((cl_long) 0x7FFFFFFFFFFFFFFFLL)) |
| TEST_VALUE_EQUAL_LITERAL( "CL_LONG_MIN", CL_LONG_MIN, ((cl_long) -0x7FFFFFFFFFFFFFFFLL - 1LL)) |
| TEST_VALUE_EQUAL_LITERAL( "CL_ULONG_MAX", CL_ULONG_MAX, ((cl_ulong) 0xFFFFFFFFFFFFFFFFULL)) |
| |
| TEST_VALUE_EQUAL_LITERAL( "CL_FLT_DIG", CL_FLT_DIG, 6) |
| TEST_VALUE_EQUAL_LITERAL( "CL_FLT_MANT_DIG", CL_FLT_MANT_DIG, 24) |
| TEST_VALUE_EQUAL_LITERAL( "CL_FLT_MAX_10_EXP", CL_FLT_MAX_10_EXP, +38) |
| TEST_VALUE_EQUAL_LITERAL( "CL_FLT_MAX_EXP", CL_FLT_MAX_EXP, +128) |
| TEST_VALUE_EQUAL_LITERAL( "CL_FLT_MIN_10_EXP", CL_FLT_MIN_10_EXP, -37) |
| TEST_VALUE_EQUAL_LITERAL( "CL_FLT_MIN_EXP", CL_FLT_MIN_EXP, -125) |
| TEST_VALUE_EQUAL_LITERAL( "CL_FLT_RADIX", CL_FLT_RADIX, 2) |
| TEST_VALUE_EQUAL_LITERAL( "CL_FLT_MAX", CL_FLT_MAX, MAKE_HEX_FLOAT( 0x1.fffffep127f, 0x1fffffeL, 103)) |
| TEST_VALUE_EQUAL_LITERAL( "CL_FLT_MIN", CL_FLT_MIN, MAKE_HEX_FLOAT(0x1.0p-126f, 0x1L, -126)) |
| TEST_VALUE_EQUAL_LITERAL( "CL_FLT_EPSILON", CL_FLT_EPSILON, MAKE_HEX_FLOAT(0x1.0p-23f, 0x1L, -23)) |
| |
| TEST_VALUE_EQUAL_LITERAL( "CL_DBL_DIG", CL_DBL_DIG, 15) |
| TEST_VALUE_EQUAL_LITERAL( "CL_DBL_MANT_DIG", CL_DBL_MANT_DIG, 53) |
| TEST_VALUE_EQUAL_LITERAL( "CL_DBL_MAX_10_EXP", CL_DBL_MAX_10_EXP, +308) |
| TEST_VALUE_EQUAL_LITERAL( "CL_DBL_MAX_EXP", CL_DBL_MAX_EXP, +1024) |
| TEST_VALUE_EQUAL_LITERAL( "CL_DBL_MIN_10_EXP", CL_DBL_MIN_10_EXP, -307) |
| TEST_VALUE_EQUAL_LITERAL( "CL_DBL_MIN_EXP", CL_DBL_MIN_EXP, -1021) |
| TEST_VALUE_EQUAL_LITERAL( "CL_DBL_RADIX", CL_DBL_RADIX, 2) |
| TEST_VALUE_EQUAL( "CL_DBL_MAX", CL_DBL_MAX, MAKE_HEX_DOUBLE(0x1.fffffffffffffp1023, 0x1fffffffffffffLL, 971)) |
| TEST_VALUE_EQUAL( "CL_DBL_MIN", CL_DBL_MIN, MAKE_HEX_DOUBLE(0x1.0p-1022, 0x1LL, -1022)) |
| TEST_VALUE_EQUAL( "CL_DBL_EPSILON", CL_DBL_EPSILON, MAKE_HEX_DOUBLE(0x1.0p-52, 0x1LL, -52)) |
| |
| TEST_VALUE_EQUAL( "CL_M_E", CL_M_E, MAKE_HEX_DOUBLE(0x1.5bf0a8b145769p+1, 0x15bf0a8b145769LL, -51) ); |
| TEST_VALUE_EQUAL( "CL_M_LOG2E", CL_M_LOG2E, MAKE_HEX_DOUBLE(0x1.71547652b82fep+0, 0x171547652b82feLL, -52) ); |
| TEST_VALUE_EQUAL( "CL_M_LOG10E", CL_M_LOG10E, MAKE_HEX_DOUBLE(0x1.bcb7b1526e50ep-2, 0x1bcb7b1526e50eLL, -54) ); |
| TEST_VALUE_EQUAL( "CL_M_LN2", CL_M_LN2, MAKE_HEX_DOUBLE(0x1.62e42fefa39efp-1, 0x162e42fefa39efLL, -53) ); |
| TEST_VALUE_EQUAL( "CL_M_LN10", CL_M_LN10, MAKE_HEX_DOUBLE(0x1.26bb1bbb55516p+1, 0x126bb1bbb55516LL, -51) ); |
| TEST_VALUE_EQUAL( "CL_M_PI", CL_M_PI, MAKE_HEX_DOUBLE(0x1.921fb54442d18p+1, 0x1921fb54442d18LL, -51) ); |
| TEST_VALUE_EQUAL( "CL_M_PI_2", CL_M_PI_2, MAKE_HEX_DOUBLE(0x1.921fb54442d18p+0, 0x1921fb54442d18LL, -52) ); |
| TEST_VALUE_EQUAL( "CL_M_PI_4", CL_M_PI_4, MAKE_HEX_DOUBLE(0x1.921fb54442d18p-1, 0x1921fb54442d18LL, -53) ); |
| TEST_VALUE_EQUAL( "CL_M_1_PI", CL_M_1_PI, MAKE_HEX_DOUBLE(0x1.45f306dc9c883p-2, 0x145f306dc9c883LL, -54) ); |
| TEST_VALUE_EQUAL( "CL_M_2_PI", CL_M_2_PI, MAKE_HEX_DOUBLE(0x1.45f306dc9c883p-1, 0x145f306dc9c883LL, -53) ); |
| TEST_VALUE_EQUAL( "CL_M_2_SQRTPI", CL_M_2_SQRTPI, MAKE_HEX_DOUBLE(0x1.20dd750429b6dp+0, 0x120dd750429b6dLL, -52) ); |
| TEST_VALUE_EQUAL( "CL_M_SQRT2", CL_M_SQRT2, MAKE_HEX_DOUBLE(0x1.6a09e667f3bcdp+0, 0x16a09e667f3bcdLL, -52) ); |
| TEST_VALUE_EQUAL( "CL_M_SQRT1_2", CL_M_SQRT1_2, MAKE_HEX_DOUBLE(0x1.6a09e667f3bcdp-1, 0x16a09e667f3bcdLL, -53) ); |
| |
| TEST_VALUE_EQUAL( "CL_M_E_F", CL_M_E_F, MAKE_HEX_FLOAT(0x1.5bf0a8p+1f, 0x15bf0a8L, -23)); |
| TEST_VALUE_EQUAL( "CL_M_LOG2E_F", CL_M_LOG2E_F, MAKE_HEX_FLOAT(0x1.715476p+0f, 0x1715476L, -24)); |
| TEST_VALUE_EQUAL( "CL_M_LOG10E_F", CL_M_LOG10E_F, MAKE_HEX_FLOAT(0x1.bcb7b2p-2f, 0x1bcb7b2L, -26)); |
| TEST_VALUE_EQUAL( "CL_M_LN2_F", CL_M_LN2_F, MAKE_HEX_FLOAT(0x1.62e43p-1f, 0x162e43L, -21) ); |
| TEST_VALUE_EQUAL( "CL_M_LN10_F", CL_M_LN10_F, MAKE_HEX_FLOAT(0x1.26bb1cp+1f, 0x126bb1cL, -23)); |
| TEST_VALUE_EQUAL( "CL_M_PI_F", CL_M_PI_F, MAKE_HEX_FLOAT(0x1.921fb6p+1f, 0x1921fb6L, -23)); |
| TEST_VALUE_EQUAL( "CL_M_PI_2_F", CL_M_PI_2_F, MAKE_HEX_FLOAT(0x1.921fb6p+0f, 0x1921fb6L, -24)); |
| TEST_VALUE_EQUAL( "CL_M_PI_4_F", CL_M_PI_4_F, MAKE_HEX_FLOAT(0x1.921fb6p-1f, 0x1921fb6L, -25)); |
| TEST_VALUE_EQUAL( "CL_M_1_PI_F", CL_M_1_PI_F, MAKE_HEX_FLOAT(0x1.45f306p-2f, 0x145f306L, -26)); |
| TEST_VALUE_EQUAL( "CL_M_2_PI_F", CL_M_2_PI_F, MAKE_HEX_FLOAT(0x1.45f306p-1f, 0x145f306L, -25)); |
| TEST_VALUE_EQUAL( "CL_M_2_SQRTPI_F", CL_M_2_SQRTPI_F,MAKE_HEX_FLOAT(0x1.20dd76p+0f, 0x120dd76L, -24)); |
| TEST_VALUE_EQUAL( "CL_M_SQRT2_F", CL_M_SQRT2_F, MAKE_HEX_FLOAT(0x1.6a09e6p+0f, 0x16a09e6L, -24)); |
| TEST_VALUE_EQUAL( "CL_M_SQRT1_2_F", CL_M_SQRT1_2_F, MAKE_HEX_FLOAT(0x1.6a09e6p-1f, 0x16a09e6L, -25)); |
| |
| return errors; |
| } |
| |
| |
| const char *kernel_int_float[] = { |
| "__kernel void test( __global float *float_out, __global int *int_out, __global uint *uint_out) \n" |
| "{\n" |
| " int_out[0] = CHAR_BIT;\n" |
| " int_out[1] = SCHAR_MAX;\n" |
| " int_out[2] = SCHAR_MIN;\n" |
| " int_out[3] = CHAR_MAX;\n" |
| " int_out[4] = CHAR_MIN;\n" |
| " int_out[5] = UCHAR_MAX;\n" |
| " int_out[6] = SHRT_MAX;\n" |
| " int_out[7] = SHRT_MIN;\n" |
| " int_out[8] = USHRT_MAX;\n" |
| " int_out[9] = INT_MAX;\n" |
| " int_out[10] = INT_MIN;\n" |
| " uint_out[0] = UINT_MAX;\n" |
| |
| " int_out[11] = FLT_DIG;\n" |
| " int_out[12] = FLT_MANT_DIG;\n" |
| " int_out[13] = FLT_MAX_10_EXP;\n" |
| " int_out[14] = FLT_MAX_EXP;\n" |
| " int_out[15] = FLT_MIN_10_EXP;\n" |
| " int_out[16] = FLT_MIN_EXP;\n" |
| " int_out[17] = FLT_RADIX;\n" |
| "#ifdef __IMAGE_SUPPORT__\n" |
| " int_out[18] = __IMAGE_SUPPORT__;\n" |
| "#else\n" |
| " int_out[18] = 0xf00baa;\n" |
| "#endif\n" |
| " float_out[0] = FLT_MAX;\n" |
| " float_out[1] = FLT_MIN;\n" |
| " float_out[2] = FLT_EPSILON;\n" |
| " float_out[3] = M_E_F;\n" |
| " float_out[4] = M_LOG2E_F;\n" |
| " float_out[5] = M_LOG10E_F;\n" |
| " float_out[6] = M_LN2_F;\n" |
| " float_out[7] = M_LN10_F;\n" |
| " float_out[8] = M_PI_F;\n" |
| " float_out[9] = M_PI_2_F;\n" |
| " float_out[10] = M_PI_4_F;\n" |
| " float_out[11] = M_1_PI_F;\n" |
| " float_out[12] = M_2_PI_F;\n" |
| " float_out[13] = M_2_SQRTPI_F;\n" |
| " float_out[14] = M_SQRT2_F;\n" |
| " float_out[15] = M_SQRT1_2_F;\n" |
| "}\n" |
| }; |
| |
| const char *kernel_long[] = { |
| "__kernel void test(__global long *long_out, __global ulong *ulong_out) \n" |
| "{\n" |
| " long_out[0] = LONG_MAX;\n" |
| " long_out[1] = LONG_MIN;\n" |
| " ulong_out[0] = ULONG_MAX;\n" |
| "}\n" |
| }; |
| |
| const char *kernel_double[] = { |
| "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n" |
| "__kernel void test( __global double *double_out, __global long *long_out ) \n " |
| "{\n" |
| " long_out[0] = DBL_DIG;\n" |
| " long_out[1] = DBL_MANT_DIG;\n" |
| " long_out[2] = DBL_MAX_10_EXP;\n" |
| " long_out[3] = DBL_MAX_EXP;\n" |
| " long_out[4] = DBL_MIN_10_EXP;\n" |
| " long_out[5] = DBL_MIN_EXP;\n" |
| " long_out[6] = DBL_RADIX;\n" |
| " double_out[0] = DBL_MAX;\n" |
| " double_out[1] = DBL_MIN;\n" |
| " double_out[2] = DBL_EPSILON;\n" |
| " double_out[3] = M_E;\n" |
| " double_out[4] = M_LOG2E;\n" |
| " double_out[5] = M_LOG10E;\n" |
| " double_out[6] = M_LN2;\n" |
| " double_out[7] = M_LN10;\n" |
| " double_out[8] = M_PI;\n" |
| " double_out[9] = M_PI_2;\n" |
| " double_out[10] = M_PI_4;\n" |
| " double_out[11] = M_1_PI;\n" |
| " double_out[12] = M_2_PI;\n" |
| " double_out[13] = M_2_SQRTPI;\n" |
| " double_out[14] = M_SQRT2;\n" |
| " double_out[15] = M_SQRT1_2;\n" |
| "}\n" |
| }; |
| |
| |
| int test_kernel_numeric_constants(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| int error, errors = 0; |
| // clProgramWrapper program; |
| // clKernelWrapper kernel; |
| // clMemWrapper streams[3]; |
| cl_program program; |
| cl_kernel kernel; |
| cl_mem streams[3]; |
| |
| size_t threads[] = {1,1,1}; |
| cl_float float_out[16]; |
| cl_int int_out[19]; |
| cl_uint uint_out[1]; |
| cl_long long_out[7]; |
| cl_ulong ulong_out[1]; |
| cl_double double_out[16]; |
| |
| /** INTs and FLOATs **/ |
| |
| // Create the kernel |
| if( create_single_kernel_helper( context, &program, &kernel, 1, kernel_int_float, "test" ) != 0 ) |
| { |
| return -1; |
| } |
| |
| /* Create some I/O streams */ |
| streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float_out), |
| NULL, &error); |
| test_error( error, "Creating test array failed" ); |
| streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(int_out), |
| NULL, &error); |
| test_error( error, "Creating test array failed" ); |
| streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(uint_out), |
| NULL, &error); |
| test_error( error, "Creating test array failed" ); |
| |
| error = clSetKernelArg(kernel, 1, sizeof( streams[1] ), &streams[1]); |
| test_error( error, "Unable to set indexed kernel arguments" ); |
| error = clSetKernelArg(kernel, 0, sizeof( streams[0] ), &streams[0]); |
| test_error( error, "Unable to set indexed kernel arguments" ); |
| error = clSetKernelArg(kernel, 2, sizeof( streams[2] ), &streams[2]); |
| test_error( error, "Unable to set indexed kernel arguments" ); |
| |
| error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL ); |
| test_error( error, "Kernel execution failed" ); |
| |
| error = clEnqueueReadBuffer( queue, streams[0], CL_TRUE, 0, sizeof(float_out), (void*)float_out, 0, NULL, NULL ); |
| test_error( error, "Unable to get result data" ); |
| error = clEnqueueReadBuffer( queue, streams[1], CL_TRUE, 0, sizeof(int_out), (void*)int_out, 0, NULL, NULL ); |
| test_error( error, "Unable to get result data" ); |
| error = clEnqueueReadBuffer( queue, streams[2], CL_TRUE, 0, sizeof(uint_out), (void*)uint_out, 0, NULL, NULL ); |
| test_error( error, "Unable to get result data" ); |
| |
| TEST_VALUE_EQUAL_LITERAL( "CHAR_BIT", int_out[0], 8) |
| TEST_VALUE_EQUAL_LITERAL( "SCHAR_MAX", int_out[1], 127) |
| TEST_VALUE_EQUAL_LITERAL( "SCHAR_MIN", int_out[2], (-127-1)) |
| TEST_VALUE_EQUAL_LITERAL( "CHAR_MAX", int_out[3], CL_SCHAR_MAX) |
| TEST_VALUE_EQUAL_LITERAL( "CHAR_MIN", int_out[4], CL_SCHAR_MIN) |
| TEST_VALUE_EQUAL_LITERAL( "UCHAR_MAX", int_out[5], 255) |
| TEST_VALUE_EQUAL_LITERAL( "SHRT_MAX", int_out[6], 32767) |
| TEST_VALUE_EQUAL_LITERAL( "SHRT_MIN",int_out[7], (-32767-1)) |
| TEST_VALUE_EQUAL_LITERAL( "USHRT_MAX", int_out[8], 65535) |
| TEST_VALUE_EQUAL_LITERAL( "INT_MAX", int_out[9], 2147483647) |
| TEST_VALUE_EQUAL_LITERAL( "INT_MIN", int_out[10], (-2147483647-1)) |
| TEST_VALUE_EQUAL_LITERAL( "UINT_MAX", uint_out[0], 0xffffffffU) |
| |
| TEST_VALUE_EQUAL_LITERAL( "FLT_DIG", int_out[11], 6) |
| TEST_VALUE_EQUAL_LITERAL( "FLT_MANT_DIG", int_out[12], 24) |
| TEST_VALUE_EQUAL_LITERAL( "FLT_MAX_10_EXP", int_out[13], +38) |
| TEST_VALUE_EQUAL_LITERAL( "FLT_MAX_EXP", int_out[14], +128) |
| TEST_VALUE_EQUAL_LITERAL( "FLT_MIN_10_EXP", int_out[15], -37) |
| TEST_VALUE_EQUAL_LITERAL( "FLT_MIN_EXP", int_out[16], -125) |
| TEST_VALUE_EQUAL_LITERAL( "FLT_RADIX", int_out[17], 2) |
| TEST_VALUE_EQUAL( "FLT_MAX", float_out[0], MAKE_HEX_FLOAT(0x1.fffffep127f, 0x1fffffeL, 103)) |
| TEST_VALUE_EQUAL( "FLT_MIN", float_out[1], MAKE_HEX_FLOAT(0x1.0p-126f, 0x1L, -126)) |
| TEST_VALUE_EQUAL( "FLT_EPSILON", float_out[2], MAKE_HEX_FLOAT(0x1.0p-23f, 0x1L, -23)) |
| TEST_VALUE_EQUAL( "M_E_F", float_out[3], CL_M_E_F ) |
| TEST_VALUE_EQUAL( "M_LOG2E_F", float_out[4], CL_M_LOG2E_F ) |
| TEST_VALUE_EQUAL( "M_LOG10E_F", float_out[5], CL_M_LOG10E_F ) |
| TEST_VALUE_EQUAL( "M_LN2_F", float_out[6], CL_M_LN2_F ) |
| TEST_VALUE_EQUAL( "M_LN10_F", float_out[7], CL_M_LN10_F ) |
| TEST_VALUE_EQUAL( "M_PI_F", float_out[8], CL_M_PI_F ) |
| TEST_VALUE_EQUAL( "M_PI_2_F", float_out[9], CL_M_PI_2_F ) |
| TEST_VALUE_EQUAL( "M_PI_4_F", float_out[10], CL_M_PI_4_F ) |
| TEST_VALUE_EQUAL( "M_1_PI_F", float_out[11], CL_M_1_PI_F ) |
| TEST_VALUE_EQUAL( "M_2_PI_F", float_out[12], CL_M_2_PI_F ) |
| TEST_VALUE_EQUAL( "M_2_SQRTPI_F", float_out[13], CL_M_2_SQRTPI_F ) |
| TEST_VALUE_EQUAL( "M_SQRT2_F", float_out[14], CL_M_SQRT2_F ) |
| TEST_VALUE_EQUAL( "M_SQRT1_2_F", float_out[15], CL_M_SQRT1_2_F ) |
| |
| // We need to check these values against what we know is supported on the device |
| if( checkForImageSupport( deviceID ) == 0 ) |
| { // has images |
| // If images are supported, the constant should have been defined to the value 1 |
| if( int_out[18] == 0xf00baa ) |
| { |
| log_error( "FAILURE: __IMAGE_SUPPORT__ undefined even though images are supported\n" ); |
| return -1; |
| } |
| else if( int_out[18] != 1 ) |
| { |
| log_error( "FAILURE: __IMAGE_SUPPORT__ defined, but to the wrong value (defined as %d, spec states it should be 1)\n", int_out[18] ); |
| return -1; |
| } |
| } |
| else |
| { // no images |
| // If images aren't supported, the constant should be undefined |
| if( int_out[18] != 0xf00baa ) |
| { |
| log_error( "FAILURE: __IMAGE_SUPPORT__ defined to value %d even though images aren't supported", int_out[18] ); |
| return -1; |
| } |
| } |
| log_info( "\t__IMAGE_SUPPORT__: %d\n", int_out[18]); |
| |
| clReleaseMemObject(streams[0]); streams[0] = NULL; |
| clReleaseMemObject(streams[1]); streams[1] = NULL; |
| clReleaseMemObject(streams[2]); streams[2] = NULL; |
| clReleaseKernel(kernel); kernel = NULL; |
| clReleaseProgram(program); program = NULL; |
| |
| /** LONGs **/ |
| |
| if(!gHasLong) { |
| log_info("Longs not supported; skipping long tests.\n"); |
| } |
| else |
| { |
| // Create the kernel |
| if( create_single_kernel_helper( context, &program, &kernel, 1, kernel_long, "test" ) != 0 ) |
| { |
| return -1; |
| } |
| |
| streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, |
| sizeof(long_out), NULL, &error); |
| test_error( error, "Creating test array failed" ); |
| streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, |
| sizeof(ulong_out), NULL, &error); |
| test_error( error, "Creating test array failed" ); |
| |
| error = clSetKernelArg(kernel, 1, sizeof( streams[1] ), &streams[1]); |
| test_error( error, "Unable to set indexed kernel arguments" ); |
| error = clSetKernelArg(kernel, 0, sizeof( streams[0] ), &streams[0]); |
| test_error( error, "Unable to set indexed kernel arguments" ); |
| |
| error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL ); |
| test_error( error, "Kernel execution failed" ); |
| |
| error = clEnqueueReadBuffer( queue, streams[0], CL_TRUE, 0, sizeof(long_out), &long_out, 0, NULL, NULL ); |
| test_error( error, "Unable to get result data" ); |
| error = clEnqueueReadBuffer( queue, streams[1], CL_TRUE, 0, sizeof(ulong_out), &ulong_out, 0, NULL, NULL ); |
| test_error( error, "Unable to get result data" ); |
| |
| TEST_VALUE_EQUAL_LITERAL( "LONG_MAX", long_out[0], ((cl_long) 0x7FFFFFFFFFFFFFFFLL)) |
| TEST_VALUE_EQUAL_LITERAL( "LONG_MIN", long_out[1], ((cl_long) -0x7FFFFFFFFFFFFFFFLL - 1LL)) |
| TEST_VALUE_EQUAL_LITERAL( "ULONG_MAX", ulong_out[0], ((cl_ulong) 0xFFFFFFFFFFFFFFFFULL)) |
| |
| clReleaseMemObject(streams[0]); streams[0] = NULL; |
| clReleaseMemObject(streams[1]); streams[1] = NULL; |
| clReleaseKernel(kernel); kernel = NULL; |
| clReleaseProgram(program); program = NULL; |
| } |
| |
| /** DOUBLEs **/ |
| |
| if(!is_extension_available(deviceID, "cl_khr_fp64")) { |
| log_info("Extension cl_khr_fp64 not supported; skipping double tests.\n"); |
| } |
| else |
| { |
| // Create the kernel |
| if( create_single_kernel_helper( context, &program, &kernel, 1, kernel_double, "test" ) != 0 ) |
| { |
| return -1; |
| } |
| |
| streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, |
| sizeof(double_out), NULL, &error); |
| test_error( error, "Creating test array failed" ); |
| streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, |
| sizeof(long_out), NULL, &error); |
| test_error( error, "Creating test array failed" ); |
| |
| error = clSetKernelArg(kernel, 1, sizeof( streams[1] ), &streams[1]); |
| test_error( error, "Unable to set indexed kernel arguments" ); |
| error = clSetKernelArg(kernel, 0, sizeof( streams[0] ), &streams[0]); |
| test_error( error, "Unable to set indexed kernel arguments" ); |
| |
| error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL ); |
| test_error( error, "Kernel execution failed" ); |
| |
| error = clEnqueueReadBuffer( queue, streams[0], CL_TRUE, 0, sizeof(double_out), &double_out, 0, NULL, NULL ); |
| test_error( error, "Unable to get result data" ); |
| error = clEnqueueReadBuffer( queue, streams[1], CL_TRUE, 0, sizeof(long_out), &long_out, 0, NULL, NULL ); |
| test_error( error, "Unable to get result data" ); |
| |
| TEST_VALUE_EQUAL_LITERAL( "DBL_DIG", long_out[0], 15) |
| TEST_VALUE_EQUAL_LITERAL( "DBL_MANT_DIG", long_out[1], 53) |
| TEST_VALUE_EQUAL_LITERAL( "DBL_MAX_10_EXP", long_out[2], +308) |
| TEST_VALUE_EQUAL_LITERAL( "DBL_MAX_EXP", long_out[3], +1024) |
| TEST_VALUE_EQUAL_LITERAL( "DBL_MIN_10_EXP", long_out[4], -307) |
| TEST_VALUE_EQUAL_LITERAL( "DBL_MIN_EXP", long_out[5], -1021) |
| TEST_VALUE_EQUAL_LITERAL( "DBL_RADIX", long_out[6], 2) |
| TEST_VALUE_EQUAL( "DBL_MAX", double_out[0], MAKE_HEX_DOUBLE(0x1.fffffffffffffp1023, 0x1fffffffffffffLL, 971)) |
| TEST_VALUE_EQUAL( "DBL_MIN", double_out[1], MAKE_HEX_DOUBLE(0x1.0p-1022, 0x1LL, -1022)) |
| TEST_VALUE_EQUAL( "DBL_EPSILON", double_out[2], MAKE_HEX_DOUBLE(0x1.0p-52, 0x1LL, -52)) |
| //TEST_VALUE_EQUAL( "M_E", double_out[3], CL_M_E ) |
| TEST_VALUE_EQUAL( "M_LOG2E", double_out[4], CL_M_LOG2E ) |
| TEST_VALUE_EQUAL( "M_LOG10E", double_out[5], CL_M_LOG10E ) |
| TEST_VALUE_EQUAL( "M_LN2", double_out[6], CL_M_LN2 ) |
| TEST_VALUE_EQUAL( "M_LN10", double_out[7], CL_M_LN10 ) |
| TEST_VALUE_EQUAL( "M_PI", double_out[8], CL_M_PI ) |
| TEST_VALUE_EQUAL( "M_PI_2", double_out[9], CL_M_PI_2 ) |
| TEST_VALUE_EQUAL( "M_PI_4", double_out[10], CL_M_PI_4 ) |
| TEST_VALUE_EQUAL( "M_1_PI", double_out[11], CL_M_1_PI ) |
| TEST_VALUE_EQUAL( "M_2_PI", double_out[12], CL_M_2_PI ) |
| TEST_VALUE_EQUAL( "M_2_SQRTPI", double_out[13], CL_M_2_SQRTPI ) |
| TEST_VALUE_EQUAL( "M_SQRT2", double_out[14], CL_M_SQRT2 ) |
| TEST_VALUE_EQUAL( "M_SQRT1_2", double_out[15], CL_M_SQRT1_2 ) |
| |
| clReleaseMemObject(streams[0]); streams[0] = NULL; |
| clReleaseMemObject(streams[1]); streams[1] = NULL; |
| clReleaseKernel(kernel); kernel = NULL; |
| clReleaseProgram(program); program = NULL; |
| } |
| |
| error = clFinish(queue); |
| test_error(error, "clFinish failed"); |
| |
| return errors; |
| } |
| |
| |
| const char *kernel_constant_limits[] = { |
| "__kernel void test( __global int *intOut, __global float *floatOut ) \n" |
| "{\n" |
| " intOut[0] = isinf( MAXFLOAT ) ? 1 : 0;\n" |
| " intOut[1] = isnormal( MAXFLOAT ) ? 1 : 0;\n" |
| " intOut[2] = isnan( MAXFLOAT ) ? 1 : 0;\n" |
| " intOut[3] = sizeof( MAXFLOAT );\n" |
| " intOut[4] = ( MAXFLOAT == FLT_MAX ) ? 1 : 0;\n" |
| // " intOut[5] = ( MAXFLOAT == CL_FLT_MAX ) ? 1 : 0;\n" |
| " intOut[6] = ( MAXFLOAT == MAXFLOAT ) ? 1 : 0;\n" |
| " intOut[7] = ( MAXFLOAT == 0x1.fffffep127f ) ? 1 : 0;\n" |
| " floatOut[0] = MAXFLOAT;\n" |
| "}\n" |
| }; |
| |
| const char *kernel_constant_extended_limits[] = { |
| "__kernel void test( __global int *intOut, __global float *floatOut ) \n" |
| "{\n" |
| " intOut[0] = ( INFINITY == HUGE_VALF ) ? 1 : 0;\n" |
| " intOut[1] = sizeof( INFINITY );\n" |
| " intOut[2] = isinf( INFINITY ) ? 1 : 0;\n" |
| " intOut[3] = isnormal( INFINITY ) ? 1 : 0;\n" |
| " intOut[4] = isnan( INFINITY ) ? 1 : 0;\n" |
| " intOut[5] = ( INFINITY > MAXFLOAT ) ? 1 : 0;\n" |
| " intOut[6] = ( -INFINITY < -MAXFLOAT ) ? 1 : 0;\n" |
| " intOut[7] = ( ( MAXFLOAT + MAXFLOAT ) == INFINITY ) ? 1 : 0;\n" |
| " intOut[8] = ( nextafter( MAXFLOAT, INFINITY ) == INFINITY ) ? 1 : 0;\n" |
| " intOut[9] = ( nextafter( -MAXFLOAT, -INFINITY ) == -INFINITY ) ? 1 : 0;\n" |
| " intOut[10] = ( INFINITY == INFINITY ) ? 1 : 0;\n" |
| " intOut[11] = ( as_uint( INFINITY ) == 0x7f800000 ) ? 1 : 0;\n" |
| " floatOut[0] = INFINITY;\n" |
| "\n" |
| " intOut[12] = sizeof( HUGE_VALF );\n" |
| " intOut[13] = ( HUGE_VALF == INFINITY ) ? 1 : 0;\n" |
| " floatOut[1] = HUGE_VALF;\n" |
| "\n" |
| " intOut[14] = ( NAN == NAN ) ? 1 : 0;\n" |
| " intOut[15] = ( NAN != NAN ) ? 1 : 0;\n" |
| " intOut[16] = isnan( NAN ) ? 1 : 0;\n" |
| " intOut[17] = isinf( NAN ) ? 1 : 0;\n" |
| " intOut[18] = isnormal( NAN ) ? 1 : 0;\n" |
| " intOut[19] = ( ( as_uint( NAN ) & 0x7fffffff ) > 0x7f800000 ) ? 1 : 0;\n" |
| " intOut[20] = sizeof( NAN );\n" |
| " floatOut[2] = NAN;\n" |
| "\n" |
| " intOut[21] = isnan( INFINITY / INFINITY ) ? 1 : 0;\n" |
| " intOut[22] = isnan( INFINITY - INFINITY ) ? 1 : 0;\n" |
| " intOut[23] = isnan( 0.f / 0.f ) ? 1 : 0;\n" |
| " intOut[24] = isnan( INFINITY * 0.f ) ? 1 : 0;\n" |
| " intOut[25] = ( INFINITY == NAN ); \n" |
| " intOut[26] = ( -INFINITY == NAN ); \n" |
| " intOut[27] = ( INFINITY > NAN ); \n" |
| " intOut[28] = ( -INFINITY < NAN ); \n" |
| " intOut[29] = ( INFINITY != NAN ); \n" |
| " intOut[30] = ( NAN > INFINITY ); \n" |
| " intOut[31] = ( NAN < -INFINITY ); \n" |
| |
| "}\n" |
| }; |
| |
| const char *kernel_constant_double_limits[] = { |
| "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n" |
| "__kernel void test( __global int *intOut, __global double *doubleOut ) \n" |
| "{\n" |
| " intOut[0] = sizeof( HUGE_VAL );\n" |
| " intOut[1] = ( HUGE_VAL == INFINITY ) ? 1 : 0;\n" |
| " intOut[2] = isinf( HUGE_VAL ) ? 1 : 0;\n" |
| " intOut[3] = isnormal( HUGE_VAL ) ? 1 : 0;\n" |
| " intOut[4] = isnan( HUGE_VAL ) ? 1 : 0;\n" |
| " intOut[5] = ( HUGE_VAL == HUGE_VALF ) ? 1 : 0;\n" |
| " intOut[6] = ( as_ulong( HUGE_VAL ) == 0x7ff0000000000000UL ) ? 1 : 0;\n" |
| " doubleOut[0] = HUGE_VAL;\n" |
| "}\n" |
| }; |
| |
| #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; } |
| #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; } |
| |
| int test_kernel_limit_constants(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| int error; |
| size_t threads[] = {1,1,1}; |
| clMemWrapper intStream, floatStream, doubleStream; |
| cl_int intOut[ 32 ]; |
| cl_float floatOut[ 3 ]; |
| cl_double doubleOut[ 1 ]; |
| |
| |
| /* Create some I/O streams */ |
| intStream = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(intOut), NULL, |
| &error); |
| test_error( error, "Creating test array failed" ); |
| floatStream = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(floatOut), |
| NULL, &error); |
| test_error( error, "Creating test array failed" ); |
| |
| // Stage 1: basic limits on MAXFLOAT |
| { |
| clProgramWrapper program; |
| clKernelWrapper kernel; |
| |
| if( create_single_kernel_helper( context, &program, &kernel, 1, kernel_constant_limits, "test" ) != 0 ) |
| { |
| return -1; |
| } |
| |
| error = clSetKernelArg( kernel, 0, sizeof( intStream ), &intStream ); |
| test_error( error, "Unable to set indexed kernel arguments" ); |
| error = clSetKernelArg( kernel, 1, sizeof( floatStream ), &floatStream ); |
| test_error( error, "Unable to set indexed kernel arguments" ); |
| |
| error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL ); |
| test_error( error, "Kernel execution failed" ); |
| |
| error = clEnqueueReadBuffer( queue, intStream, CL_TRUE, 0, sizeof(intOut), intOut, 0, NULL, NULL ); |
| test_error( error, "Unable to get result data" ); |
| error = clEnqueueReadBuffer( queue, floatStream, CL_TRUE, 0, sizeof(floatOut), floatOut, 0, NULL, NULL ); |
| test_error( error, "Unable to get result data" ); |
| |
| // Test MAXFLOAT properties |
| TEST_FLOAT_ASSERTION( intOut[0] == 0, "isinf( MAXFLOAT ) = false", floatOut[0] ) |
| TEST_FLOAT_ASSERTION( intOut[1] == 1, "isnormal( MAXFLOAT ) = true", floatOut[0] ) |
| TEST_FLOAT_ASSERTION( intOut[2] == 0, "isnan( MAXFLOAT ) = false", floatOut[0] ) |
| TEST_FLOAT_ASSERTION( intOut[3] == 4, "sizeof( MAXFLOAT ) = 4", floatOut[0] ) |
| TEST_FLOAT_ASSERTION( intOut[4] == 1, "MAXFLOAT = FLT_MAX", floatOut[0] ) |
| TEST_FLOAT_ASSERTION( floatOut[0] == CL_FLT_MAX, "MAXFLOAT = CL_FLT_MAX", floatOut[0] ) |
| TEST_FLOAT_ASSERTION( intOut[6] == 1, "MAXFLOAT = MAXFLOAT", floatOut[0] ) |
| TEST_FLOAT_ASSERTION( floatOut[0] == MAKE_HEX_FLOAT( 0x1.fffffep127f, 0x1fffffeL, 103), "MAXFLOAT = 0x1.fffffep127f", floatOut[0] ) |
| } |
| |
| // Stage 2: INFINITY and NAN |
| char profileStr[128] = ""; |
| error = clGetDeviceInfo( deviceID, CL_DEVICE_PROFILE, sizeof( profileStr ), &profileStr, NULL ); |
| test_error( error, "Unable to run INFINITY/NAN tests (unable to get CL_DEVICE_PROFILE" ); |
| |
| bool testInfNan = true; |
| if( strcmp( profileStr, "EMBEDDED_PROFILE" ) == 0 ) |
| { |
| // We test if we're not an embedded profile, OR if the inf/nan flag in the config is set |
| cl_device_fp_config single = 0; |
| error = clGetDeviceInfo( deviceID, CL_DEVICE_SINGLE_FP_CONFIG, sizeof( single ), &single, NULL ); |
| test_error( error, "Unable to run INFINITY/NAN tests (unable to get FP_CONFIG bits)" ); |
| |
| if( ( single & CL_FP_INF_NAN ) == 0 ) |
| { |
| log_info( "Skipping INFINITY and NAN tests on embedded device (INF/NAN not supported on this device)" ); |
| testInfNan = false; |
| } |
| } |
| |
| if( testInfNan ) |
| { |
| clProgramWrapper program; |
| clKernelWrapper kernel; |
| |
| if( create_single_kernel_helper( context, &program, &kernel, 1, kernel_constant_extended_limits, "test" ) != 0 ) |
| { |
| return -1; |
| } |
| |
| error = clSetKernelArg( kernel, 0, sizeof( intStream ), &intStream ); |
| test_error( error, "Unable to set indexed kernel arguments" ); |
| error = clSetKernelArg( kernel, 1, sizeof( floatStream ), &floatStream ); |
| test_error( error, "Unable to set indexed kernel arguments" ); |
| |
| error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL ); |
| test_error( error, "Kernel execution failed" ); |
| |
| error = clEnqueueReadBuffer( queue, intStream, CL_TRUE, 0, sizeof(intOut), intOut, 0, NULL, NULL ); |
| test_error( error, "Unable to get result data" ); |
| error = clEnqueueReadBuffer( queue, floatStream, CL_TRUE, 0, sizeof(floatOut), floatOut, 0, NULL, NULL ); |
| test_error( error, "Unable to get result data" ); |
| |
| TEST_FLOAT_ASSERTION( intOut[0] == 1, "INFINITY == HUGE_VALF", intOut[0] ) |
| TEST_FLOAT_ASSERTION( intOut[1] == 4, "sizeof( INFINITY ) == 4", intOut[1] ) |
| TEST_FLOAT_ASSERTION( intOut[2] == 1, "isinf( INFINITY ) == true", intOut[2] ) |
| TEST_FLOAT_ASSERTION( intOut[3] == 0, "isnormal( INFINITY ) == false", intOut[3] ) |
| TEST_FLOAT_ASSERTION( intOut[4] == 0, "isnan( INFINITY ) == false", intOut[4] ) |
| TEST_FLOAT_ASSERTION( intOut[5] == 1, "INFINITY > MAXFLOAT", intOut[5] ) |
| TEST_FLOAT_ASSERTION( intOut[6] == 1, "-INFINITY < -MAXFLOAT", intOut[6] ) |
| TEST_FLOAT_ASSERTION( intOut[7] == 1, "( MAXFLOAT + MAXFLOAT ) == INFINITY", intOut[7] ) |
| TEST_FLOAT_ASSERTION( intOut[8] == 1, "nextafter( MAXFLOAT, INFINITY ) == INFINITY", intOut[8] ) |
| TEST_FLOAT_ASSERTION( intOut[9] == 1, "nextafter( -MAXFLOAT, -INFINITY ) == -INFINITY", intOut[9] ) |
| TEST_FLOAT_ASSERTION( intOut[10] == 1, "INFINITY = INFINITY", intOut[10] ) |
| TEST_FLOAT_ASSERTION( intOut[11] == 1, "asuint( INFINITY ) == 0x7f800000", intOut[11] ) |
| TEST_FLOAT_ASSERTION( *( (uint32_t *)&floatOut[0] ) == 0x7f800000, "asuint( INFINITY ) == 0x7f800000", floatOut[0] ) |
| TEST_FLOAT_ASSERTION( floatOut[1] == INFINITY, "INFINITY == INFINITY", floatOut[1] ) |
| |
| TEST_FLOAT_ASSERTION( intOut[12] == 4, "sizeof( HUGE_VALF ) == 4", intOut[12] ) |
| TEST_FLOAT_ASSERTION( intOut[13] == 1, "HUGE_VALF == INFINITY", intOut[13] ) |
| TEST_FLOAT_ASSERTION( floatOut[1] == HUGE_VALF, "HUGE_VALF == HUGE_VALF", floatOut[1] ) |
| |
| TEST_FLOAT_ASSERTION( intOut[14] == 0, "(NAN == NAN) = false", intOut[14] ) |
| TEST_FLOAT_ASSERTION( intOut[15] == 1, "(NAN != NAN) = true", intOut[15] ) |
| TEST_FLOAT_ASSERTION( intOut[16] == 1, "isnan( NAN ) = true", intOut[16] ) |
| TEST_FLOAT_ASSERTION( intOut[17] == 0, "isinf( NAN ) = false", intOut[17] ) |
| TEST_FLOAT_ASSERTION( intOut[18] == 0, "isnormal( NAN ) = false", intOut[18] ) |
| TEST_FLOAT_ASSERTION( intOut[19] == 1, "( as_uint( NAN ) & 0x7fffffff ) > 0x7f800000", intOut[19] ) |
| TEST_FLOAT_ASSERTION( intOut[20] == 4, "sizeof( NAN ) = 4", intOut[20] ) |
| TEST_FLOAT_ASSERTION( ( *( (uint32_t *)&floatOut[2] ) & 0x7fffffff ) > 0x7f800000, "( as_uint( NAN ) & 0x7fffffff ) > 0x7f800000", floatOut[2] ) |
| |
| TEST_FLOAT_ASSERTION( intOut[ 21 ] == 1, "isnan( INFINITY / INFINITY ) = true", intOut[ 21 ] ) |
| TEST_FLOAT_ASSERTION( intOut[ 22 ] == 1, "isnan( INFINITY - INFINITY ) = true", intOut[ 22 ] ) |
| TEST_FLOAT_ASSERTION( intOut[ 23 ] == 1, "isnan( 0.f / 0.f ) = true", intOut[ 23 ] ) |
| TEST_FLOAT_ASSERTION( intOut[ 24 ] == 1, "isnan( INFINITY * 0.f ) = true", intOut[ 24 ] ) |
| TEST_FLOAT_ASSERTION( intOut[ 25 ] == 0, "( INFINITY == NAN ) = false", intOut[ 25 ] ) |
| TEST_FLOAT_ASSERTION( intOut[ 26 ] == 0, "(-INFINITY == NAN ) = false", intOut[ 26 ] ) |
| TEST_FLOAT_ASSERTION( intOut[ 27 ] == 0, "( INFINITY > NAN ) = false", intOut[ 27 ] ) |
| TEST_FLOAT_ASSERTION( intOut[ 28 ] == 0, "(-INFINITY < NAN ) = false", intOut[ 28 ] ) |
| TEST_FLOAT_ASSERTION( intOut[ 29 ] == 1, "( INFINITY != NAN ) = true", intOut[ 29 ] ) |
| TEST_FLOAT_ASSERTION( intOut[ 30 ] == 0, "( NAN < INFINITY ) = false", intOut[ 30 ] ) |
| TEST_FLOAT_ASSERTION( intOut[ 31 ] == 0, "( NAN > -INFINITY ) = false", intOut[ 31 ] ) |
| } |
| |
| // Stage 3: limits on HUGE_VAL (double) |
| if( !is_extension_available( deviceID, "cl_khr_fp64" ) ) |
| log_info( "Note: Skipping double HUGE_VAL tests (doubles unsupported on device)\n" ); |
| else |
| { |
| cl_device_fp_config config = 0; |
| error = clGetDeviceInfo( deviceID, CL_DEVICE_DOUBLE_FP_CONFIG, sizeof( config ), &config, NULL ); |
| test_error( error, "Unable to run INFINITY/NAN tests (unable to get double FP_CONFIG bits)" ); |
| |
| if( ( config & CL_FP_INF_NAN ) == 0 ) |
| log_info( "Skipping HUGE_VAL tests (INF/NAN not supported on this device)" ); |
| else |
| { |
| clProgramWrapper program; |
| clKernelWrapper kernel; |
| |
| if( create_single_kernel_helper( context, &program, &kernel, 1, kernel_constant_double_limits, "test" ) != 0 ) |
| { |
| return -1; |
| } |
| |
| doubleStream = clCreateBuffer(context, CL_MEM_READ_WRITE, |
| sizeof(doubleOut), NULL, &error); |
| test_error( error, "Creating test array failed" ); |
| |
| error = clSetKernelArg( kernel, 0, sizeof( intStream ), &intStream ); |
| test_error( error, "Unable to set indexed kernel arguments" ); |
| error = clSetKernelArg( kernel, 1, sizeof( doubleStream ), &doubleStream ); |
| test_error( error, "Unable to set indexed kernel arguments" ); |
| |
| error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL ); |
| test_error( error, "Kernel execution failed" ); |
| |
| error = clEnqueueReadBuffer( queue, intStream, CL_TRUE, 0, sizeof(intOut), intOut, 0, NULL, NULL ); |
| test_error( error, "Unable to get result data" ); |
| error = clEnqueueReadBuffer( queue, doubleStream, CL_TRUE, 0, sizeof(doubleOut), doubleOut, 0, NULL, NULL ); |
| test_error( error, "Unable to get result data" ); |
| |
| TEST_DOUBLE_ASSERTION( intOut[0] == 8, "sizeof( HUGE_VAL ) = 8", intOut[0] ) |
| TEST_DOUBLE_ASSERTION( intOut[1] == 1, "HUGE_VAL = INFINITY", intOut[1] ) |
| TEST_DOUBLE_ASSERTION( intOut[2] == 1, "isinf( HUGE_VAL ) = true", intOut[2] ) |
| TEST_DOUBLE_ASSERTION( intOut[3] == 0, "isnormal( HUGE_VAL ) = false", intOut[3] ) |
| TEST_DOUBLE_ASSERTION( intOut[4] == 0, "isnan( HUGE_VAL ) = false", intOut[4] ) |
| TEST_DOUBLE_ASSERTION( intOut[5] == 1, "HUGE_VAL = HUGE_VAL", intOut[5] ) |
| TEST_DOUBLE_ASSERTION( intOut[6] == 1, "as_ulong( HUGE_VAL ) = 0x7ff0000000000000UL", intOut[6] ) |
| TEST_DOUBLE_ASSERTION( *( (uint64_t *)&doubleOut[0] ) == 0x7ff0000000000000ULL, "as_ulong( HUGE_VAL ) = 0x7ff0000000000000UL", doubleOut[0] ) |
| } |
| } |
| |
| return 0; |
| } |
| |
| |
