From c8bd3c8dad877f5e935d7e93ae31b649b03e68e3 Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Thu, 27 Apr 2023 12:38:17 +0200 Subject: [PATCH 01/12] Added cl_khr_fp16 support for test_explicit_s2v from basic (issue #142, basic) --- test_conformance/basic/main.cpp | 12 +- test_conformance/basic/procs.h | 12 +- test_conformance/basic/test_explicit_s2v.cpp | 441 +++++++++---------- 3 files changed, 213 insertions(+), 252 deletions(-) diff --git a/test_conformance/basic/main.cpp b/test_conformance/basic/main.cpp index 86c3cec359..f9306f61eb 100644 --- a/test_conformance/basic/main.cpp +++ b/test_conformance/basic/main.cpp @@ -91,16 +91,8 @@ test_definition test_list[] = { ADD_TEST(image_param), ADD_TEST(image_multipass_integer_coord), ADD_TEST(image_multipass_float_coord), - ADD_TEST(explicit_s2v_char), - ADD_TEST(explicit_s2v_uchar), - ADD_TEST(explicit_s2v_short), - ADD_TEST(explicit_s2v_ushort), - ADD_TEST(explicit_s2v_int), - ADD_TEST(explicit_s2v_uint), - ADD_TEST(explicit_s2v_long), - ADD_TEST(explicit_s2v_ulong), - ADD_TEST(explicit_s2v_float), - ADD_TEST(explicit_s2v_double), + + ADD_TEST(explicit_s2v), ADD_TEST(enqueue_map_buffer), ADD_TEST(enqueue_map_image), diff --git a/test_conformance/basic/procs.h b/test_conformance/basic/procs.h index c14340de34..0c31564d34 100644 --- a/test_conformance/basic/procs.h +++ b/test_conformance/basic/procs.h @@ -89,16 +89,8 @@ extern int test_vstore_global(cl_device_id deviceID, cl_context context, cl extern int test_vstore_local(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); extern int test_vstore_private(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_explicit_s2v_char(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_explicit_s2v_uchar(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_explicit_s2v_short(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_explicit_s2v_ushort(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_explicit_s2v_int(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_explicit_s2v_uint(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_explicit_s2v_long(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_explicit_s2v_ulong(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_explicit_s2v_float(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_explicit_s2v_double(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); +extern int test_explicit_s2v(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); extern int test_enqueue_map_buffer(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); extern int test_enqueue_map_image(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); diff --git a/test_conformance/basic/test_explicit_s2v.cpp b/test_conformance/basic/test_explicit_s2v.cpp index bf38636afd..dc72be649a 100644 --- a/test_conformance/basic/test_explicit_s2v.cpp +++ b/test_conformance/basic/test_explicit_s2v.cpp @@ -1,6 +1,6 @@ // -// Copyright (c) 2017 The Khronos Group Inc. -// +// Copyright (c) 2023 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 @@ -20,131 +20,99 @@ #include #include #include - +#include #include "procs.h" #include "harness/conversions.h" #include "harness/typeWrappers.h" +// clang-format off + #define DECLARE_S2V_IDENT_KERNEL(srctype,dsttype,size) \ "__kernel void test_conversion(__global " srctype " *sourceValues, __global " dsttype #size " *destValues )\n" \ -"{\n" \ -" int tid = get_global_id(0);\n" \ -" " srctype " src = sourceValues[tid];\n" \ -"\n" \ -" destValues[tid] = (" dsttype #size ")src;\n" \ -"\n" \ +"{\n" \ +" int tid = get_global_id(0);\n" \ +" " srctype " src = sourceValues[tid];\n" \ +"\n" \ +" destValues[tid] = (" dsttype #size ")src;\n" \ +"\n" \ "}\n" -#define DECLARE_S2V_IDENT_KERNELS(srctype,dsttype) \ -{ \ -DECLARE_S2V_IDENT_KERNEL(srctype,#dsttype,2), \ -DECLARE_S2V_IDENT_KERNEL(srctype,#dsttype,4), \ -DECLARE_S2V_IDENT_KERNEL(srctype,#dsttype,8), \ -DECLARE_S2V_IDENT_KERNEL(srctype,#dsttype,16) \ -} +#define DECLARE_S2V_IDENT_KERNELS(srctype, dsttype) \ + { \ + DECLARE_S2V_IDENT_KERNEL(srctype, #dsttype, 2), \ + DECLARE_S2V_IDENT_KERNEL(srctype, #dsttype, 4), \ + DECLARE_S2V_IDENT_KERNEL(srctype, #dsttype, 8), \ + DECLARE_S2V_IDENT_KERNEL(srctype, #dsttype, 16) \ + } -#define DECLARE_EMPTY { NULL, NULL, NULL, NULL, NULL } - -/* Note: the next four arrays all must match in order and size to the ExplicitTypes enum in conversions.h!!! */ - -#define DECLARE_S2V_IDENT_KERNELS_SET(srctype) \ -{ \ -DECLARE_S2V_IDENT_KERNELS(#srctype,bool), \ - DECLARE_S2V_IDENT_KERNELS(#srctype,char), \ - DECLARE_S2V_IDENT_KERNELS(#srctype,uchar), \ - DECLARE_S2V_IDENT_KERNELS(#srctype,unsigned char), \ -DECLARE_S2V_IDENT_KERNELS(#srctype,short), \ -DECLARE_S2V_IDENT_KERNELS(#srctype,ushort), \ -DECLARE_S2V_IDENT_KERNELS(#srctype,unsigned short), \ -DECLARE_S2V_IDENT_KERNELS(#srctype,int), \ -DECLARE_S2V_IDENT_KERNELS(#srctype,uint), \ -DECLARE_S2V_IDENT_KERNELS(#srctype,unsigned int), \ -DECLARE_S2V_IDENT_KERNELS(#srctype,long), \ -DECLARE_S2V_IDENT_KERNELS(#srctype,ulong), \ -DECLARE_S2V_IDENT_KERNELS(#srctype,unsigned long), \ -DECLARE_S2V_IDENT_KERNELS(#srctype,float), \ -DECLARE_EMPTY \ -} +#define DECLARE_EMPTY \ + { \ + NULL, NULL, NULL, NULL, NULL \ + } -#define DECLARE_EMPTY_SET \ -{ \ -DECLARE_EMPTY, \ -DECLARE_EMPTY, \ -DECLARE_EMPTY, \ -DECLARE_EMPTY, \ -DECLARE_EMPTY, \ -DECLARE_EMPTY, \ -DECLARE_EMPTY, \ -DECLARE_EMPTY, \ -DECLARE_EMPTY, \ -DECLARE_EMPTY, \ -DECLARE_EMPTY, \ -DECLARE_EMPTY, \ -DECLARE_EMPTY, \ -DECLARE_EMPTY, \ -DECLARE_EMPTY \ -} +/* Note: the next four arrays all must match in order and size to the + * ExplicitTypes enum in conversions.h!!! */ + +#define DECLARE_S2V_IDENT_KERNELS_SET(srctype) \ + { \ + DECLARE_S2V_IDENT_KERNELS(#srctype, char), \ + DECLARE_S2V_IDENT_KERNELS(#srctype, uchar), \ + DECLARE_S2V_IDENT_KERNELS(#srctype, short), \ + DECLARE_S2V_IDENT_KERNELS(#srctype, ushort), \ + DECLARE_S2V_IDENT_KERNELS(#srctype, int), \ + DECLARE_S2V_IDENT_KERNELS(#srctype, uint), \ + DECLARE_S2V_IDENT_KERNELS(#srctype, long), \ + DECLARE_S2V_IDENT_KERNELS(#srctype, ulong), \ + DECLARE_S2V_IDENT_KERNELS(#srctype, float), \ + DECLARE_S2V_IDENT_KERNELS(#srctype, half), \ + DECLARE_S2V_IDENT_KERNELS(#srctype, double) \ + } + +#define DECLARE_EMPTY_SET \ + { \ + DECLARE_EMPTY, DECLARE_EMPTY, DECLARE_EMPTY, DECLARE_EMPTY, \ + DECLARE_EMPTY, DECLARE_EMPTY, DECLARE_EMPTY, DECLARE_EMPTY, \ + DECLARE_EMPTY, DECLARE_EMPTY, DECLARE_EMPTY, DECLARE_EMPTY, \ + DECLARE_EMPTY, DECLARE_EMPTY, DECLARE_EMPTY \ + } /* The overall array */ -const char * kernel_explicit_s2v_set[kNumExplicitTypes][kNumExplicitTypes][5] = { - DECLARE_S2V_IDENT_KERNELS_SET(bool), +const char * kernel_explicit_s2v_set[11][11][5] = { DECLARE_S2V_IDENT_KERNELS_SET(char), DECLARE_S2V_IDENT_KERNELS_SET(uchar), - DECLARE_S2V_IDENT_KERNELS_SET(unsigned char), DECLARE_S2V_IDENT_KERNELS_SET(short), DECLARE_S2V_IDENT_KERNELS_SET(ushort), - DECLARE_S2V_IDENT_KERNELS_SET(unsigned short), DECLARE_S2V_IDENT_KERNELS_SET(int), DECLARE_S2V_IDENT_KERNELS_SET(uint), - DECLARE_S2V_IDENT_KERNELS_SET(unsigned int), DECLARE_S2V_IDENT_KERNELS_SET(long), DECLARE_S2V_IDENT_KERNELS_SET(ulong), - DECLARE_S2V_IDENT_KERNELS_SET(unsigned long), DECLARE_S2V_IDENT_KERNELS_SET(float), - DECLARE_EMPTY_SET + DECLARE_S2V_IDENT_KERNELS_SET(half), + DECLARE_S2V_IDENT_KERNELS_SET(double) }; -int test_explicit_s2v_function(cl_device_id deviceID, cl_context context, cl_command_queue queue, const char *programSrc, - ExplicitType srcType, unsigned int count, ExplicitType destType, unsigned int vecSize, void *inputData ) +// clang-format on + +int test_explicit_s2v_function(cl_context context, cl_command_queue queue, + cl_kernel kernel, ExplicitType srcType, + unsigned int count, ExplicitType destType, + unsigned int vecSize, void *inputData) { - clProgramWrapper program; - clKernelWrapper kernel; int error; clMemWrapper streams[2]; - void *outData; - unsigned char convertedData[ 8 ]; /* Max type size is 8 bytes */ size_t threadSize[3], groupSize[3]; + unsigned char convertedData[8]; /* Max type size is 8 bytes */ unsigned int i, s; unsigned char *inPtr, *outPtr; size_t paramSize, destTypeSize; - const char* finalProgramSrc[2] = { - "", // optional pragma - programSrc - }; - - if (srcType == kDouble || destType == kDouble) { - finalProgramSrc[0] = "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n"; - } - - - if( programSrc == NULL ) - return 0; - paramSize = get_explicit_type_size( srcType ); destTypeSize = get_explicit_type_size( destType ); size_t destStride = destTypeSize * vecSize; - - outData = malloc( destStride * count ); - - if( create_single_kernel_helper( context, &program, &kernel, 2, finalProgramSrc, "test_conversion" ) ) - { - log_info( "****** %s%s *******\n", finalProgramSrc[0], finalProgramSrc[1] ); - return -1; - } + std::vector outData(destStride * count); streams[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, paramSize * count, inputData, &error); @@ -170,11 +138,13 @@ int test_explicit_s2v_function(cl_device_id deviceID, cl_context context, cl_com /* Now verify the results. Each value should have been duplicated four times, and we should be able to just do a memcpy instead of relying on the actual type of data */ - error = clEnqueueReadBuffer( queue, streams[1], CL_TRUE, 0, destStride * count, outData, 0, NULL, NULL ); + error = + clEnqueueReadBuffer(queue, streams[1], CL_TRUE, 0, destStride * count, + outData.data(), 0, NULL, NULL); test_error( error, "Unable to read output values!" ); inPtr = (unsigned char *)inputData; - outPtr = (unsigned char *)outData; + outPtr = (unsigned char *)outData.data(); for( i = 0; i < count; i++ ) { @@ -196,179 +166,186 @@ int test_explicit_s2v_function(cl_device_id deviceID, cl_context context, cl_com inPtr += paramSize; outPtr += destStride; } - - free( outData ); - return 0; } -int test_explicit_s2v_function_set(cl_device_id deviceID, cl_context context, cl_command_queue queue, ExplicitType srcType, - unsigned int count, void *inputData ) +struct TypesIterator { - unsigned int sizes[] = { 2, 4, 8, 16, 0 }; - int i, dstType, failed = 0; - - - for( dstType = kBool; dstType < kNumExplicitTypes; dstType++ ) + // in sync with ExplicitTypes, skip bools + using TypeIter = + std::tuple; + + TypesIterator(cl_device_id deviceID, cl_context context, + cl_command_queue queue) + : dstType(0), srcType(0), context(context), queue(queue) { - if( dstType == kDouble && !is_extension_available( deviceID, "cl_khr_fp64" ) ) - continue; - - if (( dstType == kLong || dstType == kULong ) && !gHasLong ) - continue; + vecTypes = { kChar, kUChar, kShort, kUShort, kInt, kUInt, + kLong, kULong, kFloat, kHalf, kDouble }; + fp16Support = is_extension_available(deviceID, "cl_khr_fp16"); + fp64Support = is_extension_available(deviceID, "cl_khr_fp64"); - for( i = 0; sizes[i] != 0; i++ ) - { - if( dstType != srcType ) - continue; - if( strchr( get_explicit_type_name( (ExplicitType)srcType ), ' ' ) != NULL || - strchr( get_explicit_type_name( (ExplicitType)dstType ), ' ' ) != NULL ) - continue; - - if( test_explicit_s2v_function( deviceID, context, queue, kernel_explicit_s2v_set[ srcType ][ dstType ][ i ], - srcType, count, (ExplicitType)dstType, sizes[ i ], inputData ) != 0 ) - { - log_error( "ERROR: Explicit cast of scalar %s to vector %s%d FAILED; skipping other %s vector tests\n", - get_explicit_type_name(srcType), get_explicit_type_name((ExplicitType)dstType), sizes[i], get_explicit_type_name((ExplicitType)dstType) ); - failed = -1; - break; - } - } + for_each_src_elem(it); } - return failed; -} - -int test_explicit_s2v_char(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - char data[128]; - RandomSeed seed(gRandomSeed); - - generate_random_data( kChar, 128, seed, data ); - - return test_explicit_s2v_function_set( deviceID, context, queue, kChar, 128, data ); -} - -int test_explicit_s2v_uchar(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - unsigned char data[128]; - RandomSeed seed(gRandomSeed); - - generate_random_data( kUChar, 128, seed, data ); + bool skip_type(ExplicitType type) + { + if ((type == kLong || type == kULong) && !gHasLong) + return true; + else if (type == kDouble && !fp64Support) + return true; + else if (type == kHalf && !fp16Support) + return true; + else if (strchr(get_explicit_type_name(type), ' ') != 0) + return true; + return false; + } - if( test_explicit_s2v_function_set( deviceID, context, queue, kUChar, 128, data ) != 0 ) - return -1; - if( test_explicit_s2v_function_set( deviceID, context, queue, kUnsignedChar, 128, data ) != 0 ) - return -1; - return 0; -} + template + void iterate_src_type(const SrcType &t) + { + bool doTest = !skip_type(vecTypes[srcType]); + if (doTest) + { + SrcType inputData[sample_count]; + RandomSeed seed(gRandomSeed); + generate_random_data(vecTypes[srcType], 128, seed, inputData); -int test_explicit_s2v_short(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - short data[128]; - RandomSeed seed(gRandomSeed); + for_each_dst_elem<0, Src, SrcType>(it, inputData); + } - generate_random_data( kShort, 128, seed, data ); + srcType++; + dstType = 0; + } - if( test_explicit_s2v_function_set( deviceID, context, queue, kShort, 128, data ) != 0 ) - return -1; - return 0; -} + // crucial to keep it in sync with ExplicitType + bool isExplicitTypeFloating(ExplicitType type) { return (type >= kFloat); } -int test_explicit_s2v_ushort(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - unsigned short data[128]; - RandomSeed seed(gRandomSeed); + template + void iterate_dst_type(const DstType &t, SrcType *inputData) + { + bool doTest = !skip_type(vecTypes[dstType]); - generate_random_data( kUShort, 128, seed, data ); + doTest = doTest + && ((isExplicitTypeFloating(vecTypes[srcType]) + && isExplicitTypeFloating(vecTypes[dstType])) + || (!isExplicitTypeFloating(vecTypes[srcType]) + && !isExplicitTypeFloating(vecTypes[dstType]))); - if( test_explicit_s2v_function_set( deviceID, context, queue, kUShort, 128, data ) != 0 ) - return -1; - if( test_explicit_s2v_function_set( deviceID, context, queue, kUnsignedShort, 128, data ) != 0 ) - return -1; - return 0; -} + if (doTest /*&& dstType != srcType*/) + { + test_explicit_s2v_function_set( + vecTypes[srcType], vecTypes[dstType], inputData); + } + dstType++; + } -int test_explicit_s2v_int(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - int data[128]; - RandomSeed seed(gRandomSeed); + template + inline typename std::enable_if::type + for_each_src_elem( + const std::tuple &) // Unused arguments are given no names. + {} - generate_random_data( kInt, 128, seed, data ); + template + inline typename std::enable_if < Out::type + for_each_src_elem(const std::tuple &t) + { + iterate_src_type(std::get(t)); + for_each_src_elem(t); + } - if( test_explicit_s2v_function_set( deviceID, context, queue, kInt, 128, data ) != 0 ) - return -1; - return 0; -} + template + inline typename std::enable_if::type + for_each_dst_elem(const std::tuple &, SrcType *) + {} -int test_explicit_s2v_uint(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - unsigned int data[128]; - RandomSeed seed(gRandomSeed); + template + inline typename std::enable_if < In::type + for_each_dst_elem(const std::tuple &t, SrcType *inputData) + { + iterate_dst_type(std::get(t), inputData); + for_each_dst_elem(t, inputData); + } - generate_random_data( kUInt, 128, seed, data ); + template + void test_explicit_s2v_function_set(ExplicitType srcT, ExplicitType dstT, + SrcType *inputData) + { + unsigned int sizes[] = { 2, 4, 8, 16, 0 }; - if( test_explicit_s2v_function_set( deviceID, context, queue, kUInt, 128, data ) != 0 ) - return -1; - if( test_explicit_s2v_function_set( deviceID, context, queue, kUnsignedInt, 128, data ) != 0 ) - return -1; - return 0; -} + for (int i = 0; sizes[i] != 0; i++) + { + clProgramWrapper program; + clKernelWrapper kernel; -int test_explicit_s2v_long(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - cl_long data[128]; - RandomSeed seed(gRandomSeed); + char pragma[256] = { 0 }; + const char *finalProgramSrc[2] = { + pragma, // optional pragma + kernel_explicit_s2v_set[srcType][dstType][i] + }; - generate_random_data( kLong, 128, seed, data ); + std::stringstream sstr; + if (srcType == kDouble || dstType == kDouble) + sstr << "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n"; - if( test_explicit_s2v_function_set( deviceID, context, queue, kLong, 128, data ) != 0 ) - return -1; - return 0; -} + if (srcType == kHalf || dstType == kHalf) + sstr << "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n"; -int test_explicit_s2v_ulong(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - cl_ulong data[128]; - RandomSeed seed(gRandomSeed); + snprintf(pragma, sizeof(pragma), "%s", sstr.str().c_str()); - generate_random_data( kULong, 128, seed, data ); + if (create_single_kernel_helper(context, &program, &kernel, 2, + finalProgramSrc, "test_conversion")) + { + log_info("****** %s%s *******\n", finalProgramSrc[0], + finalProgramSrc[1]); + throw std::runtime_error( + "create_single_kernel_helper failed\n"); + } - if( test_explicit_s2v_function_set( deviceID, context, queue, kULong, 128, data ) != 0 ) - return -1; - if( test_explicit_s2v_function_set( deviceID, context, queue, kUnsignedLong, 128, data ) != 0 ) - return -1; - return 0; -} + if (test_explicit_s2v_function(context, queue, kernel, srcT, + sample_count, dstT, sizes[i], + inputData) + != 0) + { + log_error("ERROR: Explicit cast of scalar %s to vector %s%d " + "FAILED; skipping other %s vector tests\n", + get_explicit_type_name(srcT), + get_explicit_type_name(dstT), sizes[i], + get_explicit_type_name(dstT)); + throw std::runtime_error("test_explicit_s2v_function failed\n"); + } + } + } -int test_explicit_s2v_float(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - float data[128]; - RandomSeed seed(gRandomSeed); +protected: + bool fp16Support; + bool fp64Support; - generate_random_data( kFloat, 128, seed, data ); + TypeIter it; + unsigned int dstType, srcType; + cl_context context; + cl_command_queue queue; - if( test_explicit_s2v_function_set( deviceID, context, queue, kFloat, 128, data ) != 0 ) - return -1; - return 0; -} + std::vector vecTypes; + constexpr static unsigned int sample_count = + 128; // hardcoded in original test +}; -int test_explicit_s2v_double(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) +int test_explicit_s2v(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) { - double data[128]; - RandomSeed seed(gRandomSeed); - - if( !is_extension_available( deviceID, "cl_khr_fp64" ) ) { - log_info("Extension cl_khr_fp64 not supported. Skipping test.\n"); - return 0; + try + { + TypesIterator(deviceID, context, queue); + } catch (const std::runtime_error &e) + { + log_error("%s", e.what()); + return TEST_FAIL; } - generate_random_data( kDouble, 128, seed, data ); - - if( test_explicit_s2v_function_set( deviceID, context, queue, kDouble, 128, data ) != 0 ) - return -1; - return 0; + return TEST_PASS; } - - From 6f4afbeede3b8ede5b8b02fc4d0765c4f96fc276 Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Thu, 27 Apr 2023 12:46:34 +0200 Subject: [PATCH 02/12] Cosmetic corrections --- test_conformance/basic/test_explicit_s2v.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/test_conformance/basic/test_explicit_s2v.cpp b/test_conformance/basic/test_explicit_s2v.cpp index dc72be649a..841757025c 100644 --- a/test_conformance/basic/test_explicit_s2v.cpp +++ b/test_conformance/basic/test_explicit_s2v.cpp @@ -77,9 +77,10 @@ DECLARE_EMPTY, DECLARE_EMPTY, DECLARE_EMPTY \ } +#define NUM_VEC_TYPES 11 /* The overall array */ -const char * kernel_explicit_s2v_set[11][11][5] = { +const char * kernel_explicit_s2v_set[NUM_VEC_TYPES][NUM_VEC_TYPES][5] = { DECLARE_S2V_IDENT_KERNELS_SET(char), DECLARE_S2V_IDENT_KERNELS_SET(uchar), DECLARE_S2V_IDENT_KERNELS_SET(short), @@ -171,7 +172,6 @@ int test_explicit_s2v_function(cl_context context, cl_command_queue queue, struct TypesIterator { - // in sync with ExplicitTypes, skip bools using TypeIter = std::tuple; @@ -218,7 +218,7 @@ struct TypesIterator dstType = 0; } - // crucial to keep it in sync with ExplicitType + // crucial to keep it in-sync with ExplicitType bool isExplicitTypeFloating(ExplicitType type) { return (type >= kFloat); } template Date: Thu, 27 Apr 2023 13:27:18 +0200 Subject: [PATCH 03/12] cosmetic fix --- test_conformance/basic/test_explicit_s2v.cpp | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/test_conformance/basic/test_explicit_s2v.cpp b/test_conformance/basic/test_explicit_s2v.cpp index 841757025c..2f0fc86d95 100644 --- a/test_conformance/basic/test_explicit_s2v.cpp +++ b/test_conformance/basic/test_explicit_s2v.cpp @@ -233,11 +233,9 @@ struct TypesIterator || (!isExplicitTypeFloating(vecTypes[srcType]) && !isExplicitTypeFloating(vecTypes[dstType]))); - if (doTest /*&& dstType != srcType*/) - { + if (doTest) test_explicit_s2v_function_set( vecTypes[srcType], vecTypes[dstType], inputData); - } dstType++; } From 52c8047f43c59cff29c85ca59e927571fb31a2dd Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Thu, 11 May 2023 09:47:33 +0200 Subject: [PATCH 04/12] Added correction to distinguish signed and unsigned char types for ARM architecture tests --- test_common/harness/conversions.cpp | 61 +++++++++++++++-------------- 1 file changed, 31 insertions(+), 30 deletions(-) diff --git a/test_common/harness/conversions.cpp b/test_common/harness/conversions.cpp index d52a2ac612..03c371863c 100644 --- a/test_common/harness/conversions.cpp +++ b/test_common/harness/conversions.cpp @@ -386,6 +386,7 @@ static Long sLowerLimits[kNumExplicitTypes] = { break; typedef unsigned char uchar; +typedef signed char schar; typedef unsigned short ushort; typedef unsigned int uint; typedef unsigned long ulong; @@ -395,7 +396,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, ExplicitType outType) { bool *boolPtr; - char *charPtr; + schar *scharPtr; uchar *ucharPtr; short *shortPtr; ushort *ushortPtr; @@ -449,29 +450,29 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, break; case kChar: - charPtr = (char *)inRaw; + scharPtr = (schar *)inRaw; switch (outType) { - BOOL_CASE(char) + BOOL_CASE(schar) case kChar: memcpy(outRaw, inRaw, get_explicit_type_size(inType)); break; - DOWN_CAST_CASE(char, kUChar, uchar, saturate) - SIMPLE_CAST_CASE(char, kUnsignedChar, uchar) - SIMPLE_CAST_CASE(char, kShort, short) - SIMPLE_CAST_CASE(char, kUShort, ushort) - SIMPLE_CAST_CASE(char, kUnsignedShort, ushort) - SIMPLE_CAST_CASE(char, kInt, int) - SIMPLE_CAST_CASE(char, kUInt, uint) - SIMPLE_CAST_CASE(char, kUnsignedInt, uint) - SIMPLE_CAST_CASE(char, kLong, Long) - SIMPLE_CAST_CASE(char, kULong, ULong) - SIMPLE_CAST_CASE(char, kUnsignedLong, ULong) - - TO_FLOAT_CASE(char) - TO_DOUBLE_CASE(char) + DOWN_CAST_CASE(schar, kUChar, uchar, saturate) + SIMPLE_CAST_CASE(schar, kUnsignedChar, uchar) + SIMPLE_CAST_CASE(schar, kShort, short) + SIMPLE_CAST_CASE(schar, kUShort, ushort) + SIMPLE_CAST_CASE(schar, kUnsignedShort, ushort) + SIMPLE_CAST_CASE(schar, kInt, int) + SIMPLE_CAST_CASE(schar, kUInt, uint) + SIMPLE_CAST_CASE(schar, kUnsignedInt, uint) + SIMPLE_CAST_CASE(schar, kLong, Long) + SIMPLE_CAST_CASE(schar, kULong, ULong) + SIMPLE_CAST_CASE(schar, kUnsignedLong, ULong) + + TO_FLOAT_CASE(schar) + TO_DOUBLE_CASE(schar) default: log_error("ERROR: Invalid type given to " @@ -491,7 +492,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, memcpy(outRaw, inRaw, get_explicit_type_size(inType)); break; - DOWN_CAST_CASE(uchar, kChar, char, saturate) + DOWN_CAST_CASE(uchar, kChar, schar, saturate) SIMPLE_CAST_CASE(uchar, kShort, short) SIMPLE_CAST_CASE(uchar, kUShort, ushort) SIMPLE_CAST_CASE(uchar, kUnsignedShort, ushort) @@ -523,7 +524,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, memcpy(outRaw, inRaw, get_explicit_type_size(inType)); break; - DOWN_CAST_CASE(uchar, kChar, char, saturate) + DOWN_CAST_CASE(uchar, kChar, schar, saturate) SIMPLE_CAST_CASE(uchar, kShort, short) SIMPLE_CAST_CASE(uchar, kUShort, ushort) SIMPLE_CAST_CASE(uchar, kUnsignedShort, ushort) @@ -554,7 +555,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, memcpy(outRaw, inRaw, get_explicit_type_size(inType)); break; - DOWN_CAST_CASE(short, kChar, char, saturate) + DOWN_CAST_CASE(short, kChar, schar, saturate) DOWN_CAST_CASE(short, kUChar, uchar, saturate) DOWN_CAST_CASE(short, kUnsignedChar, uchar, saturate) DOWN_CAST_CASE(short, kUShort, ushort, saturate) @@ -587,7 +588,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, memcpy(outRaw, inRaw, get_explicit_type_size(inType)); break; - DOWN_CAST_CASE(ushort, kChar, char, saturate) + DOWN_CAST_CASE(ushort, kChar, schar, saturate) DOWN_CAST_CASE(ushort, kUChar, uchar, saturate) DOWN_CAST_CASE(ushort, kUnsignedChar, uchar, saturate) DOWN_CAST_CASE(ushort, kShort, short, saturate) @@ -619,7 +620,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, memcpy(outRaw, inRaw, get_explicit_type_size(inType)); break; - DOWN_CAST_CASE(ushort, kChar, char, saturate) + DOWN_CAST_CASE(ushort, kChar, schar, saturate) DOWN_CAST_CASE(ushort, kUChar, uchar, saturate) DOWN_CAST_CASE(ushort, kUnsignedChar, uchar, saturate) DOWN_CAST_CASE(ushort, kShort, short, saturate) @@ -650,7 +651,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, memcpy(outRaw, inRaw, get_explicit_type_size(inType)); break; - DOWN_CAST_CASE(int, kChar, char, saturate) + DOWN_CAST_CASE(int, kChar, schar, saturate) DOWN_CAST_CASE(int, kUChar, uchar, saturate) DOWN_CAST_CASE(int, kUnsignedChar, uchar, saturate) DOWN_CAST_CASE(int, kShort, short, saturate) @@ -683,7 +684,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, memcpy(outRaw, inRaw, get_explicit_type_size(inType)); break; - DOWN_CAST_CASE(uint, kChar, char, saturate) + DOWN_CAST_CASE(uint, kChar, schar, saturate) DOWN_CAST_CASE(uint, kUChar, uchar, saturate) DOWN_CAST_CASE(uint, kUnsignedChar, uchar, saturate) DOWN_CAST_CASE(uint, kShort, short, saturate) @@ -715,7 +716,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, memcpy(outRaw, inRaw, get_explicit_type_size(inType)); break; - DOWN_CAST_CASE(uint, kChar, char, saturate) + DOWN_CAST_CASE(uint, kChar, schar, saturate) DOWN_CAST_CASE(uint, kUChar, uchar, saturate) DOWN_CAST_CASE(uint, kUnsignedChar, uchar, saturate) DOWN_CAST_CASE(uint, kShort, short, saturate) @@ -746,7 +747,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, memcpy(outRaw, inRaw, get_explicit_type_size(inType)); break; - DOWN_CAST_CASE(Long, kChar, char, saturate) + DOWN_CAST_CASE(Long, kChar, schar, saturate) DOWN_CAST_CASE(Long, kUChar, uchar, saturate) DOWN_CAST_CASE(Long, kUnsignedChar, uchar, saturate) DOWN_CAST_CASE(Long, kShort, short, saturate) @@ -779,7 +780,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, memcpy(outRaw, inRaw, get_explicit_type_size(inType)); break; - U_DOWN_CAST_CASE(ULong, kChar, char, saturate) + U_DOWN_CAST_CASE(ULong, kChar, schar, saturate) U_DOWN_CAST_CASE(ULong, kUChar, uchar, saturate) U_DOWN_CAST_CASE(ULong, kUnsignedChar, uchar, saturate) U_DOWN_CAST_CASE(ULong, kShort, short, saturate) @@ -811,7 +812,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, memcpy(outRaw, inRaw, get_explicit_type_size(inType)); break; - U_DOWN_CAST_CASE(ULong, kChar, char, saturate) + U_DOWN_CAST_CASE(ULong, kChar, schar, saturate) U_DOWN_CAST_CASE(ULong, kUChar, uchar, saturate) U_DOWN_CAST_CASE(ULong, kUnsignedChar, uchar, saturate) U_DOWN_CAST_CASE(ULong, kShort, short, saturate) @@ -838,7 +839,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, { BOOL_CASE(float) - FLOAT_ROUND_CASE(kChar, char, roundType, saturate) + FLOAT_ROUND_CASE(kChar, schar, roundType, saturate) FLOAT_ROUND_CASE(kUChar, uchar, roundType, saturate) FLOAT_ROUND_CASE(kUnsignedChar, uchar, roundType, saturate) FLOAT_ROUND_CASE(kShort, short, roundType, saturate) @@ -870,7 +871,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, { BOOL_CASE(double) - DOUBLE_ROUND_CASE(kChar, char, roundType, saturate) + DOUBLE_ROUND_CASE(kChar, schar, roundType, saturate) DOUBLE_ROUND_CASE(kUChar, uchar, roundType, saturate) DOUBLE_ROUND_CASE(kUnsignedChar, uchar, roundType, saturate) DOUBLE_ROUND_CASE(kShort, short, roundType, saturate) From b067494272417ce2faac63d6ec8e236988bfbfbb Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Wed, 14 Jun 2023 13:19:12 +0200 Subject: [PATCH 05/12] Added missing pieces of convertion procedure to support half --- test_common/harness/conversions.cpp | 121 ++++++++++++++++++++++++++++ 1 file changed, 121 insertions(+) diff --git a/test_common/harness/conversions.cpp b/test_common/harness/conversions.cpp index 03c371863c..300d7d4648 100644 --- a/test_common/harness/conversions.cpp +++ b/test_common/harness/conversions.cpp @@ -21,6 +21,8 @@ #include "mt19937.h" #include "compat.h" +#include + #if defined(__SSE__) || defined(_MSC_VER) #include #endif @@ -261,6 +263,11 @@ static Long sLowerLimits[kNumExplicitTypes] = { } \ break; +#define TO_HALF_CASE(inType) \ + case kHalf: \ + halfPtr = (cl_half *)outRaw; \ + *halfPtr = cl_half_from_float((float)(*inType##Ptr), CL_HALF_RTE); \ + break; #define TO_FLOAT_CASE(inType) \ case kFloat: \ floatPtr = (float *)outRaw; \ @@ -281,6 +288,59 @@ static Long sLowerLimits[kNumExplicitTypes] = { *outType##Ptr = (outType)lrintf_clamped(*floatPtr); \ break; +#define HALF_ROUND_CASE(outEnum, outType, rounding, sat) \ + case outEnum: { \ + outType##Ptr = (outType *)outRaw; \ + /* Get the tens digit */ \ + float fltEq = (Long)cl_half_to_float(*halfPtr); \ + Long wholeValue = (Long)fltEq; \ + float largeRemainder = (fltEq - (float)wholeValue) * 10.f; \ + /* What do we do based on that? */ \ + if (rounding == kRoundToEven) \ + { \ + if (wholeValue & 1LL) /*between 1 and 1.99 */ \ + wholeValue += 1LL; /* round up to even */ \ + } \ + else if (rounding == kRoundToZero) \ + { \ + /* Nothing to do, round-to-zero is what C casting does */ \ + } \ + else if (rounding == kRoundToPosInf) \ + { \ + /* Only positive numbers are wrong */ \ + if (largeRemainder != 0.f && wholeValue >= 0) wholeValue++; \ + } \ + else if (rounding == kRoundToNegInf) \ + { \ + /* Only negative numbers are off */ \ + if (largeRemainder != 0.f && wholeValue < 0) wholeValue--; \ + } \ + else \ + { /* Default is round-to-nearest */ \ + wholeValue = (Long)lrintf_clamped(fltEq); \ + } \ + /* Now apply saturation rules */ \ + if (sat) \ + { \ + if ((sLowerLimits[outEnum] < 0 \ + && wholeValue > (Long)sUpperLimits[outEnum]) \ + || (sLowerLimits[outEnum] == 0 \ + && (ULong)wholeValue > sUpperLimits[outEnum])) \ + *outType##Ptr = (outType)sUpperLimits[outEnum]; \ + else if (wholeValue < sLowerLimits[outEnum]) \ + *outType##Ptr = (outType)sLowerLimits[outEnum]; \ + else \ + *outType##Ptr = (outType)wholeValue; \ + } \ + else \ + { \ + *outType##Ptr = (outType)( \ + wholeValue \ + & (0xffffffffffffffffLL >> (64 - (sizeof(outType) * 8)))); \ + } \ + } \ + break; + #define FLOAT_ROUND_CASE(outEnum, outType, rounding, sat) \ case outEnum: { \ outType##Ptr = (outType *)outRaw; \ @@ -404,6 +464,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, uint *uintPtr; Long *LongPtr; ULong *ULongPtr; + cl_half *halfPtr; float *floatPtr; double *doublePtr; @@ -434,6 +495,11 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, get_explicit_type_size(outType)); break; + case kHalf: + halfPtr = (cl_half *)outRaw; + *halfPtr = + (*boolPtr) ? cl_half_from_float(-1.f, CL_HALF_RTE) : 0; + break; case kFloat: floatPtr = (float *)outRaw; *floatPtr = (*boolPtr) ? -1.f : 0.f; @@ -471,6 +537,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, SIMPLE_CAST_CASE(schar, kULong, ULong) SIMPLE_CAST_CASE(schar, kUnsignedLong, ULong) + TO_HALF_CASE(schar) TO_FLOAT_CASE(schar) TO_DOUBLE_CASE(schar) @@ -503,6 +570,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, SIMPLE_CAST_CASE(uchar, kULong, ULong) SIMPLE_CAST_CASE(uchar, kUnsignedLong, ULong) + TO_HALF_CASE(uchar) TO_FLOAT_CASE(uchar) TO_DOUBLE_CASE(uchar) @@ -535,6 +603,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, SIMPLE_CAST_CASE(uchar, kULong, ULong) SIMPLE_CAST_CASE(uchar, kUnsignedLong, ULong) + TO_HALF_CASE(uchar) TO_FLOAT_CASE(uchar) TO_DOUBLE_CASE(uchar) @@ -567,6 +636,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, SIMPLE_CAST_CASE(short, kULong, ULong) SIMPLE_CAST_CASE(short, kUnsignedLong, ULong) + TO_HALF_CASE(short) TO_FLOAT_CASE(short) TO_DOUBLE_CASE(short) @@ -599,6 +669,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, SIMPLE_CAST_CASE(ushort, kULong, ULong) SIMPLE_CAST_CASE(ushort, kUnsignedLong, ULong) + TO_HALF_CASE(ushort) TO_FLOAT_CASE(ushort) TO_DOUBLE_CASE(ushort) @@ -631,6 +702,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, SIMPLE_CAST_CASE(ushort, kULong, ULong) SIMPLE_CAST_CASE(ushort, kUnsignedLong, ULong) + TO_HALF_CASE(ushort) TO_FLOAT_CASE(ushort) TO_DOUBLE_CASE(ushort) @@ -663,6 +735,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, SIMPLE_CAST_CASE(int, kULong, ULong) SIMPLE_CAST_CASE(int, kUnsignedLong, ULong) + TO_HALF_CASE(int) TO_FLOAT_CASE(int) TO_DOUBLE_CASE(int) @@ -695,6 +768,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, SIMPLE_CAST_CASE(uint, kULong, ULong) SIMPLE_CAST_CASE(uint, kUnsignedLong, ULong) + TO_HALF_CASE(uint) TO_FLOAT_CASE(uint) TO_DOUBLE_CASE(uint) @@ -727,6 +801,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, SIMPLE_CAST_CASE(uint, kULong, ULong) SIMPLE_CAST_CASE(uint, kUnsignedLong, ULong) + TO_HALF_CASE(uint) TO_FLOAT_CASE(uint) TO_DOUBLE_CASE(uint) @@ -759,6 +834,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, DOWN_CAST_CASE(Long, kULong, ULong, saturate) DOWN_CAST_CASE(Long, kUnsignedLong, ULong, saturate) + TO_HALF_CASE(Long) TO_FLOAT_CASE(Long) TO_DOUBLE_CASE(Long) @@ -791,6 +867,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, U_DOWN_CAST_CASE(ULong, kUnsignedInt, uint, saturate) U_DOWN_CAST_CASE(ULong, kLong, Long, saturate) + TO_HALF_CASE(ULong) TO_FLOAT_CASE(ULong) TO_DOUBLE_CASE(ULong) @@ -823,6 +900,7 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, U_DOWN_CAST_CASE(ULong, kUnsignedInt, uint, saturate) U_DOWN_CAST_CASE(ULong, kLong, Long, saturate) + TO_HALF_CASE(ULong) TO_FLOAT_CASE(ULong) TO_DOUBLE_CASE(ULong) @@ -833,6 +911,45 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, } break; + case kHalf: + halfPtr = (cl_half *)inRaw; + switch (outType) + { + BOOL_CASE(half) + + HALF_ROUND_CASE(kChar, schar, roundType, saturate) + HALF_ROUND_CASE(kUChar, uchar, roundType, saturate) + HALF_ROUND_CASE(kUnsignedChar, uchar, roundType, saturate) + HALF_ROUND_CASE(kShort, short, roundType, saturate) + HALF_ROUND_CASE(kUShort, ushort, roundType, saturate) + HALF_ROUND_CASE(kUnsignedShort, ushort, roundType, saturate) + HALF_ROUND_CASE(kInt, int, roundType, saturate) + HALF_ROUND_CASE(kUInt, uint, roundType, saturate) + HALF_ROUND_CASE(kUnsignedInt, uint, roundType, saturate) + HALF_ROUND_CASE(kLong, Long, roundType, saturate) + HALF_ROUND_CASE(kULong, ULong, roundType, saturate) + HALF_ROUND_CASE(kUnsignedLong, ULong, roundType, saturate) + + case kHalf: + memcpy(outRaw, inRaw, get_explicit_type_size(inType)); + break; + + case kFloat: + floatPtr = (float *)outRaw; + *floatPtr = cl_half_to_float(*halfPtr); + break; + case kDouble: + doublePtr = (double *)outRaw; + *doublePtr = cl_half_to_float(*halfPtr); + break; + + default: + log_error("ERROR: Invalid type given to " + "convert_explicit_value!!\n"); + break; + } + break; + case kFloat: floatPtr = (float *)inRaw; switch (outType) @@ -852,6 +969,8 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, FLOAT_ROUND_CASE(kULong, ULong, roundType, saturate) FLOAT_ROUND_CASE(kUnsignedLong, ULong, roundType, saturate) + TO_HALF_CASE(float) + case kFloat: memcpy(outRaw, inRaw, get_explicit_type_size(inType)); break; @@ -884,6 +1003,8 @@ void convert_explicit_value(void *inRaw, void *outRaw, ExplicitType inType, DOUBLE_ROUND_CASE(kULong, ULong, roundType, saturate) DOUBLE_ROUND_CASE(kUnsignedLong, ULong, roundType, saturate) + TO_HALF_CASE(double) + TO_FLOAT_CASE(double); case kDouble: From 242c767f73c5f76fcc2c50a8cda86097d2891b0d Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Thu, 6 Jul 2023 12:58:27 +0200 Subject: [PATCH 06/12] Corrected condition to verify if additional pragma is necessary (issue #142, basic) --- test_conformance/basic/test_explicit_s2v.cpp | 8 ++++++-- test_conformance/basic/test_fpmath.cpp | 10 +++++++--- 2 files changed, 13 insertions(+), 5 deletions(-) diff --git a/test_conformance/basic/test_explicit_s2v.cpp b/test_conformance/basic/test_explicit_s2v.cpp index 2f0fc86d95..238b6e06e8 100644 --- a/test_conformance/basic/test_explicit_s2v.cpp +++ b/test_conformance/basic/test_explicit_s2v.cpp @@ -26,6 +26,8 @@ #include "harness/conversions.h" #include "harness/typeWrappers.h" +namespace { + // clang-format off #define DECLARE_S2V_IDENT_KERNEL(srctype,dsttype,size) \ @@ -286,10 +288,10 @@ struct TypesIterator }; std::stringstream sstr; - if (srcType == kDouble || dstType == kDouble) + if (srcT == kDouble || dstT == kDouble) sstr << "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n"; - if (srcType == kHalf || dstType == kHalf) + if (srcT == kHalf || dstT == kHalf) sstr << "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n"; snprintf(pragma, sizeof(pragma), "%s", sstr.str().c_str()); @@ -333,6 +335,8 @@ struct TypesIterator 128; // hardcoded in original test }; +} // anonymous namespace + int test_explicit_s2v(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { diff --git a/test_conformance/basic/test_fpmath.cpp b/test_conformance/basic/test_fpmath.cpp index 6719e72816..8a2210fbcf 100644 --- a/test_conformance/basic/test_fpmath.cpp +++ b/test_conformance/basic/test_fpmath.cpp @@ -33,7 +33,11 @@ #include "procs.h" -static const char *fp_kernel_code = R"( +extern cl_half_rounding_mode halfRoundingMode; + +namespace { + +const char *fp_kernel_code = R"( %s __kernel void test_fp(__global TYPE *srcA, __global TYPE *srcB, __global TYPE *dst) { @@ -42,8 +46,6 @@ __kernel void test_fp(__global TYPE *srcA, __global TYPE *srcB, __global TYPE *d dst[tid] = srcA[tid] OP srcB[tid]; })"; -extern cl_half_rounding_mode halfRoundingMode; - #define HFF(num) cl_half_from_float(num, halfRoundingMode) #define HTF(num) cl_half_to_float(num) @@ -370,6 +372,8 @@ struct TypesIterator std::map type2name; }; +} // anonymous namespace + int test_fpmath(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) { From ffb7738aeb45b5de7d73f39c3955a377d4ae1cfc Mon Sep 17 00:00:00 2001 From: Vasu Penugonda Date: Tue, 1 Aug 2023 13:12:16 +0530 Subject: [PATCH 07/12] Add NaN check for half to float conversion --- test_conformance/basic/test_explicit_s2v.cpp | 27 ++++++++++++++++++++ 1 file changed, 27 insertions(+) diff --git a/test_conformance/basic/test_explicit_s2v.cpp b/test_conformance/basic/test_explicit_s2v.cpp index 238b6e06e8..c23ad76046 100644 --- a/test_conformance/basic/test_explicit_s2v.cpp +++ b/test_conformance/basic/test_explicit_s2v.cpp @@ -98,6 +98,26 @@ const char * kernel_explicit_s2v_set[NUM_VEC_TYPES][NUM_VEC_TYPES][5] = { // clang-format on +int IsFloatNaN(double x) +{ + union { + cl_float d; + cl_uint u; + } u; + u.d = (cl_float)x; + return ((u.u & 0x7fffffffU) > 0x7F800000U); +} + +bool IsHalfNaN(cl_half v) +{ + // Extract FP16 exponent and mantissa + uint16_t h_exp = (((cl_half)v) >> (CL_HALF_MANT_DIG - 1)) & 0x1F; + uint16_t h_mant = ((cl_half)v) & 0x3FF; + + // NaN test + return (h_exp == 0x1F && h_mant != 0); +} + int test_explicit_s2v_function(cl_context context, cl_command_queue queue, cl_kernel kernel, ExplicitType srcType, unsigned int count, ExplicitType destType, @@ -159,6 +179,13 @@ int test_explicit_s2v_function(cl_context context, cl_command_queue queue, { if( memcmp( convertedData, outPtr + destTypeSize * s, destTypeSize ) != 0 ) { + if ((srcType == kHalf) && (destType == kFloat) + && IsHalfNaN(*reinterpret_cast(inPtr)) + && IsFloatNaN(*reinterpret_cast(outPtr + destTypeSize * s))) + { + continue; + } + unsigned int *p = (unsigned int *)outPtr; log_error( "ERROR: Output value %d:%d does not validate for size %d:%d!\n", i, s, vecSize, (int)destTypeSize ); log_error( " Input: 0x%0*x\n", (int)( paramSize * 2 ), *(unsigned int *)inPtr & ( 0xffffffff >> ( 32 - paramSize * 8 ) ) ); From c7055fd9053a820f4945c5d706f7d52d5176828c Mon Sep 17 00:00:00 2001 From: Sreelakshmi Haridas Date: Tue, 5 Sep 2023 15:27:26 -0600 Subject: [PATCH 08/12] check-format fixes --- test_conformance/basic/test_explicit_s2v.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/test_conformance/basic/test_explicit_s2v.cpp b/test_conformance/basic/test_explicit_s2v.cpp index c23ad76046..473b828283 100644 --- a/test_conformance/basic/test_explicit_s2v.cpp +++ b/test_conformance/basic/test_explicit_s2v.cpp @@ -181,7 +181,8 @@ int test_explicit_s2v_function(cl_context context, cl_command_queue queue, { if ((srcType == kHalf) && (destType == kFloat) && IsHalfNaN(*reinterpret_cast(inPtr)) - && IsFloatNaN(*reinterpret_cast(outPtr + destTypeSize * s))) + && IsFloatNaN(*reinterpret_cast( + outPtr + destTypeSize * s))) { continue; } From 7d78bcd592de74966af044b5b30e06b6da0b31ce Mon Sep 17 00:00:00 2001 From: Vasu Penugonda Date: Wed, 4 Oct 2023 11:29:50 +0530 Subject: [PATCH 09/12] Add NaN check for all float types Use std::isnan for float/double types. Change-Id: I005bddccaa3f8490ac59b2aa431ed315733ad143 --- test_conformance/basic/test_explicit_s2v.cpp | 24 ++++++++------------ 1 file changed, 10 insertions(+), 14 deletions(-) diff --git a/test_conformance/basic/test_explicit_s2v.cpp b/test_conformance/basic/test_explicit_s2v.cpp index 473b828283..03d67e1971 100644 --- a/test_conformance/basic/test_explicit_s2v.cpp +++ b/test_conformance/basic/test_explicit_s2v.cpp @@ -13,6 +13,8 @@ // See the License for the specific language governing permissions and // limitations under the License. // +#include +#define isnan std::isnan #include "harness/compat.h" #include @@ -98,16 +100,6 @@ const char * kernel_explicit_s2v_set[NUM_VEC_TYPES][NUM_VEC_TYPES][5] = { // clang-format on -int IsFloatNaN(double x) -{ - union { - cl_float d; - cl_uint u; - } u; - u.d = (cl_float)x; - return ((u.u & 0x7fffffffU) > 0x7F800000U); -} - bool IsHalfNaN(cl_half v) { // Extract FP16 exponent and mantissa @@ -179,10 +171,14 @@ int test_explicit_s2v_function(cl_context context, cl_command_queue queue, { if( memcmp( convertedData, outPtr + destTypeSize * s, destTypeSize ) != 0 ) { - if ((srcType == kHalf) && (destType == kFloat) - && IsHalfNaN(*reinterpret_cast(inPtr)) - && IsFloatNaN(*reinterpret_cast( - outPtr + destTypeSize * s))) + bool isSrcNaN = (((srcType == kHalf) && IsHalfNaN(*reinterpret_cast(inPtr))) + || ((srcType == kFloat) && isnan(*reinterpret_cast(inPtr))) + || ((srcType == kDouble) && isnan(*reinterpret_cast(inPtr)))); + bool isDestNaN = (((destType == kHalf) && IsHalfNaN(*reinterpret_cast(outPtr + destTypeSize * s))) + || ((destType == kFloat) && isnan(*reinterpret_cast(outPtr + destTypeSize * s))) + || ((destType == kDouble) && isnan(*reinterpret_cast(outPtr + destTypeSize * s)))); + + if (isSrcNaN && isDestNaN) { continue; } From cc6389baf421bfd33a5dc989e5108d416e34afcd Mon Sep 17 00:00:00 2001 From: Vasu Penugonda Date: Tue, 10 Oct 2023 19:53:54 +0530 Subject: [PATCH 10/12] Fix Ubuntu build error with isnan macro definition Change-Id: I671ed826a9631fbbc66d0aa9b674ab00124c7967 --- test_conformance/basic/test_explicit_s2v.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/test_conformance/basic/test_explicit_s2v.cpp b/test_conformance/basic/test_explicit_s2v.cpp index 03d67e1971..c23d976e51 100644 --- a/test_conformance/basic/test_explicit_s2v.cpp +++ b/test_conformance/basic/test_explicit_s2v.cpp @@ -14,7 +14,8 @@ // limitations under the License. // #include -#define isnan std::isnan +using std::isnan; +#define isnan isnan #include "harness/compat.h" #include From 9ad151aa341167c7dd863525f1749e3811adc289 Mon Sep 17 00:00:00 2001 From: Sreelakshmi Haridas Date: Thu, 12 Oct 2023 14:13:06 -0600 Subject: [PATCH 11/12] Check format fixes --- test_conformance/basic/test_explicit_s2v.cpp | 22 ++++++++++++++------ 1 file changed, 16 insertions(+), 6 deletions(-) diff --git a/test_conformance/basic/test_explicit_s2v.cpp b/test_conformance/basic/test_explicit_s2v.cpp index c23d976e51..ea09d13a57 100644 --- a/test_conformance/basic/test_explicit_s2v.cpp +++ b/test_conformance/basic/test_explicit_s2v.cpp @@ -172,12 +172,22 @@ int test_explicit_s2v_function(cl_context context, cl_command_queue queue, { if( memcmp( convertedData, outPtr + destTypeSize * s, destTypeSize ) != 0 ) { - bool isSrcNaN = (((srcType == kHalf) && IsHalfNaN(*reinterpret_cast(inPtr))) - || ((srcType == kFloat) && isnan(*reinterpret_cast(inPtr))) - || ((srcType == kDouble) && isnan(*reinterpret_cast(inPtr)))); - bool isDestNaN = (((destType == kHalf) && IsHalfNaN(*reinterpret_cast(outPtr + destTypeSize * s))) - || ((destType == kFloat) && isnan(*reinterpret_cast(outPtr + destTypeSize * s))) - || ((destType == kDouble) && isnan(*reinterpret_cast(outPtr + destTypeSize * s)))); + bool isSrcNaN = + (((srcType == kHalf) + && IsHalfNaN(*reinterpret_cast(inPtr))) + || ((srcType == kFloat) + && isnan(*reinterpret_cast(inPtr))) + || ((srcType == kDouble) + && isnan(*reinterpret_cast(inPtr)))); + bool isDestNaN = (((destType == kHalf) + && IsHalfNaN(*reinterpret_cast( + outPtr + destTypeSize * s))) + || ((destType == kFloat) + && isnan(*reinterpret_cast( + outPtr + destTypeSize * s))) + || ((destType == kDouble) + && isnan(*reinterpret_cast( + outPtr + destTypeSize * s)))); if (isSrcNaN && isDestNaN) { From a541f4bb3181e1ca7cfb4c27e2d3dc75687b3dd9 Mon Sep 17 00:00:00 2001 From: Sreelakshmi Haridas Date: Fri, 13 Oct 2023 15:37:43 -0600 Subject: [PATCH 12/12] NAN define not needed anymore --- test_conformance/basic/test_explicit_s2v.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/test_conformance/basic/test_explicit_s2v.cpp b/test_conformance/basic/test_explicit_s2v.cpp index ea09d13a57..067afb435b 100644 --- a/test_conformance/basic/test_explicit_s2v.cpp +++ b/test_conformance/basic/test_explicit_s2v.cpp @@ -15,7 +15,6 @@ // #include using std::isnan; -#define isnan isnan #include "harness/compat.h" #include