From 0a00a1f5b03b23e2e63482d7acbaac4e574ff764 Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Tue, 27 Aug 2024 17:41:05 +0200 Subject: [PATCH 1/7] Added printf test for double type cases (#2022) according to work plan for issue https://github.com/KhronosGroup/OpenCL-CTS/issues/1058 --- test_common/harness/kernelHelpers.cpp | 29 ++- test_common/harness/kernelHelpers.h | 4 +- test_conformance/printf/test_printf.cpp | 52 ++++- test_conformance/printf/test_printf.h | 3 + test_conformance/printf/util_printf.cpp | 279 +++++++++++++++++++++--- 5 files changed, 309 insertions(+), 58 deletions(-) diff --git a/test_common/harness/kernelHelpers.cpp b/test_common/harness/kernelHelpers.cpp index 633b05e5c7..c307fca559 100644 --- a/test_common/harness/kernelHelpers.cpp +++ b/test_common/harness/kernelHelpers.cpp @@ -1511,30 +1511,27 @@ size_t get_min_alignment(cl_context context) return align_size; } -cl_device_fp_config get_default_rounding_mode(cl_device_id device, - const cl_uint ¶m) +cl_device_fp_config get_default_rounding_mode(const cl_device_id device, + const cl_uint param) { - if (param == CL_DEVICE_DOUBLE_FP_CONFIG) - test_error_ret( - -1, - "FAILURE: CL_DEVICE_DOUBLE_FP_CONFIG not supported by this routine", - 0); - char profileStr[128] = ""; - cl_device_fp_config single = 0; - int error = clGetDeviceInfo(device, param, sizeof(single), &single, NULL); + cl_device_fp_config config = 0; + int error = clGetDeviceInfo(device, param, sizeof(config), &config, NULL); if (error) { - std::string message = std::string("Unable to get device ") - + std::string(param == CL_DEVICE_HALF_FP_CONFIG - ? "CL_DEVICE_HALF_FP_CONFIG" - : "CL_DEVICE_SINGLE_FP_CONFIG"); + std::string config_name = "CL_DEVICE_SINGLE_FP_CONFIG"; + if (param == CL_DEVICE_HALF_FP_CONFIG) + config_name = "CL_DEVICE_HALF_FP_CONFIG"; + else if (param == CL_DEVICE_DOUBLE_FP_CONFIG) + config_name = "CL_DEVICE_DOUBLE_FP_CONFIG"; + std::string message = + std::string("Unable to get device ") + config_name; test_error_ret(error, message.c_str(), 0); } - if (single & CL_FP_ROUND_TO_NEAREST) return CL_FP_ROUND_TO_NEAREST; + if (config & CL_FP_ROUND_TO_NEAREST) return CL_FP_ROUND_TO_NEAREST; - if (0 == (single & CL_FP_ROUND_TO_ZERO)) + if (0 == (config & CL_FP_ROUND_TO_ZERO)) test_error_ret(-1, "FAILURE: device must support either " "CL_FP_ROUND_TO_ZERO or CL_FP_ROUND_TO_NEAREST", diff --git a/test_common/harness/kernelHelpers.h b/test_common/harness/kernelHelpers.h index 86a6919293..518b27499a 100644 --- a/test_common/harness/kernelHelpers.h +++ b/test_common/harness/kernelHelpers.h @@ -157,8 +157,8 @@ size_t get_min_alignment(cl_context context); /* Helper to obtain the default rounding mode for single precision computation. * (Double is always CL_FP_ROUND_TO_NEAREST.) Returns 0 on error. */ cl_device_fp_config -get_default_rounding_mode(cl_device_id device, - const cl_uint ¶m = CL_DEVICE_SINGLE_FP_CONFIG); +get_default_rounding_mode(const cl_device_id device, + const cl_uint param = CL_DEVICE_SINGLE_FP_CONFIG); #define PASSIVE_REQUIRE_IMAGE_SUPPORT(device) \ if (checkForImageSupport(device)) \ diff --git a/test_conformance/printf/test_printf.cpp b/test_conformance/printf/test_printf.cpp index 3d539ed572..d59e06825b 100644 --- a/test_conformance/printf/test_printf.cpp +++ b/test_conformance/printf/test_printf.cpp @@ -698,6 +698,15 @@ int doTest(cl_command_queue queue, cl_context context, return TEST_SKIPPED_ITSELF; } + if ((allTestCase[testId]->_type == TYPE_DOUBLE + || allTestCase[testId]->_type == TYPE_DOUBLE_LIMITS) + && !is_extension_available(device, "cl_khr_fp64")) + { + log_info("Skipping double because cl_khr_fp64 extension is not " + "supported.\n"); + return TEST_SKIPPED_ITSELF; + } + auto& genParams = allTestCase[testId]->_genParameters; auto fail_count = s_test_fail; @@ -708,18 +717,25 @@ int doTest(cl_command_queue queue, cl_context context, { if (allTestCase[testId]->_type == TYPE_VECTOR) { - if ((strcmp(allTestCase[testId]->_genParameters[testNum].dataType, - "half") - == 0) - && !is_extension_available(device, "cl_khr_fp16")) - { - log_info("Skipping half because cl_khr_fp16 extension is not " - "supported.\n"); + auto is_vector_type_supported = [&](const char* type_name, + const char* ext_name) { + if ((strcmp(genParams[testNum].dataType, type_name) == 0) + && !is_extension_available(device, ext_name)) + { + log_info("Skipping %s because %s extension " + "is not supported.\n", + type_name, ext_name); - s_test_skip++; - s_test_cnt++; - continue; - } + s_test_skip++; + s_test_cnt++; + return false; + } + return true; + }; + + if (!is_vector_type_supported("half", "cl_khr_fp16")) continue; + + if (!is_vector_type_supported("double", "cl_khr_fp64")) continue; // Long support for varible type if (!strcmp(allTestCase[testId]->_genParameters[testNum].dataType, @@ -935,6 +951,18 @@ int test_float_limits(cl_device_id deviceID, cl_context context, return doTest(gQueue, gContext, TYPE_FLOAT_LIMITS, deviceID); } +int test_double(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) +{ + return doTest(gQueue, gContext, TYPE_DOUBLE, deviceID); +} + +int test_double_limits(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) +{ + return doTest(gQueue, gContext, TYPE_DOUBLE_LIMITS, deviceID); +} + int test_octal(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { @@ -1020,6 +1048,8 @@ test_definition test_list[] = { ADD_TEST(half_limits), ADD_TEST(float), ADD_TEST(float_limits), + ADD_TEST(double), + ADD_TEST(double_limits), ADD_TEST(octal), ADD_TEST(unsigned), ADD_TEST(hexadecimal), diff --git a/test_conformance/printf/test_printf.h b/test_conformance/printf/test_printf.h index a2cd9ed2be..51f351170f 100644 --- a/test_conformance/printf/test_printf.h +++ b/test_conformance/printf/test_printf.h @@ -50,6 +50,8 @@ enum PrintfTestType TYPE_HALF_LIMITS, TYPE_FLOAT, TYPE_FLOAT_LIMITS, + TYPE_DOUBLE, + TYPE_DOUBLE_LIMITS, TYPE_OCTAL, TYPE_UNSIGNED, TYPE_HEXADEC, @@ -80,6 +82,7 @@ struct printDataGenParameters static std::vector correctBufferInt; static std::vector correctBufferHalf; static std::vector correctBufferFloat; +static std::vector correctBufferDouble; static std::vector correctBufferOctal; static std::vector correctBufferUnsigned; static std::vector correctBufferHexadecimal; diff --git a/test_conformance/printf/util_printf.cpp b/test_conformance/printf/util_printf.cpp index 03d5eb171c..759a8dea1d 100644 --- a/test_conformance/printf/util_printf.cpp +++ b/test_conformance/printf/util_printf.cpp @@ -26,6 +26,8 @@ static void intRefBuilder(printDataGenParameters&, char*, const size_t); static void halfRefBuilder(printDataGenParameters&, char* rResult, const size_t); static void floatRefBuilder(printDataGenParameters&, char* rResult, const size_t); +static void doubleRefBuilder(printDataGenParameters&, char* rResult, + const size_t); static void octalRefBuilder(printDataGenParameters&, char*, const size_t); static void unsignedRefBuilder(printDataGenParameters&, char*, const size_t); static void hexRefBuilder(printDataGenParameters&, char*, const size_t); @@ -295,99 +297,99 @@ std::vector printFloatGenParameters = { // Default(right)-justified - { { "%f" }, "10.3456" }, + { { "%f" }, "10.3456f" }, // One position after the decimal,default(right)-justified - { { "%.1f" }, "10.3456" }, + { { "%.1f" }, "10.3456f" }, // Two positions after the decimal,default(right)-justified - { { "%.2f" }, "10.3456" }, + { { "%.2f" }, "10.3456f" }, //(Minimum)Eight-wide,three positions after the // decimal,default(right)-justified - { { "%8.3f" }, "10.3456" }, + { { "%8.3f" }, "10.3456f" }, //(Minimum)Eight-wide,two positions after the // decimal,zero-filled,default(right)-justified - { { "%08.2f" }, "10.3456" }, + { { "%08.2f" }, "10.3456f" }, //(Minimum)Eight-wide,two positions after the decimal,left-justified - { { "%-8.2f" }, "10.3456" }, + { { "%-8.2f" }, "10.3456f" }, //(Minimum)Eight-wide,two positions after the decimal,with // sign,default(right)-justified - { { "%+8.2f" }, "-10.3456" }, + { { "%+8.2f" }, "-10.3456f" }, // Zero positions after the // decimal([floor]rounding),default(right)-justified - { { "%.0f" }, "0.1" }, + { { "%.0f" }, "0.1f" }, // Zero positions after the decimal([ceil]rounding),default(right)-justified - { { "%.0f" }, "0.6" }, + { { "%.0f" }, "0.6f" }, // Zero-filled,default positions number after the // decimal,default(right)-justified - { { "%0f" }, "0.6" }, + { { "%0f" }, "0.6f" }, // Double argument representing floating-point,used by f // style,default(right)-justified - { { "%4g" }, "12345.6789" }, + { { "%4g" }, "12345.6789f" }, // Double argument representing floating-point,used by e // style,default(right)-justified - { { "%4.2g" }, "12345.6789" }, + { { "%4.2g" }, "12345.6789f" }, // Double argument representing floating-point,used by f // style,default(right)-justified - { { "%4G" }, "0.0000023" }, + { { "%4G" }, "0.0000023f" }, // Double argument representing floating-point,used by e // style,default(right)-justified - { { "%4G" }, "0.023" }, + { { "%4G" }, "0.023f" }, // Double argument representing floating-point,with // exponent,left-justified,default(right)-justified // Use a value that is exactly representable as 32-bit float. - { { "%-#20.15e" }, "789456128.0" }, + { { "%-#20.15e" }, "789456128.f" }, // Double argument representing floating-point,with // exponent,left-justified,with sign,capital E,default(right)-justified // Use a value that is exactly representable as 32-bit float. - { { "%+#21.15E" }, "789456128.0" }, + { { "%+#21.15E" }, "789456128.f" }, // Double argument representing floating-point,in [-]xh.hhhhpAd style - { { "%.6a" }, "0.1" }, + { { "%.6a" }, "0.1f" }, //(Minimum)Ten-wide,Double argument representing floating-point,in // xh.hhhhpAd style,default(right)-justified - { { "%10.2a" }, "9990.235" }, + { { "%10.2a" }, "9990.235f" }, //(Minimum)Ten-wide,two positions after the decimal,with // a blank space inserted before the value, default(right)-justified - { { "% 10.2f" }, "1.25" }, + { { "% 10.2f" }, "1.25f" }, //(Minimum)Eight-wide,two positions after the decimal,with // zeros inserted before the value, default(right)-justified - { { "%08.2f" }, "3.14" }, + { { "%08.2f" }, "3.14f" }, }; //--------------------------------------------------------- @@ -485,6 +487,200 @@ testCase testCaseFloatLimits = { }; +//============================================== + +// double + +//============================================== + + +//-------------------------------------------------------- + +// [string] format | [string] double-data representation | + +//-------------------------------------------------------- + +std::vector printDoubleGenParameters = { + + // Default(right)-justified + + { { "%f" }, "10.3456" }, + + // One position after the decimal,default(right)-justified + + { { "%.1f" }, "10.3456" }, + + // Two positions after the decimal,default(right)-justified + + { { "%.2f" }, "10.3456" }, + + //(Minimum)Eight-wide,three positions after the + // decimal,default(right)-justified + + { { "%8.3f" }, "10.3456" }, + + //(Minimum)Eight-wide,two positions after the + // decimal,zero-filled,default(right)-justified + + { { "%08.2f" }, "10.3456" }, + + //(Minimum)Eight-wide,two positions after the decimal,left-justified + + { { "%-8.2f" }, "10.3456" }, + + //(Minimum)Eight-wide,two positions after the decimal,with + // sign,default(right)-justified + + { { "%+8.2f" }, "-10.3456" }, + + // Zero positions after the + // decimal([floor]rounding),default(right)-justified + + { { "%.0f" }, "0.1" }, + + // Zero positions after the decimal([ceil]rounding),default(right)-justified + + { { "%.0f" }, "0.6" }, + + // Zero-filled,default positions number after the + // decimal,default(right)-justified + + { { "%0f" }, "0.6" }, + + // Double argument representing floating-point,used by f + // style,default(right)-justified + + { { "%4g" }, "12345.6789" }, + + // Double argument representing floating-point,used by e + // style,default(right)-justified + + { { "%4.2g" }, "12345.6789" }, + + // Double argument representing floating-point,used by f + // style,default(right)-justified + + { { "%4G" }, "0.0000023" }, + + // Double argument representing floating-point,used by e + // style,default(right)-justified + + { { "%4G" }, "0.023" }, + + // Double argument representing floating-point,with + // exponent,left-justified,default(right)-justified + + { { "%-#20.15e" }, "789456123.0" }, + + // Double argument representing floating-point,with + // exponent,left-justified,with sign,capital E,default(right)-justified + + { { "%+#21.15E" }, "789456123.0" }, + + // Double argument representing floating-point,in [-]xh.hhhhpAd style + + { { "%.6a" }, "0.1" }, + + //(Minimum)Ten-wide,Double argument representing floating-point,in + // xh.hhhhpAd style,default(right)-justified + + { { "%10.2a" }, "9990.235" }, +}; + +//--------------------------------------------------------- + +// Test case for double | + +//--------------------------------------------------------- + +testCase testCaseDouble = { + + TYPE_DOUBLE, + + correctBufferDouble, + + printDoubleGenParameters, + + doubleRefBuilder, + + kdouble + +}; + +//============================================== + +// double limits + +//============================================== + + +//-------------------------------------------------------- + +// [string] double | [string] double-data representation | + +//-------------------------------------------------------- + + +std::vector printDoubleLimitsGenParameters = { + + // Infinity (1.0/0.0) + { { "%f", "%e", "%g", "%a" }, "1.0/0.0" }, + + // NaN + { { "%f", "%e", "%g", "%a" }, "nan(0UL)" }, + + // NaN + { { "%f", "%e", "%g", "%a" }, "acospi(2.0)" }, + + // Infinity (1.0/0.0) + { { "%F", "%E", "%G", "%A" }, "1.0/0.0" }, + + // NaN + { { "%F", "%E", "%G", "%A" }, "nan(0UL)" }, + + // NaN + { { "%F", "%E", "%G", "%A" }, "acospi(2.0)" } +}; +//-------------------------------------------------------- + +// Lookup table - [string]double-correct buffer | + +//-------------------------------------------------------- + +std::vector correctBufferDoubleLimits = { + + "inf", + + "nan", + + "nan", + + "INF", + + "NAN", + + "NAN" + +}; + +//--------------------------------------------------------- + +// Test case for float | + +//--------------------------------------------------------- + +testCase testCaseDoubleLimits = { + + TYPE_DOUBLE_LIMITS, + + correctBufferDoubleLimits, + + printDoubleLimitsGenParameters, + + NULL + +}; + //========================================================= // octal @@ -1208,11 +1404,12 @@ testCase testCaseMixedFormat = { TYPE_MIXED_FORMAT_RANDOM, //------------------------------------------------------------------------------- std::vector allTestCase = { - &testCaseInt, &testCaseHalf, &testCaseHalfLimits, - &testCaseFloat, &testCaseFloatLimits, &testCaseOctal, - &testCaseUnsigned, &testCaseHexadecimal, &testCaseChar, - &testCaseString, &testCaseFormatString, &testCaseVector, - &testCaseAddrSpace, &testCaseMixedFormat + &testCaseInt, &testCaseHalf, &testCaseHalfLimits, + &testCaseFloat, &testCaseFloatLimits, &testCaseDouble, + &testCaseDoubleLimits, &testCaseOctal, &testCaseUnsigned, + &testCaseHexadecimal, &testCaseChar, &testCaseString, + &testCaseFormatString, &testCaseVector, &testCaseAddrSpace, + &testCaseMixedFormat }; //----------------------------------------- @@ -1344,6 +1541,13 @@ static void floatRefBuilder(printDataGenParameters& params, char* refResult, con strtof(params.dataRepresentation, NULL)); } +static void doubleRefBuilder(printDataGenParameters& params, char* refResult, + const size_t refSize) +{ + snprintf(refResult, refSize, params.genericFormats.front().c_str(), + strtod(params.dataRepresentation, NULL)); +} + static void octalRefBuilder(printDataGenParameters& params, char* refResult, const size_t refSize) { const unsigned long int data = strtoul(params.dataRepresentation, NULL, 10); @@ -1374,11 +1578,17 @@ static void hexRefBuilder(printDataGenParameters& params, char* refResult, const */ void generateRef(const cl_device_id device) { + bool fp16_supported = is_extension_available(device, "cl_khr_fp16"); + bool fp64_supported = is_extension_available(device, "cl_khr_fp64"); + const cl_device_fp_config fpConfigSingle = get_default_rounding_mode(device); const cl_device_fp_config fpConfigHalf = (half_rounding_mode == CL_HALF_RTE) ? CL_FP_ROUND_TO_NEAREST : CL_FP_ROUND_TO_ZERO; + const cl_device_fp_config fpConfigDouble = fp64_supported + ? get_default_rounding_mode(device, CL_DEVICE_DOUBLE_FP_CONFIG) + : 0; const RoundingMode hostRound = get_round(); // Map device rounding to CTS rounding type @@ -1413,10 +1623,21 @@ void generateRef(const cl_device_id device) // Make sure the reference result is empty assert(caseToTest->_correctBuffer.size() == 0); - const cl_device_fp_config* fpConfig = &fpConfigSingle; - if (caseToTest->_type == TYPE_HALF - || caseToTest->_type == TYPE_HALF_LIMITS) - fpConfig = &fpConfigHalf; + const cl_device_fp_config* fpConfig = nullptr; + switch (caseToTest->_type) + { + case TYPE_HALF: + case TYPE_HALF_LIMITS: + if (!fp16_supported) continue; + fpConfig = &fpConfigHalf; + break; + case TYPE_DOUBLE: + case TYPE_DOUBLE_LIMITS: + if (!fp64_supported) continue; + fpConfig = &fpConfigDouble; + break; + default: fpConfig = &fpConfigSingle; break; + } RoundingMode deviceRound = get_rounding(*fpConfig); // Loop through each input From c0db608eacc10741f2559300d8e0af6e460eac18 Mon Sep 17 00:00:00 2001 From: Sven van Haastregt Date: Thu, 29 Aug 2024 14:45:15 +0200 Subject: [PATCH 2/7] printf: add missing vector length modifiers (#2044) The length modifier is required with a vector specifier. Fixes https://github.com/KhronosGroup/OpenCL-CTS/issues/2039 Signed-off-by: Sven van Haastregt --- test_conformance/printf/util_printf.cpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/test_conformance/printf/util_printf.cpp b/test_conformance/printf/util_printf.cpp index 759a8dea1d..5b96e8683d 100644 --- a/test_conformance/printf/util_printf.cpp +++ b/test_conformance/printf/util_printf.cpp @@ -1175,18 +1175,18 @@ std::vector printVectorGenParameters = { // Three component vector in scientific notation - { { "" }, "(1234.25f,987654.5f,0.0005f)", "%.2", "e", "float", "3" }, + { { "" }, "(1234.25f,987654.5f,0.0005f)", "%.2", "hle", "float", "3" }, // Four component vector in hexadecimal floating point, lowercase format - { { "" }, "(0.25f,0.5f,1.f,1.5f)", "%", "a", "float", "4" }, + { { "" }, "(0.25f,0.5f,1.f,1.5f)", "%", "hla", "float", "4" }, // Eight component vector in the shortest float representation { { "" }, "(1.f,2.f,3.f,4.f,1.5f,3.14f,2.5f,3.5f)", "%", - "g", + "hlg", "float", "8" }, @@ -1195,17 +1195,17 @@ std::vector printVectorGenParameters = { { { "" }, "(1,2,3,4,5,6,7,8,9,0,32,64,128,256,512,1024)", "%", - "o", + "hlo", "uint", "16" }, // Eight component vector in signed decimal integer format - { { "" }, "(1,-2,3,-4,5,-6,7,-8)", "%+", "i", "int", "8" }, + { { "" }, "(1,-2,3,-4,5,-6,7,-8)", "%+", "hli", "int", "8" }, // Four component vector in unsigned decimal integer format - { { "" }, "(512,1024,262144,1048576)", "%05", "u", "uint", "4" }, + { { "" }, "(512,1024,262144,1048576)", "%05", "hlu", "uint", "4" }, }; From a7162188d6d58289691da33ca43308677a6a611f Mon Sep 17 00:00:00 2001 From: Haonan Yang Date: Thu, 29 Aug 2024 20:45:40 +0800 Subject: [PATCH 3/7] [NFC] Use std::vector to store large size array. (#2060) windows default stack size is 1M, BUFFER_SIZE is 2 * 1024 * 1024, use array with BUFFER_SIZE elemets on stack can exceed available stack size limits. --- test_conformance/math_brute_force/ternary_float.cpp | 2 +- test_conformance/math_brute_force/ternary_half.cpp | 2 +- test_conformance/math_brute_force/unary_two_results_float.cpp | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/test_conformance/math_brute_force/ternary_float.cpp b/test_conformance/math_brute_force/ternary_float.cpp index 3e024b1e01..efde6063d8 100644 --- a/test_conformance/math_brute_force/ternary_float.cpp +++ b/test_conformance/math_brute_force/ternary_float.cpp @@ -139,7 +139,7 @@ int TestFunc_Float_Float_Float_Float(const Func *f, MTdata d, bool relaxedMode) float maxErrorVal3 = 0.0f; uint64_t step = getTestStep(sizeof(float), BUFFER_SIZE); - cl_uchar overflow[BUFFER_SIZE / sizeof(float)]; + std::vector overflow(BUFFER_SIZE / sizeof(float)); float float_ulps; if (gIsEmbedded) diff --git a/test_conformance/math_brute_force/ternary_half.cpp b/test_conformance/math_brute_force/ternary_half.cpp index a7aa93036a..5cb858d43a 100644 --- a/test_conformance/math_brute_force/ternary_half.cpp +++ b/test_conformance/math_brute_force/ternary_half.cpp @@ -79,7 +79,7 @@ int TestFunc_Half_Half_Half_Half(const Func *f, MTdata d, bool relaxedMode) constexpr size_t bufferElements = BUFFER_SIZE / sizeof(cl_half); - cl_uchar overflow[bufferElements]; + std::vector overflow(bufferElements); float half_ulps = f->half_ulps; int skipNanInf = (0 == strcmp("fma", f->nameInCode)); diff --git a/test_conformance/math_brute_force/unary_two_results_float.cpp b/test_conformance/math_brute_force/unary_two_results_float.cpp index 983e039587..8a5d3000a6 100644 --- a/test_conformance/math_brute_force/unary_two_results_float.cpp +++ b/test_conformance/math_brute_force/unary_two_results_float.cpp @@ -51,7 +51,7 @@ int TestFunc_Float2_Float(const Func *f, MTdata d, bool relaxedMode) float maxErrorVal1 = 0.0f; uint64_t step = getTestStep(sizeof(float), BUFFER_SIZE); int scale = (int)((1ULL << 32) / (16 * BUFFER_SIZE / sizeof(float)) + 1); - cl_uchar overflow[BUFFER_SIZE / sizeof(float)]; + std::vector overflow(BUFFER_SIZE / sizeof(float)); int isFract = 0 == strcmp("fract", f->nameInCode); int skipNanInf = isFract && !gInfNanSupport; From 90f523ea57291f4494a4794d79b0c78e1efc9938 Mon Sep 17 00:00:00 2001 From: Sven van Haastregt Date: Mon, 2 Sep 2024 12:02:28 +0200 Subject: [PATCH 4/7] api: fix -Wformat warnings (#2025) The main sources of warnings were: * Printing of a `size_t` which requires the `%zu` specifier. * Printing of `cl_long`/`cl_ulong` which is now done using the `PRI*64` macros to ensure portability across 32 and 64-bit builds. Signed-off-by: Sven van Haastregt --- test_conformance/api/test_api_consistency.cpp | 7 +- test_conformance/api/test_api_min_max.cpp | 107 ++++++++++-------- test_conformance/api/test_mem_object_info.cpp | 49 ++++---- .../test_mem_object_properties_queries.cpp | 2 +- .../api/test_pipe_properties_queries.cpp | 2 +- test_conformance/api/test_platform.cpp | 36 +++--- test_conformance/api/test_queries.cpp | 10 +- .../api/test_queue_properties_queries.cpp | 7 +- 8 files changed, 126 insertions(+), 94 deletions(-) diff --git a/test_conformance/api/test_api_consistency.cpp b/test_conformance/api/test_api_consistency.cpp index cc7e190a39..56c831815e 100644 --- a/test_conformance/api/test_api_consistency.cpp +++ b/test_conformance/api/test_api_consistency.cpp @@ -13,6 +13,8 @@ // See the License for the specific language governing permissions and // limitations under the License. // +#include + #include "testBase.h" #include "harness/testHarness.h" #include "harness/deviceInfo.h" @@ -158,8 +160,9 @@ static int check_atomic_capabilities(cl_device_atomic_capabilities atomicCaps, { if ((atomicCaps & requiredCaps) != requiredCaps) { - log_error("Atomic capabilities %llx is missing support for at least " - "one required capability %llx!\n", + log_error("Atomic capabilities %" PRIx64 + " is missing support for at least " + "one required capability %" PRIx64 "!\n", atomicCaps, requiredCaps); return TEST_FAIL; } diff --git a/test_conformance/api/test_api_min_max.cpp b/test_conformance/api/test_api_min_max.cpp index 40c6d19a01..f92f224ffc 100644 --- a/test_conformance/api/test_api_min_max.cpp +++ b/test_conformance/api/test_api_min_max.cpp @@ -18,6 +18,7 @@ #include "harness/testHarness.h" #include #include +#include const char *sample_single_param_kernel[] = { "__kernel void sample_test(__global int *src)\n" @@ -227,13 +228,13 @@ int test_min_max_work_items_sizes(cl_device_id deviceID, cl_context context, { if (deviceMaxWorkItemSize[i] < 1) { - log_error("MAX_WORK_ITEM_SIZE in dimension %d is invalid: %lu\n", i, + log_error("MAX_WORK_ITEM_SIZE in dimension %d is invalid: %zu\n", i, deviceMaxWorkItemSize[i]); errors++; } else { - log_info("Dimension %d has max work item size %lu\n", i, + log_info("Dimension %d has max work item size %zu\n", i, deviceMaxWorkItemSize[i]); } } @@ -257,7 +258,7 @@ int test_min_max_work_group_size(cl_device_id deviceID, cl_context context, NULL); test_error(error, "Unable to get max work group size from device"); - log_info("Reported %ld max device work group size.\n", deviceMaxThreadSize); + log_info("Reported %zu max device work group size.\n", deviceMaxThreadSize); if (deviceMaxThreadSize == 0) { @@ -601,8 +602,9 @@ int test_min_max_mem_alloc_size(cl_device_id deviceID, cl_context context, return -1; } - log_info("Reported max allocation size of %lld bytes (%gMB) and global mem " - "size of %lld bytes (%gMB).\n", + log_info("Reported max allocation size of %" PRIu64 + " bytes (%gMB) and global mem " + "size of %" PRIu64 " bytes (%gMB).\n", maxAllocSize, maxAllocSize / (1024.0 * 1024.0), memSize, memSize / (1024.0 * 1024.0)); @@ -611,14 +613,16 @@ int test_min_max_mem_alloc_size(cl_device_id deviceID, cl_context context, while (currentSize >= maxAllocSize / MAX_REDUCTION_FACTOR) { - log_info("Trying to create a buffer of size of %lld bytes (%gMB).\n", + log_info("Trying to create a buffer of size of %" PRIu64 + " bytes (%gMB).\n", currentSize, (double)currentSize / (1024.0 * 1024.0)); memHdl = clCreateBuffer(context, CL_MEM_READ_ONLY, (size_t)currentSize, NULL, &error); if (error == CL_MEM_OBJECT_ALLOCATION_FAILURE || error == CL_OUT_OF_RESOURCES || error == CL_OUT_OF_HOST_MEMORY) { - log_info("\tAllocation failed at size of %lld bytes (%gMB).\n", + log_info("\tAllocation failed at size of %" PRIu64 + " bytes (%gMB).\n", currentSize, (double)currentSize / (1024.0 * 1024.0)); currentSize -= minSizeToTry; continue; @@ -626,8 +630,8 @@ int test_min_max_mem_alloc_size(cl_device_id deviceID, cl_context context, test_error(error, "clCreateBuffer failed for maximum sized buffer."); return 0; } - log_error("Failed to allocate even %lld bytes (%gMB).\n", currentSize, - (double)currentSize / (1024.0 * 1024.0)); + log_error("Failed to allocate even %" PRIu64 " bytes (%gMB).\n", + currentSize, (double)currentSize / (1024.0 * 1024.0)); return -1; } @@ -671,7 +675,7 @@ int test_min_max_image_2d_width(cl_device_id deviceID, cl_context context, (int)maxDimension); return -1; } - log_info("Max reported width is %ld.\n", maxDimension); + log_info("Max reported width is %zu.\n", maxDimension); /* Verify we can use the format */ image_format_desc.image_channel_data_type = CL_UNORM_INT8; @@ -688,8 +692,8 @@ int test_min_max_image_2d_width(cl_device_id deviceID, cl_context context, deviceID, MAX_DEVICE_MEMORY_SIZE_DIVISOR); if ((cl_ulong)maxDimension * 1 * 4 > maxAllocSize) { - log_error("Can not allocate a large enough image (min size: %lld " - "bytes, max allowed: %lld bytes) to test.\n", + log_error("Can not allocate a large enough image (min size: %" PRIu64 + " bytes, max allowed: %" PRIu64 " bytes) to test.\n", (cl_ulong)maxDimension * 1 * 4, maxAllocSize); return -1; } @@ -748,7 +752,7 @@ int test_min_max_image_2d_height(cl_device_id deviceID, cl_context context, (int)maxDimension); return -1; } - log_info("Max reported height is %ld.\n", maxDimension); + log_info("Max reported height is %zu.\n", maxDimension); /* Verify we can use the format */ image_format_desc.image_channel_data_type = CL_UNORM_INT8; @@ -765,8 +769,8 @@ int test_min_max_image_2d_height(cl_device_id deviceID, cl_context context, deviceID, MAX_DEVICE_MEMORY_SIZE_DIVISOR); if ((cl_ulong)maxDimension * 1 * 4 > maxAllocSize) { - log_error("Can not allocate a large enough image (min size: %lld " - "bytes, max allowed: %lld bytes) to test.\n", + log_error("Can not allocate a large enough image (min size: %" PRIu64 + " bytes, max allowed: %" PRIu64 " bytes) to test.\n", (cl_ulong)maxDimension * 1 * 4, maxAllocSize); return -1; } @@ -815,7 +819,7 @@ int test_min_max_image_3d_width(cl_device_id deviceID, cl_context context, (int)maxDimension); return -1; } - log_info("Max reported width is %ld.\n", maxDimension); + log_info("Max reported width is %zu.\n", maxDimension); /* Verify we can use the format */ image_format_desc.image_channel_data_type = CL_UNORM_INT8; @@ -832,8 +836,8 @@ int test_min_max_image_3d_width(cl_device_id deviceID, cl_context context, deviceID, MAX_DEVICE_MEMORY_SIZE_DIVISOR); if ((cl_ulong)maxDimension * 2 * 4 > maxAllocSize) { - log_error("Can not allocate a large enough image (min size: %lld " - "bytes, max allowed: %lld bytes) to test.\n", + log_error("Can not allocate a large enough image (min size: %" PRIu64 + " bytes, max allowed: %" PRIu64 " bytes) to test.\n", (cl_ulong)maxDimension * 2 * 4, maxAllocSize); return -1; } @@ -883,7 +887,7 @@ int test_min_max_image_3d_height(cl_device_id deviceID, cl_context context, (int)maxDimension); return -1; } - log_info("Max reported height is %ld.\n", maxDimension); + log_info("Max reported height is %zu.\n", maxDimension); /* Verify we can use the format */ image_format_desc.image_channel_data_type = CL_UNORM_INT8; @@ -900,8 +904,8 @@ int test_min_max_image_3d_height(cl_device_id deviceID, cl_context context, deviceID, MAX_DEVICE_MEMORY_SIZE_DIVISOR); if ((cl_ulong)maxDimension * 2 * 4 > maxAllocSize) { - log_error("Can not allocate a large enough image (min size: %lld " - "bytes, max allowed: %lld bytes) to test.\n", + log_error("Can not allocate a large enough image (min size: %" PRIu64 + " bytes, max allowed: %" PRIu64 " bytes) to test.\n", (cl_ulong)maxDimension * 2 * 4, maxAllocSize); return -1; } @@ -952,7 +956,7 @@ int test_min_max_image_3d_depth(cl_device_id deviceID, cl_context context, (int)maxDimension); return -1; } - log_info("Max reported depth is %ld.\n", maxDimension); + log_info("Max reported depth is %zu.\n", maxDimension); /* Verify we can use the format */ image_format_desc.image_channel_data_type = CL_UNORM_INT8; @@ -969,8 +973,8 @@ int test_min_max_image_3d_depth(cl_device_id deviceID, cl_context context, deviceID, MAX_DEVICE_MEMORY_SIZE_DIVISOR); if ((cl_ulong)maxDimension * 1 * 4 > maxAllocSize) { - log_error("Can not allocate a large enough image (min size: %lld " - "bytes, max allowed: %lld bytes) to test.\n", + log_error("Can not allocate a large enough image (min size: %" PRIu64 + " bytes, max allowed: %" PRIu64 " bytes) to test.\n", (cl_ulong)maxDimension * 1 * 4, maxAllocSize); return -1; } @@ -1019,7 +1023,7 @@ int test_min_max_image_array_size(cl_device_id deviceID, cl_context context, (int)maxDimension); return -1; } - log_info("Max reported image array size is %ld.\n", maxDimension); + log_info("Max reported image array size is %zu.\n", maxDimension); /* Verify we can use the format */ image_format_desc.image_channel_data_type = CL_UNORM_INT8; @@ -1037,8 +1041,8 @@ int test_min_max_image_array_size(cl_device_id deviceID, cl_context context, deviceID, MAX_DEVICE_MEMORY_SIZE_DIVISOR); if ((cl_ulong)maxDimension * 1 * 4 > maxAllocSize) { - log_error("Can not allocate a large enough image (min size: %lld " - "bytes, max allowed: %lld bytes) to test.\n", + log_error("Can not allocate a large enough image (min size: %" PRIu64 + " bytes, max allowed: %" PRIu64 " bytes) to test.\n", (cl_ulong)maxDimension * 1 * 4, maxAllocSize); return -1; } @@ -1091,7 +1095,7 @@ int test_min_max_image_buffer_size(cl_device_id deviceID, cl_context context, (int)maxDimensionPixels); return -1; } - log_info("Max reported image buffer size is %ld pixels.\n", + log_info("Max reported image buffer size is %zu pixels.\n", maxDimensionPixels); pixelBytes = maxAllocSize / maxDimensionPixels; @@ -1209,7 +1213,7 @@ int test_min_max_parameter_size(cl_device_id deviceID, cl_context context, { log_info( "Trying a kernel with %ld int arguments (%ld bytes) and one " - "cl_mem (%ld bytes) for %ld bytes total.\n", + "cl_mem (%zu bytes) for %ld bytes total.\n", numberOfIntParametersToTry, sizeof(cl_int) * numberOfIntParametersToTry, sizeof(cl_mem), sizeof(cl_mem) + numberOfIntParametersToTry * sizeof(cl_int)); @@ -1218,7 +1222,7 @@ int test_min_max_parameter_size(cl_device_id deviceID, cl_context context, { log_info( "Trying a kernel with %ld long arguments (%ld bytes) and one " - "cl_mem (%ld bytes) for %ld bytes total.\n", + "cl_mem (%zu bytes) for %ld bytes total.\n", numberOfIntParametersToTry, sizeof(cl_long) * numberOfIntParametersToTry, sizeof(cl_mem), sizeof(cl_mem) + numberOfIntParametersToTry * sizeof(cl_long)); @@ -1365,8 +1369,8 @@ int test_min_max_parameter_size(cl_device_id deviceID, cl_context context, { if (long_result != expectedResult) { - log_error("Expected result (%lld) does not equal actual result " - "(%lld).\n", + log_error("Expected result (%" PRId64 + ") does not equal actual result (%" PRId64 ").\n", expectedResult, long_result); numberOfIntParametersToTry -= decrement; continue; @@ -1383,8 +1387,8 @@ int test_min_max_parameter_size(cl_device_id deviceID, cl_context context, { if (int_result != expectedResult) { - log_error("Expected result (%lld) does not equal actual result " - "(%d).\n", + log_error("Expected result (%" PRId64 + ") does not equal actual result (%d).\n", expectedResult, int_result); numberOfIntParametersToTry -= decrement; continue; @@ -1555,7 +1559,8 @@ int test_min_max_constant_buffer_size(cl_device_id deviceID, cl_context context, return -1; } - log_info("Reported max constant buffer size of %lld bytes.\n", maxSize); + log_info("Reported max constant buffer size of %" PRIu64 " bytes.\n", + maxSize); /* We have four buffers allocations */ maxGlobalSize = get_device_info_global_mem_size( @@ -1583,7 +1588,8 @@ int test_min_max_constant_buffer_size(cl_device_id deviceID, cl_context context, d = init_genrand(gRandomSeed); while (!allocPassed && currentSize >= maxSize / MAX_REDUCTION_FACTOR) { - log_info("Attempting to allocate constant buffer of size %lld bytes\n", + log_info("Attempting to allocate constant buffer of size %" PRIu64 + " bytes\n", maxSize); /* Create some I/O streams */ @@ -1631,8 +1637,8 @@ int test_min_max_constant_buffer_size(cl_device_id deviceID, cl_context context, || (error == CL_MEM_OBJECT_ALLOCATION_FAILURE) || (error == CL_OUT_OF_HOST_MEMORY)) { - log_info("Kernel enqueue failed at size %lld, trying at a reduced " - "size.\n", + log_info("Kernel enqueue failed at size %" PRIu64 + ", trying at a reduced size.\n", currentSize); currentSize -= stepSize; free(constantData); @@ -1657,8 +1663,8 @@ int test_min_max_constant_buffer_size(cl_device_id deviceID, cl_context context, || (event_status == CL_MEM_OBJECT_ALLOCATION_FAILURE) || (event_status == CL_OUT_OF_HOST_MEMORY)) { - log_info("Kernel event indicates failure at size %lld, trying " - "at a reduced size.\n", + log_info("Kernel event indicates failure at size %" PRIu64 + ", trying at a reduced size.\n", currentSize); currentSize -= stepSize; free(constantData); @@ -1715,7 +1721,8 @@ int test_min_max_constant_buffer_size(cl_device_id deviceID, cl_context context, } else if (currentSize != maxSize) { - log_info("Passed at reduced size. (%lld of %lld bytes)\n", + log_info("Passed at reduced size. (%" PRIu64 " of %" PRIu64 + " bytes)\n", currentSize, maxSize); return 0; } @@ -1778,11 +1785,11 @@ int test_min_max_constant_args(cl_device_id deviceID, cl_context context, deviceID, MAX_DEVICE_MEMORY_SIZE_DIVISOR); individualBufferSize = ((int)maxSize / 2) / maxArgs; - log_info( - "Reported max constant arg count of %u and max constant buffer " - "size of %llu. Test will attempt to allocate half of that, or %llu " - "buffers of size %zu.\n", - maxArgs, maxSize, maxArgs, individualBufferSize); + log_info("Reported max constant arg count of %u and max constant buffer " + "size of %" PRIu64 + ". Test will attempt to allocate half of that, or %u " + "buffers of size %zu.\n", + maxArgs, maxSize, maxArgs, individualBufferSize); str2 = (char *)malloc(sizeof(char) * 32 * (maxArgs + 2)); constArgs = (char *)malloc(sizeof(char) * 32 * (maxArgs + 2)); @@ -2034,7 +2041,7 @@ int test_min_max_local_mem_size(cl_device_id deviceID, cl_context context, return -1; } - log_info("Reported max local buffer size for device: %lld bytes.\n", + log_info("Reported max local buffer size for device: %" PRIu64 " bytes.\n", maxSize); /* Create a kernel to test with */ @@ -2052,7 +2059,7 @@ int test_min_max_local_mem_size(cl_device_id deviceID, cl_context context, "clGetKernelWorkGroupInfo for CL_KERNEL_LOCAL_MEM_SIZE failed"); log_info("Reported local buffer usage for kernel " - "(CL_KERNEL_LOCAL_MEM_SIZE): %lld bytes.\n", + "(CL_KERNEL_LOCAL_MEM_SIZE): %" PRIu64 " bytes.\n", kernelLocalUsage); /* Create some I/O streams */ @@ -2167,13 +2174,13 @@ int test_min_max_kernel_preferred_work_group_size_multiple( // Since the preferred size is only a performance hint, we can only really // check that we get a sane value back - log_info("size: %ld preferred: %ld max: %ld\n", max_workgroup_size, + log_info("size: %zu preferred: %zu max: %zu\n", max_workgroup_size, preferred_workgroup_size, max_local_workgroup_size[0]); if (preferred_workgroup_size > max_workgroup_size) { log_error("ERROR: Reported preferred workgroup multiple larger than " - "max workgroup size (preferred %ld, max %ld)\n", + "max workgroup size (preferred %zu, max %zu)\n", preferred_workgroup_size, max_workgroup_size); return -1; } diff --git a/test_conformance/api/test_mem_object_info.cpp b/test_conformance/api/test_mem_object_info.cpp index b31cebb3c3..e99b125a49 100644 --- a/test_conformance/api/test_mem_object_info.cpp +++ b/test_conformance/api/test_mem_object_info.cpp @@ -18,21 +18,24 @@ #include "harness/testHarness.h" -#define TEST_MEM_OBJECT_PARAM( mem, paramName, val, expected, name, type, cast ) \ -error = clGetMemObjectInfo( mem, paramName, sizeof( val ), &val, &size ); \ -test_error( error, "Unable to get mem object " name ); \ -if( val != expected ) \ -{ \ -log_error( "ERROR: Mem object " name " did not validate! (expected " type ", got " type " from %s:%d)\n", \ -expected, (cast)val, __FILE__, __LINE__ ); \ -return -1; \ -} \ -if( size != sizeof( val ) ) \ -{ \ -log_error( "ERROR: Returned size of mem object " name " does not validate! (expected %d, got %d from %s:%d)\n", \ -(int)sizeof( val ), (int)size , __FILE__, __LINE__ ); \ -return -1; \ -} +#define TEST_MEM_OBJECT_PARAM(mem, paramName, val, expected, name, type, cast) \ + error = clGetMemObjectInfo(mem, paramName, sizeof(val), &val, &size); \ + test_error(error, "Unable to get mem object " name); \ + if (val != expected) \ + { \ + log_error("ERROR: Mem object " name \ + " did not validate! (expected " type ", got " type \ + " from %s:%d)\n", \ + (cast)expected, (cast)val, __FILE__, __LINE__); \ + return -1; \ + } \ + if (size != sizeof(val)) \ + { \ + log_error("ERROR: Returned size of mem object " name \ + " does not validate! (expected %d, got %d from %s:%d)\n", \ + (int)sizeof(val), (int)size, __FILE__, __LINE__); \ + return -1; \ + } static void CL_CALLBACK mem_obj_destructor_callback( cl_mem, void * data ) { @@ -236,7 +239,8 @@ int test_get_buffer_info( cl_device_id deviceID, cl_context context, cl_command_ TEST_MEM_OBJECT_PARAM( bufferObject, CL_MEM_FLAGS, flags, (unsigned int)bufferFlags[ i ], "flags", "%d", unsigned int ) size_t sz; - TEST_MEM_OBJECT_PARAM( bufferObject, CL_MEM_SIZE, sz, (size_t)( addressAlign * 4 ), "size", "%ld", size_t ) + TEST_MEM_OBJECT_PARAM(bufferObject, CL_MEM_SIZE, sz, + (size_t)(addressAlign * 4), "size", "%zu", size_t) cl_uint mapCount; error = clGetMemObjectInfo( bufferObject, CL_MEM_MAP_COUNT, sizeof( mapCount ), &mapCount, &size ); @@ -265,7 +269,8 @@ int test_get_buffer_info( cl_device_id deviceID, cl_context context, cl_command_ TEST_MEM_OBJECT_PARAM( bufferObject, CL_MEM_ASSOCIATED_MEMOBJECT, origObj, (void *)NULL, "associated mem object", "%p", void * ) size_t offset; - TEST_MEM_OBJECT_PARAM( bufferObject, CL_MEM_OFFSET, offset, 0L, "offset", "%ld", size_t ) + TEST_MEM_OBJECT_PARAM(bufferObject, CL_MEM_OFFSET, offset, size_t(0), + "offset", "%zu", size_t) cl_buffer_region region; region.origin = addressAlign; @@ -321,7 +326,8 @@ int test_get_buffer_info( cl_device_id deviceID, cl_context context, cl_command_ } TEST_MEM_OBJECT_PARAM( subBufferObject, CL_MEM_FLAGS, flags, (unsigned int)inheritedFlags, "flags", "%d", unsigned int ) - TEST_MEM_OBJECT_PARAM( subBufferObject, CL_MEM_SIZE, sz, (size_t)( addressAlign ), "size", "%ld", size_t ) + TEST_MEM_OBJECT_PARAM(subBufferObject, CL_MEM_SIZE, sz, + (size_t)(addressAlign), "size", "%zu", size_t) if ( bufferFlags[ i ] & CL_MEM_USE_HOST_PTR ) { @@ -356,7 +362,9 @@ int test_get_buffer_info( cl_device_id deviceID, cl_context context, cl_command_ TEST_MEM_OBJECT_PARAM( subBufferObject, CL_MEM_ASSOCIATED_MEMOBJECT, origObj, (cl_mem)bufferObject, "associated mem object", "%p", void * ) - TEST_MEM_OBJECT_PARAM( subBufferObject, CL_MEM_OFFSET, offset, (size_t)( addressAlign ), "offset", "%ld", size_t ) + TEST_MEM_OBJECT_PARAM(subBufferObject, CL_MEM_OFFSET, offset, + (size_t)(addressAlign), "offset", "%zu", + size_t) } } @@ -405,7 +413,8 @@ int test_get_imageObject_info( cl_mem * image, cl_mem_flags objectFlags, cl_imag TEST_MEM_OBJECT_PARAM( *image, CL_MEM_CONTEXT, otherCtx, context, "context", "%p", cl_context ) - TEST_MEM_OBJECT_PARAM( *image, CL_MEM_OFFSET, offset, 0L, "offset", "%ld", size_t ) + TEST_MEM_OBJECT_PARAM(*image, CL_MEM_OFFSET, offset, size_t(0), "offset", + "%zu", size_t) return CL_SUCCESS; } diff --git a/test_conformance/api/test_mem_object_properties_queries.cpp b/test_conformance/api/test_mem_object_properties_queries.cpp index 55300a62fb..7a5cb0cea2 100644 --- a/test_conformance/api/test_mem_object_properties_queries.cpp +++ b/test_conformance/api/test_mem_object_properties_queries.cpp @@ -165,7 +165,7 @@ static int create_object_and_check_properties(cl_context context, } if (set_size != test_case.properties.size() * sizeof(cl_mem_properties)) { - log_error("ERROR: CL_MEM_PROPERTIES size is %d, expected %d.\n", + log_error("ERROR: CL_MEM_PROPERTIES size is %zu, expected %zu.\n", set_size, test_case.properties.size() * sizeof(cl_queue_properties)); return TEST_FAIL; diff --git a/test_conformance/api/test_pipe_properties_queries.cpp b/test_conformance/api/test_pipe_properties_queries.cpp index db918952b8..099bb2ca22 100644 --- a/test_conformance/api/test_pipe_properties_queries.cpp +++ b/test_conformance/api/test_pipe_properties_queries.cpp @@ -60,7 +60,7 @@ static int create_pipe_and_check_array_properties( } if (set_size != test_case.properties.size() * sizeof(cl_pipe_properties)) { - log_error("ERROR: CL_PIPE_PROPERTIES size is %d, expected %d.\n", + log_error("ERROR: CL_PIPE_PROPERTIES size is %zu, expected %zu.\n", set_size, test_case.properties.size() * sizeof(cl_pipe_properties)); return TEST_FAIL; diff --git a/test_conformance/api/test_platform.cpp b/test_conformance/api/test_platform.cpp index 57a3d7534a..841612a753 100644 --- a/test_conformance/api/test_platform.cpp +++ b/test_conformance/api/test_platform.cpp @@ -189,8 +189,10 @@ int test_get_platform_ids(cl_device_id deviceID, cl_context context, cl_command_ test_error(err, "clGetPlatformInfo for CL_PLATFORM_PROFILE failed"); log_info("\tCL_PLATFORM_PROFILE: %s\n", string_returned); if (strlen(string_returned)+1 != size) { - log_error("Returned string length %ld does not equal reported one %ld.\n", strlen(string_returned)+1, size); - total_errors++; + log_error( + "Returned string length %zu does not equal reported one %zu.\n", + strlen(string_returned) + 1, size); + total_errors++; } memset(string_returned, 0, 8192); @@ -198,8 +200,10 @@ int test_get_platform_ids(cl_device_id deviceID, cl_context context, cl_command_ test_error(err, "clGetPlatformInfo for CL_PLATFORM_VERSION failed"); log_info("\tCL_PLATFORM_VERSION: %s\n", string_returned); if (strlen(string_returned)+1 != size) { - log_error("Returned string length %ld does not equal reported one %ld.\n", strlen(string_returned)+1, size); - total_errors++; + log_error( + "Returned string length %zu does not equal reported one %zu.\n", + strlen(string_returned) + 1, size); + total_errors++; } memset(string_returned, 0, 8192); @@ -207,8 +211,10 @@ int test_get_platform_ids(cl_device_id deviceID, cl_context context, cl_command_ test_error(err, "clGetPlatformInfo for CL_PLATFORM_NAME failed"); log_info("\tCL_PLATFORM_NAME: %s\n", string_returned); if (strlen(string_returned)+1 != size) { - log_error("Returned string length %ld does not equal reported one %ld.\n", strlen(string_returned)+1, size); - total_errors++; + log_error( + "Returned string length %zu does not equal reported one %zu.\n", + strlen(string_returned) + 1, size); + total_errors++; } memset(string_returned, 0, 8192); @@ -216,8 +222,10 @@ int test_get_platform_ids(cl_device_id deviceID, cl_context context, cl_command_ test_error(err, "clGetPlatformInfo for CL_PLATFORM_VENDOR failed"); log_info("\tCL_PLATFORM_VENDOR: %s\n", string_returned); if (strlen(string_returned)+1 != size) { - log_error("Returned string length %ld does not equal reported one %ld.\n", strlen(string_returned)+1, size); - total_errors++; + log_error( + "Returned string length %zu does not equal reported one %zu.\n", + strlen(string_returned) + 1, size); + total_errors++; } memset(string_returned, 0, 8192); @@ -225,8 +233,10 @@ int test_get_platform_ids(cl_device_id deviceID, cl_context context, cl_command_ test_error(err, "clGetPlatformInfo for CL_PLATFORM_EXTENSIONS failed"); log_info("\tCL_PLATFORM_EXTENSIONS: %s\n", string_returned); if (strlen(string_returned)+1 != size) { - log_error("Returned string length %ld does not equal reported one %ld.\n", strlen(string_returned)+1, size); - total_errors++; + log_error( + "Returned string length %zu does not equal reported one %zu.\n", + strlen(string_returned) + 1, size); + total_errors++; } err = clGetDeviceIDs(platforms[p], CL_DEVICE_TYPE_ALL, 0, NULL, &num_devices); @@ -258,8 +268,8 @@ int test_get_platform_ids(cl_device_id deviceID, cl_context context, cl_command_ test_error(err, "clGetDeviceInfo failed for CL_DEVICE_PLATFORM\n"); if (returned_size != sizeof(cl_platform_id)) { - log_error("Reported return size (%ld) does not match expected size " - "(%ld).\n", + log_error("Reported return size (%zu) does not match expected size " + "(%zu).\n", returned_size, sizeof(cl_platform_id)); total_errors++; } @@ -288,7 +298,7 @@ int test_get_platform_ids(cl_device_id deviceID, cl_context context, cl_command_ if (returned_size != sizeof(cl_context_properties) * 3) { log_error("Invalid size returned from clGetContextInfo for " - "CL_CONTEXT_PROPERTIES. Got %ld, expected %ld.\n", + "CL_CONTEXT_PROPERTIES. Got %zu, expected %zu.\n", returned_size, sizeof(cl_context_properties) * 3); total_errors++; } diff --git a/test_conformance/api/test_queries.cpp b/test_conformance/api/test_queries.cpp index a1d8c0218e..c82017ba60 100644 --- a/test_conformance/api/test_queries.cpp +++ b/test_conformance/api/test_queries.cpp @@ -19,6 +19,7 @@ #include #include #include +#include #include int test_get_platform_info(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) @@ -237,7 +238,7 @@ int test_sampler_params(cl_device_id deviceID, cl_context context, if (set_size != 0) { log_error( - "ERROR: CL_SAMPLER_PROPERTIES size is %d, expected 0\n", + "ERROR: CL_SAMPLER_PROPERTIES size is %zu, expected 0\n", set_size); return TEST_FAIL; } @@ -248,7 +249,7 @@ int test_sampler_params(cl_device_id deviceID, cl_context context, != test_properties.size() * sizeof(cl_sampler_properties)) { log_error( - "ERROR: CL_SAMPLER_PROPERTIES size is %d, expected %d.\n", + "ERROR: CL_SAMPLER_PROPERTIES size is %zu, expected %zu.\n", set_size, test_properties.size() * sizeof(cl_sampler_properties)); return TEST_FAIL; @@ -380,7 +381,8 @@ int check_get_command_queue_info_params(cl_device_id deviceID, clGetDeviceInfo(deviceID, host_queue_query, sizeof(host_queue_props), &host_queue_props, NULL); test_error(error, "clGetDeviceInfo failed"); - log_info("CL_DEVICE_QUEUE_ON_HOST_PROPERTIES is %d\n", host_queue_props); + log_info("CL_DEVICE_QUEUE_ON_HOST_PROPERTIES is %" PRIu64 "\n", + host_queue_props); cl_queue_properties device_queue_props = 0; if (version >= Version(2, 0)) @@ -389,7 +391,7 @@ int check_get_command_queue_info_params(cl_device_id deviceID, sizeof(device_queue_props), &device_queue_props, NULL); test_error(error, "clGetDeviceInfo failed"); - log_info("CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES is %d\n", + log_info("CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES is %" PRIu64 "\n", device_queue_props); } diff --git a/test_conformance/api/test_queue_properties_queries.cpp b/test_conformance/api/test_queue_properties_queries.cpp index 03dbb2c5df..6d9b438027 100644 --- a/test_conformance/api/test_queue_properties_queries.cpp +++ b/test_conformance/api/test_queue_properties_queries.cpp @@ -142,9 +142,10 @@ static int create_queue_and_check_array_properties( } if (set_size != test_case.properties.size() * sizeof(cl_queue_properties)) { - log_error("ERROR: CL_QUEUE_PROPERTIES_ARRAY size is %d, expected %d.\n", - set_size, - test_case.properties.size() * sizeof(cl_queue_properties)); + log_error( + "ERROR: CL_QUEUE_PROPERTIES_ARRAY size is %zu, expected %zu.\n", + set_size, + test_case.properties.size() * sizeof(cl_queue_properties)); return TEST_FAIL; } From 7131f879740f775e6fdeaf802fa84529dd985262 Mon Sep 17 00:00:00 2001 From: saurabhnv <156190705+saurabhnv@users.noreply.github.com> Date: Tue, 3 Sep 2024 22:48:39 +0530 Subject: [PATCH 5/7] Fix external semaphore test when exportability is not supported (#2045) An implementation may not support exportable semaphore, subtests available in cl_khr_external_semaphore assumes support for exportable semaphore, resulting in failure on such implementation. Allow implementations to use importable semaphore in such cases. --- .../test_external_semaphore.cpp | 194 +++++++++++------- 1 file changed, 115 insertions(+), 79 deletions(-) diff --git a/test_conformance/extensions/cl_khr_external_semaphore/test_external_semaphore.cpp b/test_conformance/extensions/cl_khr_external_semaphore/test_external_semaphore.cpp index af0314c1df..7d21284973 100644 --- a/test_conformance/extensions/cl_khr_external_semaphore/test_external_semaphore.cpp +++ b/test_conformance/extensions/cl_khr_external_semaphore/test_external_semaphore.cpp @@ -4,7 +4,7 @@ #include "opencl_vulkan_wrapper.hpp" #include #include -#include +#include #define FLUSH_DELAY_S 5 @@ -96,18 +96,26 @@ static cl_int get_device_semaphore_handle_types( num_handle_types = size_handle_types / sizeof(cl_external_semaphore_handle_type_khr); - std::vector - handle_types_query_result(num_handle_types); - err = clGetDeviceInfo(deviceID, param, - handle_types_query_result.size() - * sizeof(cl_external_semaphore_handle_type_khr), - handle_types_query_result.data(), nullptr); - test_error(err, "Failed to get exportable handle types"); - - for (auto handle_type : handle_types_query_result) + + // Empty list (size_handle_types:0) is a valid value denoting that + // implementation is incapable to import/export semaphore + if (num_handle_types > 0) { - handle_types.push_back(handle_type); + std::vector + handle_types_query_result(num_handle_types); + err = + clGetDeviceInfo(deviceID, param, + handle_types_query_result.size() + * sizeof(cl_external_semaphore_handle_type_khr), + handle_types_query_result.data(), nullptr); + test_error(err, "Failed to get exportable handle types"); + + for (auto handle_type : handle_types_query_result) + { + handle_types.push_back(handle_type); + } } + return CL_SUCCESS; } @@ -231,16 +239,21 @@ int test_external_semaphores_cross_context(cl_device_id deviceID, export_handle_types); test_error(err, "Failed to query export handle types"); + // If cl_khr_external_semaphore is reported, implementation must + // support any of import, export or maybe both. + if (import_handle_types.empty() && export_handle_types.empty()) + { + test_fail("No support for import/export semaphore.\n"); + } + // Find handles that support both import and export - std::unordered_set + std::vector import_export_handle_types; - std::copy(import_handle_types.begin(), import_handle_types.end(), - std::inserter(import_export_handle_types, - import_export_handle_types.end())); - std::copy(export_handle_types.begin(), export_handle_types.end(), - std::inserter(import_export_handle_types, - import_export_handle_types.end())); + std::set_intersection( + import_handle_types.begin(), import_handle_types.end(), + export_handle_types.begin(), export_handle_types.end(), + std::back_inserter(import_export_handle_types)); cl_context context2 = clCreateContext(NULL, 1, &deviceID, notify_callback, NULL, &err); @@ -257,7 +270,7 @@ int test_external_semaphores_cross_context(cl_device_id deviceID, if (import_export_handle_types.empty()) { log_info("Could not find a handle type that supports both import and " - "export"); + "export.\n"); return TEST_SKIPPED_ITSELF; } @@ -366,8 +379,10 @@ int test_external_semaphores_simple_1(cl_device_id deviceID, cl_context context, VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); - clExternalExportableSemaphore sema_ext( - vkVk2CLSemaphore, context, vkExternalSemaphoreHandleType, deviceID); + clExternalSemaphore *raw_sema_ext = NULL; + CREATE_OPENCL_SEMAPHORE(raw_sema_ext, vkVk2CLSemaphore, context, + vkExternalSemaphoreHandleType, deviceID, true); + std::unique_ptr sema_ext(raw_sema_ext); cl_int err = CL_SUCCESS; @@ -378,13 +393,14 @@ int test_external_semaphores_simple_1(cl_device_id deviceID, cl_context context, // Signal semaphore clEventWrapper signal_event; - err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema_ext.getCLSemaphore(), - nullptr, 0, nullptr, &signal_event); + err = + clEnqueueSignalSemaphoresKHR(queue, 1, &sema_ext->getCLSemaphore(), + nullptr, 0, nullptr, &signal_event); test_error(err, "Could not signal semaphore"); // Wait semaphore clEventWrapper wait_event; - err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_ext.getCLSemaphore(), + err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_ext->getCLSemaphore(), nullptr, 0, nullptr, &wait_event); test_error(err, "Could not wait semaphore"); @@ -442,8 +458,10 @@ int test_external_semaphores_simple_2(cl_device_id deviceID, cl_context context, VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); - clExternalExportableSemaphore sema_ext( - vkVk2CLSemaphore, context, vkExternalSemaphoreHandleType, deviceID); + clExternalSemaphore *raw_sema_ext = NULL; + CREATE_OPENCL_SEMAPHORE(raw_sema_ext, vkVk2CLSemaphore, context, + vkExternalSemaphoreHandleType, deviceID, true); + std::unique_ptr sema_ext(raw_sema_ext); cl_int err = CL_SUCCESS; @@ -470,13 +488,14 @@ int test_external_semaphores_simple_2(cl_device_id deviceID, cl_context context, // Signal semaphore clEventWrapper signal_event; - err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema_ext.getCLSemaphore(), - nullptr, 0, nullptr, &signal_event); + err = + clEnqueueSignalSemaphoresKHR(queue, 1, &sema_ext->getCLSemaphore(), + nullptr, 0, nullptr, &signal_event); test_error(err, "Could not signal semaphore"); // Wait semaphore clEventWrapper wait_event; - err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_ext.getCLSemaphore(), + err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_ext->getCLSemaphore(), nullptr, 0, nullptr, &wait_event); test_error(err, "Could not wait semaphore"); @@ -548,8 +567,10 @@ int test_external_semaphores_reuse(cl_device_id deviceID, cl_context context, VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); - clExternalExportableSemaphore sema_ext( - vkVk2CLSemaphore, context, vkExternalSemaphoreHandleType, deviceID); + clExternalSemaphore *raw_sema_ext = NULL; + CREATE_OPENCL_SEMAPHORE(raw_sema_ext, vkVk2CLSemaphore, context, + vkExternalSemaphoreHandleType, deviceID, true); + std::unique_ptr sema_ext(raw_sema_ext); cl_int err = CL_SUCCESS; @@ -575,9 +596,9 @@ int test_external_semaphores_reuse(cl_device_id deviceID, cl_context context, test_error(err, "Unable to enqueue task_1"); // Signal semaphore (dependency on task_1) - err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema_ext.getCLSemaphore(), - nullptr, 1, &task_events[0], - &signal_events[0]); + err = clEnqueueSignalSemaphoresKHR( + queue, 1, &sema_ext->getCLSemaphore(), nullptr, 1, &task_events[0], + &signal_events[0]); test_error(err, "Could not signal semaphore"); // In a loop @@ -586,7 +607,7 @@ int test_external_semaphores_reuse(cl_device_id deviceID, cl_context context, { // Wait semaphore err = clEnqueueWaitSemaphoresKHR( - queue, 1, &sema_ext.getCLSemaphore(), nullptr, 0, nullptr, + queue, 1, &sema_ext->getCLSemaphore(), nullptr, 0, nullptr, &wait_events[loop - 1]); test_error(err, "Could not wait semaphore"); @@ -601,13 +622,13 @@ int test_external_semaphores_reuse(cl_device_id deviceID, cl_context context, // Signal semaphore (dependency on task_loop) err = clEnqueueSignalSemaphoresKHR( - queue, 1, &sema_ext.getCLSemaphore(), nullptr, 1, + queue, 1, &sema_ext->getCLSemaphore(), nullptr, 1, &task_events[loop], &signal_events[loop]); test_error(err, "Could not signal semaphore"); } // Wait semaphore - err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_ext.getCLSemaphore(), + err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_ext->getCLSemaphore(), nullptr, 0, nullptr, &wait_events[loop - 1]); test_error(err, "Could not wait semaphore"); @@ -671,22 +692,25 @@ static int external_semaphore_cross_queue_helper(cl_device_id deviceID, VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); - clExternalExportableSemaphore sema_ext( - vkVk2CLSemaphore, context, vkExternalSemaphoreHandleType, deviceID); + clExternalSemaphore *raw_sema_ext = NULL; + CREATE_OPENCL_SEMAPHORE(raw_sema_ext, vkVk2CLSemaphore, context, + vkExternalSemaphoreHandleType, deviceID, true); + std::unique_ptr sema_ext(raw_sema_ext); cl_int err = CL_SUCCESS; // Signal semaphore on queue_1 clEventWrapper signal_event; - err = - clEnqueueSignalSemaphoresKHR(queue_1, 1, &sema_ext.getCLSemaphore(), - nullptr, 0, nullptr, &signal_event); + err = clEnqueueSignalSemaphoresKHR(queue_1, 1, + &sema_ext->getCLSemaphore(), nullptr, + 0, nullptr, &signal_event); test_error(err, "Could not signal semaphore"); // Wait semaphore on queue_2 clEventWrapper wait_event; - err = clEnqueueWaitSemaphoresKHR(queue_2, 1, &sema_ext.getCLSemaphore(), - nullptr, 0, nullptr, &wait_event); + err = + clEnqueueWaitSemaphoresKHR(queue_2, 1, &sema_ext->getCLSemaphore(), + nullptr, 0, nullptr, &wait_event); test_error(err, "Could not wait semaphore"); // Finish queue_1 and queue_2 @@ -797,11 +821,15 @@ int test_external_semaphores_cross_queues_io2(cl_device_id deviceID, VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); - clExternalExportableSemaphore sema_ext_1( - vkVk2CLSemaphore, context, vkExternalSemaphoreHandleType, deviceID); - clExternalExportableSemaphore sema_ext_2(vkVk2CLSemaphore, context2, - vkExternalSemaphoreHandleType, - deviceID); + clExternalSemaphore *raw_sema_ext_1 = NULL; + CREATE_OPENCL_SEMAPHORE(raw_sema_ext_1, vkVk2CLSemaphore, context, + vkExternalSemaphoreHandleType, deviceID, true); + std::unique_ptr sema_ext_1(raw_sema_ext_1); + + clExternalSemaphore *raw_sema_ext_2 = NULL; + CREATE_OPENCL_SEMAPHORE(raw_sema_ext_2, vkVk2CLSemaphore, context2, + vkExternalSemaphoreHandleType, deviceID, true); + std::unique_ptr sema_ext_2(raw_sema_ext_2); clCommandQueueWrapper queue1 = clCreateCommandQueue(context, deviceID, 0, &err); @@ -814,28 +842,28 @@ int test_external_semaphores_cross_queues_io2(cl_device_id deviceID, // Signal semaphore 1 clEventWrapper signal_1_event; err = clEnqueueSignalSemaphoresKHR( - queue1, 1, &sema_ext_1.getCLSemaphore(), nullptr, 0, nullptr, + queue1, 1, &sema_ext_1->getCLSemaphore(), nullptr, 0, nullptr, &signal_1_event); test_error(err, "Could not signal semaphore"); // Wait semaphore 1 clEventWrapper wait_1_event; err = - clEnqueueWaitSemaphoresKHR(queue1, 1, &sema_ext_1.getCLSemaphore(), + clEnqueueWaitSemaphoresKHR(queue1, 1, &sema_ext_1->getCLSemaphore(), nullptr, 0, nullptr, &wait_1_event); test_error(err, "Could not wait semaphore"); // Signal semaphore 2 clEventWrapper signal_2_event; err = clEnqueueSignalSemaphoresKHR( - queue2, 1, &sema_ext_2.getCLSemaphore(), nullptr, 0, nullptr, + queue2, 1, &sema_ext_2->getCLSemaphore(), nullptr, 0, nullptr, &signal_2_event); test_error(err, "Could not signal semaphore"); // Wait semaphore 2 clEventWrapper wait_2_event; err = - clEnqueueWaitSemaphoresKHR(queue2, 1, &sema_ext_2.getCLSemaphore(), + clEnqueueWaitSemaphoresKHR(queue2, 1, &sema_ext_2->getCLSemaphore(), nullptr, 0, nullptr, &wait_2_event); test_error(err, "Could not wait semaphore"); @@ -900,12 +928,15 @@ int test_external_semaphores_multi_signal(cl_device_id deviceID, VulkanSemaphore vkVk2CLSemaphore2(vkDevice, vkExternalSemaphoreHandleType); - clExternalExportableSemaphore sema_ext_1(vkVk2CLSemaphore1, context, - vkExternalSemaphoreHandleType, - deviceID); - clExternalExportableSemaphore sema_ext_2(vkVk2CLSemaphore2, context, - vkExternalSemaphoreHandleType, - deviceID); + clExternalSemaphore *raw_sema_ext_1 = NULL; + CREATE_OPENCL_SEMAPHORE(raw_sema_ext_1, vkVk2CLSemaphore1, context, + vkExternalSemaphoreHandleType, deviceID, true); + std::unique_ptr sema_ext_1(raw_sema_ext_1); + + clExternalSemaphore *raw_sema_ext_2 = NULL; + CREATE_OPENCL_SEMAPHORE(raw_sema_ext_2, vkVk2CLSemaphore2, context, + vkExternalSemaphoreHandleType, deviceID, true); + std::unique_ptr sema_ext_2(raw_sema_ext_2); cl_int err = CL_SUCCESS; @@ -916,22 +947,24 @@ int test_external_semaphores_multi_signal(cl_device_id deviceID, // Signal semaphore 1 and 2 clEventWrapper signal_event; - cl_semaphore_khr sema_list[] = { sema_ext_1.getCLSemaphore(), - sema_ext_2.getCLSemaphore() }; + cl_semaphore_khr sema_list[] = { sema_ext_1->getCLSemaphore(), + sema_ext_2->getCLSemaphore() }; err = clEnqueueSignalSemaphoresKHR(queue, 2, sema_list, nullptr, 0, nullptr, &signal_event); test_error(err, "Could not signal semaphore"); // Wait semaphore 1 clEventWrapper wait_1_event; - err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_ext_1.getCLSemaphore(), - nullptr, 0, nullptr, &wait_1_event); + err = + clEnqueueWaitSemaphoresKHR(queue, 1, &sema_ext_1->getCLSemaphore(), + nullptr, 0, nullptr, &wait_1_event); test_error(err, "Could not wait semaphore"); // Wait semaphore 2 clEventWrapper wait_2_event; - err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_ext_2.getCLSemaphore(), - nullptr, 0, nullptr, &wait_2_event); + err = + clEnqueueWaitSemaphoresKHR(queue, 1, &sema_ext_2->getCLSemaphore(), + nullptr, 0, nullptr, &wait_2_event); test_error(err, "Could not wait semaphore"); // Finish @@ -991,12 +1024,15 @@ int test_external_semaphores_multi_wait(cl_device_id deviceID, VulkanSemaphore vkVk2CLSemaphore2(vkDevice, vkExternalSemaphoreHandleType); - clExternalExportableSemaphore sema_ext_1(vkVk2CLSemaphore1, context, - vkExternalSemaphoreHandleType, - deviceID); - clExternalExportableSemaphore sema_ext_2(vkVk2CLSemaphore2, context, - vkExternalSemaphoreHandleType, - deviceID); + clExternalSemaphore *raw_sema_ext_1 = NULL; + CREATE_OPENCL_SEMAPHORE(raw_sema_ext_1, vkVk2CLSemaphore1, context, + vkExternalSemaphoreHandleType, deviceID, true); + std::unique_ptr sema_ext_1(raw_sema_ext_1); + + clExternalSemaphore *raw_sema_ext_2 = NULL; + CREATE_OPENCL_SEMAPHORE(raw_sema_ext_2, vkVk2CLSemaphore2, context, + vkExternalSemaphoreHandleType, deviceID, true); + std::unique_ptr sema_ext_2(raw_sema_ext_2); cl_int err = CL_SUCCESS; @@ -1007,22 +1043,22 @@ int test_external_semaphores_multi_wait(cl_device_id deviceID, // Signal semaphore 1 clEventWrapper signal_1_event; - err = - clEnqueueSignalSemaphoresKHR(queue, 1, &sema_ext_1.getCLSemaphore(), - nullptr, 0, nullptr, &signal_1_event); + err = clEnqueueSignalSemaphoresKHR( + queue, 1, &sema_ext_1->getCLSemaphore(), nullptr, 0, nullptr, + &signal_1_event); test_error(err, "Could not signal semaphore"); // Signal semaphore 2 clEventWrapper signal_2_event; - err = - clEnqueueSignalSemaphoresKHR(queue, 1, &sema_ext_2.getCLSemaphore(), - nullptr, 0, nullptr, &signal_2_event); + err = clEnqueueSignalSemaphoresKHR( + queue, 1, &sema_ext_2->getCLSemaphore(), nullptr, 0, nullptr, + &signal_2_event); test_error(err, "Could not signal semaphore"); // Wait semaphore 1 and 2 clEventWrapper wait_event; - cl_semaphore_khr sema_list[] = { sema_ext_1.getCLSemaphore(), - sema_ext_2.getCLSemaphore() }; + cl_semaphore_khr sema_list[] = { sema_ext_1->getCLSemaphore(), + sema_ext_2->getCLSemaphore() }; err = clEnqueueWaitSemaphoresKHR(queue, 2, sema_list, nullptr, 0, nullptr, &wait_event); test_error(err, "Could not wait semaphore"); From eb7a30ae42fb8a66e670d10b96dba168d707d66f Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Tue, 3 Sep 2024 19:23:22 +0200 Subject: [PATCH 6/7] Modernization of tests from test_semaphores.cpp to align with new SemaphoreTestBase infrastructure (#2029) --- .../cl_khr_semaphore/test_semaphores.cpp | 1212 ++++++++--------- 1 file changed, 599 insertions(+), 613 deletions(-) diff --git a/test_conformance/extensions/cl_khr_semaphore/test_semaphores.cpp b/test_conformance/extensions/cl_khr_semaphore/test_semaphores.cpp index ec4f752b0b..f4f2832414 100644 --- a/test_conformance/extensions/cl_khr_semaphore/test_semaphores.cpp +++ b/test_conformance/extensions/cl_khr_semaphore/test_semaphores.cpp @@ -16,13 +16,14 @@ #include "harness/typeWrappers.h" -#include "harness/extensionHelpers.h" #include "harness/errorHelpers.h" #include #include #include #include +#include "semaphore_base.h" + #define FLUSH_DELAY_S 5 #define SEMAPHORE_PARAM_TEST(param_name, param_type, expected) \ @@ -30,8 +31,8 @@ { \ param_type value; \ size_t size; \ - cl_int error = clGetSemaphoreInfoKHR(sema, param_name, sizeof(value), \ - &value, &size); \ + cl_int error = clGetSemaphoreInfoKHR(semaphore, param_name, \ + sizeof(value), &value, &size); \ test_error(error, "Unable to get " #param_name " from semaphore"); \ if (value != expected) \ { \ @@ -54,8 +55,8 @@ { \ param_type value[num_params]; \ size_t size; \ - cl_int error = clGetSemaphoreInfoKHR(sema, param_name, sizeof(value), \ - &value, &size); \ + cl_int error = clGetSemaphoreInfoKHR(semaphore, param_name, \ + sizeof(value), &value, &size); \ test_error(error, "Unable to get " #param_name " from semaphore"); \ if (size != sizeof(value)) \ { \ @@ -70,326 +71,625 @@ } \ } while (false) -static const char* source = "__kernel void empty() {}"; +namespace { + +const char* source = "__kernel void empty() {}"; -// Helper function that signals and waits on semaphore across two different -// queues. -static int semaphore_cross_queue_helper(cl_device_id deviceID, - cl_context context, - cl_command_queue queue_1, - cl_command_queue queue_2) +struct SimpleSemaphore1 : public SemaphoreTestBase { - cl_int err; + SimpleSemaphore1(cl_device_id device, cl_context context, + cl_command_queue queue) + : SemaphoreTestBase(device, context, queue) + {} - if (!is_extension_available(deviceID, "cl_khr_semaphore")) + cl_int Run() override { - log_info("cl_khr_semaphore is not supported on this platform. " - "Skipping test.\n"); - return TEST_SKIPPED_ITSELF; - } + cl_int err = CL_SUCCESS; + // Create ooo queue + clCommandQueueWrapper queue = clCreateCommandQueue( + context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); + test_error(err, "Could not create command queue"); + + // Create semaphore + cl_semaphore_properties_khr sema_props[] = { + static_cast(CL_SEMAPHORE_TYPE_KHR), + static_cast( + CL_SEMAPHORE_TYPE_BINARY_KHR), + 0 + }; + semaphore = + clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err); + test_error(err, "Could not create semaphore"); + + // Signal semaphore + clEventWrapper signal_event; + err = clEnqueueSignalSemaphoresKHR(queue, 1, semaphore, nullptr, 0, + nullptr, &signal_event); + test_error(err, "Could not signal semaphore"); - // Obtain pointers to semaphore's API - GET_PFN(deviceID, clCreateSemaphoreWithPropertiesKHR); - GET_PFN(deviceID, clEnqueueSignalSemaphoresKHR); - GET_PFN(deviceID, clEnqueueWaitSemaphoresKHR); - GET_PFN(deviceID, clReleaseSemaphoreKHR); - - // Create semaphore - cl_semaphore_properties_khr sema_props[] = { - static_cast(CL_SEMAPHORE_TYPE_KHR), - static_cast(CL_SEMAPHORE_TYPE_BINARY_KHR), - 0 - }; - cl_semaphore_khr sema = - clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err); - test_error(err, "Could not create semaphore"); - - // Signal semaphore on queue_1 - clEventWrapper signal_event; - err = clEnqueueSignalSemaphoresKHR(queue_1, 1, &sema, nullptr, 0, nullptr, - &signal_event); - test_error(err, "Could not signal semaphore"); - - // Wait semaphore on queue_2 - clEventWrapper wait_event; - err = clEnqueueWaitSemaphoresKHR(queue_2, 1, &sema, nullptr, 0, nullptr, - &wait_event); - test_error(err, "Could not wait semaphore"); - - // Finish queue_1 and queue_2 - err = clFinish(queue_1); - test_error(err, "Could not finish queue"); - - err = clFinish(queue_2); - test_error(err, "Could not finish queue"); - - // Ensure all events are completed - test_assert_event_complete(signal_event); - test_assert_event_complete(wait_event); - - // Release semaphore - err = clReleaseSemaphoreKHR(sema); - test_error(err, "Could not release semaphore"); - - return TEST_PASS; -} + // Wait semaphore + clEventWrapper wait_event; + err = clEnqueueWaitSemaphoresKHR(queue, 1, semaphore, nullptr, 0, + nullptr, &wait_event); + test_error(err, "Could not wait semaphore"); -// Confirm that a signal followed by a wait will complete successfully -int test_semaphores_simple_1(cl_device_id deviceID, cl_context context, - cl_command_queue defaultQueue, int num_elements) + // Finish + err = clFinish(queue); + test_error(err, "Could not finish queue"); + + // Ensure all events are completed + test_assert_event_complete(signal_event); + test_assert_event_complete(wait_event); + + return CL_SUCCESS; + } +}; + +struct SimpleSemaphore2 : public SemaphoreTestBase { - cl_int err; + SimpleSemaphore2(cl_device_id device, cl_context context, + cl_command_queue queue) + : SemaphoreTestBase(device, context, queue) + {} - if (!is_extension_available(deviceID, "cl_khr_semaphore")) + cl_int Run() override { - log_info("cl_khr_semaphore is not supported on this platform. " - "Skipping test.\n"); - return TEST_SKIPPED_ITSELF; + cl_int err = CL_SUCCESS; + // Create ooo queue + clCommandQueueWrapper queue = clCreateCommandQueue( + context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); + test_error(err, "Could not create command queue"); + + // Create semaphore + cl_semaphore_properties_khr sema_props[] = { + static_cast(CL_SEMAPHORE_TYPE_KHR), + static_cast( + CL_SEMAPHORE_TYPE_BINARY_KHR), + 0 + }; + semaphore = + clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err); + test_error(err, "Could not create semaphore"); + + // Create user event + clEventWrapper user_event = clCreateUserEvent(context, &err); + test_error(err, "Could not create user event"); + + // Create Kernel + clProgramWrapper program; + clKernelWrapper kernel; + err = create_single_kernel_helper(context, &program, &kernel, 1, + &source, "empty"); + test_error(err, "Could not create kernel"); + + // Enqueue task_1 (dependency on user_event) + clEventWrapper task_1_event; + err = clEnqueueTask(queue, kernel, 1, &user_event, &task_1_event); + test_error(err, "Could not enqueue task 1"); + + // Signal semaphore + clEventWrapper signal_event; + err = clEnqueueSignalSemaphoresKHR(queue, 1, semaphore, nullptr, 0, + nullptr, &signal_event); + test_error(err, "Could not signal semaphore"); + + // Wait semaphore + clEventWrapper wait_event; + err = clEnqueueWaitSemaphoresKHR(queue, 1, semaphore, nullptr, 0, + nullptr, &wait_event); + test_error(err, "Could not wait semaphore"); + + // Flush and delay + err = clFlush(queue); + test_error(err, "Could not flush queue"); + std::this_thread::sleep_for(std::chrono::seconds(FLUSH_DELAY_S)); + + // Ensure all events are completed except for task_1 + test_assert_event_inprogress(task_1_event); + test_assert_event_complete(signal_event); + test_assert_event_complete(wait_event); + + // Complete user_event + err = clSetUserEventStatus(user_event, CL_COMPLETE); + test_error(err, "Could not set user event to CL_COMPLETE"); + + // Finish + err = clFinish(queue); + test_error(err, "Could not finish queue"); + + // Ensure all events are completed + test_assert_event_complete(task_1_event); + test_assert_event_complete(signal_event); + test_assert_event_complete(wait_event); + + return CL_SUCCESS; } +}; - // Obtain pointers to semaphore's API - GET_PFN(deviceID, clCreateSemaphoreWithPropertiesKHR); - GET_PFN(deviceID, clEnqueueSignalSemaphoresKHR); - GET_PFN(deviceID, clEnqueueWaitSemaphoresKHR); - GET_PFN(deviceID, clReleaseSemaphoreKHR); - - // Create ooo queue - clCommandQueueWrapper queue = clCreateCommandQueue( - context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); - test_error(err, "Could not create command queue"); - - // Create semaphore - cl_semaphore_properties_khr sema_props[] = { - static_cast(CL_SEMAPHORE_TYPE_KHR), - static_cast(CL_SEMAPHORE_TYPE_BINARY_KHR), - 0 - }; - cl_semaphore_khr sema = - clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err); - test_error(err, "Could not create semaphore"); - - // Signal semaphore - clEventWrapper signal_event; - err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema, nullptr, 0, nullptr, - &signal_event); - test_error(err, "Could not signal semaphore"); - - // Wait semaphore - clEventWrapper wait_event; - err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema, nullptr, 0, nullptr, - &wait_event); - test_error(err, "Could not wait semaphore"); - - // Finish - err = clFinish(queue); - test_error(err, "Could not finish queue"); - - // Ensure all events are completed - test_assert_event_complete(signal_event); - test_assert_event_complete(wait_event); - - // Release semaphore - err = clReleaseSemaphoreKHR(sema); - test_error(err, "Could not release semaphore"); - - return TEST_PASS; -} +struct SemaphoreReuse : public SemaphoreTestBase +{ + SemaphoreReuse(cl_device_id device, cl_context context, + cl_command_queue queue) + : SemaphoreTestBase(device, context, queue) + {} -// Confirm that signal a semaphore with no event dependencies will not result -// in an implicit dependency on everything previously submitted -int test_semaphores_simple_2(cl_device_id deviceID, cl_context context, - cl_command_queue defaultQueue, int num_elements) + cl_int Run() override + { + cl_int err = CL_SUCCESS; + // Create ooo queue + clCommandQueueWrapper queue = clCreateCommandQueue( + context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); + test_error(err, "Could not create command queue"); + + // Create semaphore + cl_semaphore_properties_khr sema_props[] = { + static_cast(CL_SEMAPHORE_TYPE_KHR), + static_cast( + CL_SEMAPHORE_TYPE_BINARY_KHR), + 0 + }; + semaphore = + clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err); + test_error(err, "Could not create semaphore"); + + // Create Kernel + clProgramWrapper program; + clKernelWrapper kernel; + err = create_single_kernel_helper(context, &program, &kernel, 1, + &source, "empty"); + test_error(err, "Could not create kernel"); + + constexpr size_t loop_count = 10; + clEventWrapper signal_events[loop_count]; + clEventWrapper wait_events[loop_count]; + clEventWrapper task_events[loop_count]; + + // Enqueue task_1 + err = clEnqueueTask(queue, kernel, 0, nullptr, &task_events[0]); + test_error(err, "Unable to enqueue task_1"); + + // Signal semaphore (dependency on task_1) + err = clEnqueueSignalSemaphoresKHR(queue, 1, semaphore, nullptr, 1, + &task_events[0], &signal_events[0]); + test_error(err, "Could not signal semaphore"); + + // In a loop + size_t loop; + for (loop = 1; loop < loop_count; ++loop) + { + // Wait semaphore + err = clEnqueueWaitSemaphoresKHR(queue, 1, semaphore, nullptr, 0, + nullptr, &wait_events[loop - 1]); + test_error(err, "Could not wait semaphore"); + + // Enqueue task_loop (dependency on wait) + err = clEnqueueTask(queue, kernel, 1, &wait_events[loop - 1], + &task_events[loop]); + test_error(err, "Unable to enqueue task_loop"); + + // Wait for the "wait semaphore" to complete + err = clWaitForEvents(1, &wait_events[loop - 1]); + test_error(err, "Unable to wait for wait semaphore to complete"); + + // Signal semaphore (dependency on task_loop) + err = clEnqueueSignalSemaphoresKHR(queue, 1, semaphore, nullptr, 1, + &task_events[loop], + &signal_events[loop]); + test_error(err, "Could not signal semaphore"); + } + + // Wait semaphore + err = clEnqueueWaitSemaphoresKHR(queue, 1, semaphore, nullptr, 0, + nullptr, &wait_events[loop - 1]); + test_error(err, "Could not wait semaphore"); + + // Finish + err = clFinish(queue); + test_error(err, "Could not finish queue"); + + // Ensure all events are completed + for (loop = 0; loop < loop_count; ++loop) + { + test_assert_event_complete(wait_events[loop]); + test_assert_event_complete(signal_events[loop]); + test_assert_event_complete(task_events[loop]); + } + + return CL_SUCCESS; + } +}; + +template struct SemaphoreCrossQueue : public SemaphoreTestBase { - cl_int err; + SemaphoreCrossQueue(cl_device_id device, cl_context context, + cl_command_queue queue) + : SemaphoreTestBase(device, context, queue) + {} + + // Helper function that signals and waits on semaphore across two different + // queues. + int semaphore_cross_queue_helper(cl_device_id deviceID, cl_context context, + cl_command_queue queue_1, + cl_command_queue queue_2) + { + cl_int err = CL_SUCCESS; + // Create semaphore + cl_semaphore_properties_khr sema_props[] = { + static_cast(CL_SEMAPHORE_TYPE_KHR), + static_cast( + CL_SEMAPHORE_TYPE_BINARY_KHR), + 0 + }; + semaphore = + clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err); + test_error(err, "Could not create semaphore"); + + // Signal semaphore on queue_1 + clEventWrapper signal_event; + err = clEnqueueSignalSemaphoresKHR(queue_1, 1, semaphore, nullptr, 0, + nullptr, &signal_event); + test_error(err, "Could not signal semaphore"); - if (!is_extension_available(deviceID, "cl_khr_semaphore")) + // Wait semaphore on queue_2 + clEventWrapper wait_event; + err = clEnqueueWaitSemaphoresKHR(queue_2, 1, semaphore, nullptr, 0, + nullptr, &wait_event); + test_error(err, "Could not wait semaphore"); + + // Finish queue_1 and queue_2 + err = clFinish(queue_1); + test_error(err, "Could not finish queue"); + + err = clFinish(queue_2); + test_error(err, "Could not finish queue"); + + // Ensure all events are completed + test_assert_event_complete(signal_event); + test_assert_event_complete(wait_event); + + return TEST_PASS; + } + + cl_int run_in_order() { - log_info("cl_khr_semaphore is not supported on this platform. " - "Skipping test.\n"); - return TEST_SKIPPED_ITSELF; + cl_int err = CL_SUCCESS; + // Create in-order queues + clCommandQueueWrapper queue_1 = + clCreateCommandQueue(context, device, 0, &err); + test_error(err, "Could not create command queue"); + + clCommandQueueWrapper queue_2 = + clCreateCommandQueue(context, device, 0, &err); + test_error(err, "Could not create command queue"); + + return semaphore_cross_queue_helper(device, context, queue_1, queue_2); } - // Obtain pointers to semaphore's API - GET_PFN(deviceID, clCreateSemaphoreWithPropertiesKHR); - GET_PFN(deviceID, clEnqueueSignalSemaphoresKHR); - GET_PFN(deviceID, clEnqueueWaitSemaphoresKHR); - GET_PFN(deviceID, clReleaseSemaphoreKHR); - - // Create ooo queue - clCommandQueueWrapper queue = clCreateCommandQueue( - context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); - test_error(err, "Could not create command queue"); - - // Create semaphore - cl_semaphore_properties_khr sema_props[] = { - static_cast(CL_SEMAPHORE_TYPE_KHR), - static_cast(CL_SEMAPHORE_TYPE_BINARY_KHR), - 0 - }; - cl_semaphore_khr sema = - clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err); - test_error(err, "Could not create semaphore"); - - // Create user event - clEventWrapper user_event = clCreateUserEvent(context, &err); - test_error(err, "Could not create user event"); - - // Create Kernel - clProgramWrapper program; - clKernelWrapper kernel; - err = create_single_kernel_helper(context, &program, &kernel, 1, &source, - "empty"); - test_error(err, "Could not create kernel"); - - // Enqueue task_1 (dependency on user_event) - clEventWrapper task_1_event; - err = clEnqueueTask(queue, kernel, 1, &user_event, &task_1_event); - test_error(err, "Could not enqueue task 1"); - - // Signal semaphore - clEventWrapper signal_event; - err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema, nullptr, 0, nullptr, - &signal_event); - test_error(err, "Could not signal semaphore"); - - // Wait semaphore - clEventWrapper wait_event; - err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema, nullptr, 0, nullptr, - &wait_event); - test_error(err, "Could not wait semaphore"); - - // Flush and delay - err = clFlush(queue); - test_error(err, "Could not flush queue"); - std::this_thread::sleep_for(std::chrono::seconds(FLUSH_DELAY_S)); - - // Ensure all events are completed except for task_1 - test_assert_event_inprogress(task_1_event); - test_assert_event_complete(signal_event); - test_assert_event_complete(wait_event); - - // Complete user_event - err = clSetUserEventStatus(user_event, CL_COMPLETE); - test_error(err, "Could not set user event to CL_COMPLETE"); - - // Finish - err = clFinish(queue); - test_error(err, "Could not finish queue"); - - // Ensure all events are completed - test_assert_event_complete(task_1_event); - test_assert_event_complete(signal_event); - test_assert_event_complete(wait_event); - - // Release semaphore - err = clReleaseSemaphoreKHR(sema); - test_error(err, "Could not release semaphore"); - - return TEST_PASS; -} + cl_int run_out_of_order() + { + cl_int err = CL_SUCCESS; + // Create ooo queues + clCommandQueueWrapper queue_1 = clCreateCommandQueue( + context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); + test_error(err, "Could not create command queue"); -// Confirm that a semaphore can be reused multiple times -int test_semaphores_reuse(cl_device_id deviceID, cl_context context, - cl_command_queue defaultQueue, int num_elements) -{ - cl_int err; + clCommandQueueWrapper queue_2 = clCreateCommandQueue( + context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); + test_error(err, "Could not create command queue"); + + return semaphore_cross_queue_helper(device, context, queue_1, queue_2); + } - if (!is_extension_available(deviceID, "cl_khr_semaphore")) + cl_int Run() override { - log_info("cl_khr_semaphore is not supported on this platform. " - "Skipping test.\n"); - return TEST_SKIPPED_ITSELF; + if (in_order) + return run_in_order(); + else + return run_out_of_order(); } +}; + +struct SemaphoreMultiSignal : public SemaphoreTestBase +{ + SemaphoreMultiSignal(cl_device_id device, cl_context context, + cl_command_queue queue) + : SemaphoreTestBase(device, context, queue), semaphore_second(this) + {} - // Obtain pointers to semaphore's API - GET_PFN(deviceID, clCreateSemaphoreWithPropertiesKHR); - GET_PFN(deviceID, clEnqueueSignalSemaphoresKHR); - GET_PFN(deviceID, clEnqueueWaitSemaphoresKHR); - GET_PFN(deviceID, clReleaseSemaphoreKHR); - - // Create ooo queue - clCommandQueueWrapper queue = clCreateCommandQueue( - context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); - test_error(err, "Could not create command queue"); - - // Create semaphore - cl_semaphore_properties_khr sema_props[] = { - static_cast(CL_SEMAPHORE_TYPE_KHR), - static_cast(CL_SEMAPHORE_TYPE_BINARY_KHR), - 0 - }; - cl_semaphore_khr sema = - clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err); - test_error(err, "Could not create semaphore"); - - // Create Kernel - clProgramWrapper program; - clKernelWrapper kernel; - err = create_single_kernel_helper(context, &program, &kernel, 1, &source, - "empty"); - test_error(err, "Could not create kernel"); - - constexpr size_t loop_count = 10; - clEventWrapper signal_events[loop_count]; - clEventWrapper wait_events[loop_count]; - clEventWrapper task_events[loop_count]; - - // Enqueue task_1 - err = clEnqueueTask(queue, kernel, 0, nullptr, &task_events[0]); - test_error(err, "Unable to enqueue task_1"); - - // Signal semaphore (dependency on task_1) - err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema, nullptr, 1, - &task_events[0], &signal_events[0]); - test_error(err, "Could not signal semaphore"); - - // In a loop - size_t loop; - for (loop = 1; loop < loop_count; ++loop) + cl_int Run() override { - // Wait semaphore - err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema, nullptr, 0, nullptr, - &wait_events[loop - 1]); + cl_int err = CL_SUCCESS; + // Create ooo queue + clCommandQueueWrapper queue = clCreateCommandQueue( + context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); + test_error(err, "Could not create command queue"); + + // Create semaphore + cl_semaphore_properties_khr sema_props[] = { + static_cast(CL_SEMAPHORE_TYPE_KHR), + static_cast( + CL_SEMAPHORE_TYPE_BINARY_KHR), + 0 + }; + semaphore = + clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err); + test_error(err, "Could not create semaphore"); + + semaphore_second = + clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err); + test_error(err, "Could not create semaphore"); + + // Signal semaphore 1 and 2 + clEventWrapper signal_event; + cl_semaphore_khr sema_list[] = { semaphore, semaphore_second }; + err = clEnqueueSignalSemaphoresKHR(queue, 2, sema_list, nullptr, 0, + nullptr, &signal_event); + test_error(err, "Could not signal semaphore"); + + // Wait semaphore 1 + clEventWrapper wait_1_event; + err = clEnqueueWaitSemaphoresKHR(queue, 1, semaphore, nullptr, 0, + nullptr, &wait_1_event); + test_error(err, "Could not wait semaphore"); + + // Wait semaphore 2 + clEventWrapper wait_2_event; + err = clEnqueueWaitSemaphoresKHR(queue, 1, semaphore_second, nullptr, 0, + nullptr, &wait_2_event); test_error(err, "Could not wait semaphore"); - // Enqueue task_loop (dependency on wait) - err = clEnqueueTask(queue, kernel, 1, &wait_events[loop - 1], - &task_events[loop]); - test_error(err, "Unable to enqueue task_loop"); + // Finish + err = clFinish(queue); + test_error(err, "Could not finish queue"); - // Wait for the "wait semaphore" to complete - err = clWaitForEvents(1, &wait_events[loop - 1]); - test_error(err, "Unable to wait for wait semaphore to complete"); + // Ensure all events are completed + test_assert_event_complete(signal_event); + test_assert_event_complete(wait_1_event); + test_assert_event_complete(wait_2_event); - // Signal semaphore (dependency on task_loop) - err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema, nullptr, 1, - &task_events[loop], - &signal_events[loop]); + return CL_SUCCESS; + } + clSemaphoreWrapper semaphore_second = nullptr; +}; + +struct SemaphoreMultiWait : public SemaphoreTestBase +{ + SemaphoreMultiWait(cl_device_id device, cl_context context, + cl_command_queue queue) + : SemaphoreTestBase(device, context, queue), semaphore_second(this) + {} + + cl_int Run() override + { + cl_int err = CL_SUCCESS; + // Create ooo queue + clCommandQueueWrapper queue = clCreateCommandQueue( + context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); + test_error(err, "Could not create command queue"); + + // Create semaphores + cl_semaphore_properties_khr sema_props[] = { + static_cast(CL_SEMAPHORE_TYPE_KHR), + static_cast( + CL_SEMAPHORE_TYPE_BINARY_KHR), + 0 + }; + semaphore = + clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err); + test_error(err, "Could not create semaphore"); + + semaphore_second = + clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err); + test_error(err, "Could not create semaphore"); + + // Signal semaphore 1 + clEventWrapper signal_1_event; + err = clEnqueueSignalSemaphoresKHR(queue, 1, semaphore, nullptr, 0, + nullptr, &signal_1_event); test_error(err, "Could not signal semaphore"); + + // Signal semaphore 2 + clEventWrapper signal_2_event; + err = clEnqueueSignalSemaphoresKHR(queue, 1, semaphore_second, nullptr, + 0, nullptr, &signal_2_event); + test_error(err, "Could not signal semaphore"); + + // Wait semaphore 1 and 2 + clEventWrapper wait_event; + cl_semaphore_khr sema_list[] = { semaphore, semaphore_second }; + err = clEnqueueWaitSemaphoresKHR(queue, 2, sema_list, nullptr, 0, + nullptr, &wait_event); + test_error(err, "Could not wait semaphore"); + + // Finish + err = clFinish(queue); + test_error(err, "Could not finish queue"); + + // Ensure all events are completed + test_assert_event_complete(signal_1_event); + test_assert_event_complete(signal_2_event); + test_assert_event_complete(wait_event); + + return CL_SUCCESS; } + clSemaphoreWrapper semaphore_second = nullptr; +}; + +struct SemaphoreQueries : public SemaphoreTestBase +{ + SemaphoreQueries(cl_device_id device, cl_context context, + cl_command_queue queue) + : SemaphoreTestBase(device, context, queue) + {} - // Wait semaphore - err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema, nullptr, 0, nullptr, - &wait_events[loop - 1]); - test_error(err, "Could not wait semaphore"); + cl_int Run() override + { + cl_int err = CL_SUCCESS; + // Create binary semaphore + cl_semaphore_properties_khr sema_props[] = { + static_cast(CL_SEMAPHORE_TYPE_KHR), + static_cast( + CL_SEMAPHORE_TYPE_BINARY_KHR), + static_cast( + CL_SEMAPHORE_DEVICE_HANDLE_LIST_KHR), + (cl_semaphore_properties_khr)device, + CL_SEMAPHORE_DEVICE_HANDLE_LIST_END_KHR, + 0 + }; + semaphore = + clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err); + test_error(err, "Could not create semaphore"); + + // Confirm that querying CL_SEMAPHORE_TYPE_KHR returns + // CL_SEMAPHORE_TYPE_BINARY_KHR + SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_TYPE_KHR, cl_semaphore_type_khr, + CL_SEMAPHORE_TYPE_BINARY_KHR); + + // Confirm that querying CL_SEMAPHORE_CONTEXT_KHR returns the right + // context + SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_CONTEXT_KHR, cl_context, context); + + // Confirm that querying CL_SEMAPHORE_REFERENCE_COUNT_KHR returns the + // right value + SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_REFERENCE_COUNT_KHR, cl_uint, 1); + + err = clRetainSemaphoreKHR(semaphore); + test_error(err, "Could not retain semaphore"); + SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_REFERENCE_COUNT_KHR, cl_uint, 2); + + err = clReleaseSemaphoreKHR(semaphore); + test_error(err, "Could not release semaphore"); + SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_REFERENCE_COUNT_KHR, cl_uint, 1); + + // Confirm that querying CL_SEMAPHORE_DEVICE_HANDLE_LIST_KHR returns the + // same device id the semaphore was created with + SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_DEVICE_HANDLE_LIST_KHR, cl_device_id, + device); + + // Confirm that querying CL_SEMAPHORE_PROPERTIES_KHR returns the same + // properties the semaphore was created with + SEMAPHORE_PARAM_TEST_ARRAY(CL_SEMAPHORE_PROPERTIES_KHR, + cl_semaphore_properties_khr, 6, sema_props); + + // Confirm that querying CL_SEMAPHORE_PAYLOAD_KHR returns the unsignaled + // state + SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_PAYLOAD_KHR, cl_semaphore_payload_khr, + 0); + + return CL_SUCCESS; + } +}; - // Finish - err = clFinish(queue); - test_error(err, "Could not finish queue"); +struct SemaphoreImportExportFD : public SemaphoreTestBase +{ + SemaphoreImportExportFD(cl_device_id device, cl_context context, + cl_command_queue queue) + : SemaphoreTestBase(device, context, queue), semaphore_second(this) + {} - // Ensure all events are completed - for (loop = 0; loop < loop_count; ++loop) + cl_int Run() override { - test_assert_event_complete(wait_events[loop]); - test_assert_event_complete(signal_events[loop]); - test_assert_event_complete(task_events[loop]); + cl_int err = CL_SUCCESS; + if (!is_extension_available(device, + "cl_khr_external_semaphore_sync_fd")) + { + log_info( + "cl_khr_external_semaphore_sync_fd is not supported on this " + "platform. Skipping test.\n"); + return TEST_SKIPPED_ITSELF; + } + + // Create ooo queue + clCommandQueueWrapper queue = clCreateCommandQueue( + context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); + test_error(err, "Could not create command queue"); + + // Create semaphore + cl_semaphore_properties_khr sema_1_props[] = { + static_cast(CL_SEMAPHORE_TYPE_KHR), + static_cast( + CL_SEMAPHORE_TYPE_BINARY_KHR), + static_cast( + CL_SEMAPHORE_EXPORT_HANDLE_TYPES_KHR), + static_cast( + CL_SEMAPHORE_HANDLE_SYNC_FD_KHR), + static_cast( + CL_SEMAPHORE_EXPORT_HANDLE_TYPES_LIST_END_KHR), + 0 + }; + semaphore = + clCreateSemaphoreWithPropertiesKHR(context, sema_1_props, &err); + test_error(err, "Could not create semaphore"); + + // Signal semaphore + clEventWrapper signal_event; + err = clEnqueueSignalSemaphoresKHR(queue, 1, semaphore, nullptr, 0, + nullptr, &signal_event); + test_error(err, "Could not signal semaphore"); + + // Extract sync fd + int handle = -1; + size_t handle_size; + err = clGetSemaphoreHandleForTypeKHR( + semaphore, device, CL_SEMAPHORE_HANDLE_SYNC_FD_KHR, sizeof(handle), + &handle, &handle_size); + test_error(err, "Could not extract semaphore handle"); + test_assert_error(sizeof(handle) == handle_size, "Invalid handle size"); + test_assert_error(handle >= 0, "Invalid handle"); + + // Create semaphore from sync fd + cl_semaphore_properties_khr sema_2_props[] = { + static_cast(CL_SEMAPHORE_TYPE_KHR), + static_cast( + CL_SEMAPHORE_TYPE_BINARY_KHR), + CL_SEMAPHORE_HANDLE_SYNC_FD_KHR, + static_cast(handle), 0 + }; + + semaphore_second = + clCreateSemaphoreWithPropertiesKHR(context, sema_2_props, &err); + test_error(err, "Could not create semaphore"); + + // Wait semaphore + clEventWrapper wait_event; + err = clEnqueueWaitSemaphoresKHR(queue, 1, semaphore_second, nullptr, 0, + nullptr, &wait_event); + test_error(err, "Could not wait semaphore"); + + // Finish + err = clFinish(queue); + test_error(err, "Could not finish queue"); + + // Check all events are completed + test_assert_event_complete(signal_event); + test_assert_event_complete(wait_event); + + return CL_SUCCESS; } + clSemaphoreWrapper semaphore_second = nullptr; +}; +} // anonymous namespace - // Release semaphore - err = clReleaseSemaphoreKHR(sema); - test_error(err, "Could not release semaphore"); +// Confirm that a signal followed by a wait will complete successfully +int test_semaphores_simple_1(cl_device_id deviceID, cl_context context, + cl_command_queue defaultQueue, int num_elements) +{ + return MakeAndRunTest(deviceID, context, defaultQueue); +} + +// Confirm that signal a semaphore with no event dependencies will not result +// in an implicit dependency on everything previously submitted +int test_semaphores_simple_2(cl_device_id deviceID, cl_context context, + cl_command_queue defaultQueue, int num_elements) +{ + return MakeAndRunTest(deviceID, context, defaultQueue); +} - return TEST_PASS; +// Confirm that a semaphore can be reused multiple times +int test_semaphores_reuse(cl_device_id deviceID, cl_context context, + cl_command_queue defaultQueue, int num_elements) +{ + return MakeAndRunTest(deviceID, context, defaultQueue); } // Confirm that a semaphore works across different ooo queues @@ -397,18 +697,8 @@ int test_semaphores_cross_queues_ooo(cl_device_id deviceID, cl_context context, cl_command_queue defaultQueue, int num_elements) { - cl_int err; - - // Create ooo queues - clCommandQueueWrapper queue_1 = clCreateCommandQueue( - context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); - test_error(err, "Could not create command queue"); - - clCommandQueueWrapper queue_2 = clCreateCommandQueue( - context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); - test_error(err, "Could not create command queue"); - - return semaphore_cross_queue_helper(deviceID, context, queue_1, queue_2); + return MakeAndRunTest>(deviceID, context, + defaultQueue); } // Confirm that a semaphore works across different in-order queues @@ -416,18 +706,8 @@ int test_semaphores_cross_queues_io(cl_device_id deviceID, cl_context context, cl_command_queue defaultQueue, int num_elements) { - cl_int err; - - // Create in-order queues - clCommandQueueWrapper queue_1 = - clCreateCommandQueue(context, deviceID, 0, &err); - test_error(err, "Could not create command queue"); - - clCommandQueueWrapper queue_2 = - clCreateCommandQueue(context, deviceID, 0, &err); - test_error(err, "Could not create command queue"); - - return semaphore_cross_queue_helper(deviceID, context, queue_1, queue_2); + return MakeAndRunTest>(deviceID, context, + defaultQueue); } // Confirm that we can signal multiple semaphores with one command @@ -435,225 +715,22 @@ int test_semaphores_multi_signal(cl_device_id deviceID, cl_context context, cl_command_queue defaultQueue, int num_elements) { - cl_int err; - - if (!is_extension_available(deviceID, "cl_khr_semaphore")) - { - log_info("cl_khr_semaphore is not supported on this platform. " - "Skipping test.\n"); - return TEST_SKIPPED_ITSELF; - } - - // Obtain pointers to semaphore's API - GET_PFN(deviceID, clCreateSemaphoreWithPropertiesKHR); - GET_PFN(deviceID, clEnqueueSignalSemaphoresKHR); - GET_PFN(deviceID, clEnqueueWaitSemaphoresKHR); - GET_PFN(deviceID, clReleaseSemaphoreKHR); - - // Create ooo queue - clCommandQueueWrapper queue = clCreateCommandQueue( - context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); - test_error(err, "Could not create command queue"); - - // Create semaphore - cl_semaphore_properties_khr sema_props[] = { - static_cast(CL_SEMAPHORE_TYPE_KHR), - static_cast(CL_SEMAPHORE_TYPE_BINARY_KHR), - 0 - }; - cl_semaphore_khr sema_1 = - clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err); - test_error(err, "Could not create semaphore"); - - cl_semaphore_khr sema_2 = - clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err); - test_error(err, "Could not create semaphore"); - - // Signal semaphore 1 and 2 - clEventWrapper signal_event; - cl_semaphore_khr sema_list[] = { sema_1, sema_2 }; - err = clEnqueueSignalSemaphoresKHR(queue, 2, sema_list, nullptr, 0, nullptr, - &signal_event); - test_error(err, "Could not signal semaphore"); - - // Wait semaphore 1 - clEventWrapper wait_1_event; - err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_1, nullptr, 0, nullptr, - &wait_1_event); - test_error(err, "Could not wait semaphore"); - - // Wait semaphore 2 - clEventWrapper wait_2_event; - err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_2, nullptr, 0, nullptr, - &wait_2_event); - test_error(err, "Could not wait semaphore"); - - // Finish - err = clFinish(queue); - test_error(err, "Could not finish queue"); - - // Ensure all events are completed - test_assert_event_complete(signal_event); - test_assert_event_complete(wait_1_event); - test_assert_event_complete(wait_2_event); - - // Release semaphores - err = clReleaseSemaphoreKHR(sema_1); - test_error(err, "Could not release semaphore"); - - err = clReleaseSemaphoreKHR(sema_2); - test_error(err, "Could not release semaphore"); - - return TEST_PASS; + return MakeAndRunTest(deviceID, context, + defaultQueue); } // Confirm that we can wait for multiple semaphores with one command int test_semaphores_multi_wait(cl_device_id deviceID, cl_context context, cl_command_queue defaultQueue, int num_elements) { - cl_int err; - - if (!is_extension_available(deviceID, "cl_khr_semaphore")) - { - log_info("cl_khr_semaphore is not supported on this platform. " - "Skipping test.\n"); - return TEST_SKIPPED_ITSELF; - } - - // Obtain pointers to semaphore's API - GET_PFN(deviceID, clCreateSemaphoreWithPropertiesKHR); - GET_PFN(deviceID, clEnqueueSignalSemaphoresKHR); - GET_PFN(deviceID, clEnqueueWaitSemaphoresKHR); - GET_PFN(deviceID, clReleaseSemaphoreKHR); - - // Create ooo queue - clCommandQueueWrapper queue = clCreateCommandQueue( - context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); - test_error(err, "Could not create command queue"); - - // Create semaphores - cl_semaphore_properties_khr sema_props[] = { - static_cast(CL_SEMAPHORE_TYPE_KHR), - static_cast(CL_SEMAPHORE_TYPE_BINARY_KHR), - 0 - }; - cl_semaphore_khr sema_1 = - clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err); - test_error(err, "Could not create semaphore"); - - cl_semaphore_khr sema_2 = - clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err); - test_error(err, "Could not create semaphore"); - - // Signal semaphore 1 - clEventWrapper signal_1_event; - err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema_1, nullptr, 0, nullptr, - &signal_1_event); - test_error(err, "Could not signal semaphore"); - - // Signal semaphore 2 - clEventWrapper signal_2_event; - err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema_2, nullptr, 0, nullptr, - &signal_2_event); - test_error(err, "Could not signal semaphore"); - - // Wait semaphore 1 and 2 - clEventWrapper wait_event; - cl_semaphore_khr sema_list[] = { sema_1, sema_2 }; - err = clEnqueueWaitSemaphoresKHR(queue, 2, sema_list, nullptr, 0, nullptr, - &wait_event); - test_error(err, "Could not wait semaphore"); - - // Finish - err = clFinish(queue); - test_error(err, "Could not finish queue"); - - // Ensure all events are completed - test_assert_event_complete(signal_1_event); - test_assert_event_complete(signal_2_event); - test_assert_event_complete(wait_event); - - // Release semaphores - err = clReleaseSemaphoreKHR(sema_1); - test_error(err, "Could not release semaphore"); - - err = clReleaseSemaphoreKHR(sema_2); - test_error(err, "Could not release semaphore"); - - return TEST_PASS; + return MakeAndRunTest(deviceID, context, defaultQueue); } // Confirm the semaphores can be successfully queried int test_semaphores_queries(cl_device_id deviceID, cl_context context, cl_command_queue defaultQueue, int num_elements) { - cl_int err = CL_SUCCESS; - - if (!is_extension_available(deviceID, "cl_khr_semaphore")) - { - log_info("cl_khr_semaphore is not supported on this platform. " - "Skipping test.\n"); - return TEST_SKIPPED_ITSELF; - } - - // Obtain pointers to semaphore's API - GET_PFN(deviceID, clCreateSemaphoreWithPropertiesKHR); - GET_PFN(deviceID, clGetSemaphoreInfoKHR); - GET_PFN(deviceID, clRetainSemaphoreKHR); - GET_PFN(deviceID, clReleaseSemaphoreKHR); - - // Create binary semaphore - cl_semaphore_properties_khr sema_props[] = { - static_cast(CL_SEMAPHORE_TYPE_KHR), - static_cast(CL_SEMAPHORE_TYPE_BINARY_KHR), - static_cast( - CL_SEMAPHORE_DEVICE_HANDLE_LIST_KHR), - (cl_semaphore_properties_khr)deviceID, - CL_SEMAPHORE_DEVICE_HANDLE_LIST_END_KHR, - 0 - }; - cl_semaphore_khr sema = - clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err); - test_error(err, "Could not create semaphore"); - - // Confirm that querying CL_SEMAPHORE_TYPE_KHR returns - // CL_SEMAPHORE_TYPE_BINARY_KHR - SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_TYPE_KHR, cl_semaphore_type_khr, - CL_SEMAPHORE_TYPE_BINARY_KHR); - - // Confirm that querying CL_SEMAPHORE_CONTEXT_KHR returns the right context - SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_CONTEXT_KHR, cl_context, context); - - // Confirm that querying CL_SEMAPHORE_REFERENCE_COUNT_KHR returns the right - // value - SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_REFERENCE_COUNT_KHR, cl_uint, 1); - - // Confirm that querying CL_SEMAPHORE_DEVICE_HANDLE_LIST_KHR returns the - // same device id the semaphore was created with - SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_DEVICE_HANDLE_LIST_KHR, cl_device_id, - deviceID); - - err = clRetainSemaphoreKHR(sema); - test_error(err, "Could not retain semaphore"); - SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_REFERENCE_COUNT_KHR, cl_uint, 2); - - err = clReleaseSemaphoreKHR(sema); - test_error(err, "Could not release semaphore"); - SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_REFERENCE_COUNT_KHR, cl_uint, 1); - - // Confirm that querying CL_SEMAPHORE_PROPERTIES_KHR returns the same - // properties the semaphore was created with - SEMAPHORE_PARAM_TEST_ARRAY(CL_SEMAPHORE_PROPERTIES_KHR, - cl_semaphore_properties_khr, 6, sema_props); - - // Confirm that querying CL_SEMAPHORE_PAYLOAD_KHR returns the unsignaled - // state - SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_PAYLOAD_KHR, cl_semaphore_payload_khr, 0); - - err = clReleaseSemaphoreKHR(sema); - test_error(err, "Could not release semaphore"); - - return TEST_PASS; + return MakeAndRunTest(deviceID, context, defaultQueue); } // Test it is possible to export a semaphore to a sync fd and import the same @@ -662,97 +739,6 @@ int test_semaphores_import_export_fd(cl_device_id deviceID, cl_context context, cl_command_queue defaultQueue, int num_elements) { - cl_int err; - - if (!is_extension_available(deviceID, "cl_khr_semaphore")) - { - log_info("cl_khr_semaphore is not supported on this platform. " - "Skipping test.\n"); - return TEST_SKIPPED_ITSELF; - } - - if (!is_extension_available(deviceID, "cl_khr_external_semaphore_sync_fd")) - { - log_info("cl_khr_external_semaphore_sync_fd is not supported on this " - "platform. Skipping test.\n"); - return TEST_SKIPPED_ITSELF; - } - - // Obtain pointers to semaphore's API - GET_PFN(deviceID, clCreateSemaphoreWithPropertiesKHR); - GET_PFN(deviceID, clEnqueueSignalSemaphoresKHR); - GET_PFN(deviceID, clEnqueueWaitSemaphoresKHR); - GET_PFN(deviceID, clGetSemaphoreHandleForTypeKHR); - GET_PFN(deviceID, clReleaseSemaphoreKHR); - - // Create ooo queue - clCommandQueueWrapper queue = clCreateCommandQueue( - context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); - test_error(err, "Could not create command queue"); - - // Create semaphore - cl_semaphore_properties_khr sema_1_props[] = { - static_cast(CL_SEMAPHORE_TYPE_KHR), - static_cast(CL_SEMAPHORE_TYPE_BINARY_KHR), - static_cast( - CL_SEMAPHORE_EXPORT_HANDLE_TYPES_KHR), - static_cast( - CL_SEMAPHORE_HANDLE_SYNC_FD_KHR), - static_cast( - CL_SEMAPHORE_EXPORT_HANDLE_TYPES_LIST_END_KHR), - 0 - }; - cl_semaphore_khr sema_1 = - clCreateSemaphoreWithPropertiesKHR(context, sema_1_props, &err); - test_error(err, "Could not create semaphore"); - - // Signal semaphore - clEventWrapper signal_event; - err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema_1, nullptr, 0, nullptr, - &signal_event); - test_error(err, "Could not signal semaphore"); - - // Extract sync fd - int handle = -1; - size_t handle_size; - err = clGetSemaphoreHandleForTypeKHR(sema_1, deviceID, - CL_SEMAPHORE_HANDLE_SYNC_FD_KHR, - sizeof(handle), &handle, &handle_size); - test_error(err, "Could not extract semaphore handle"); - test_assert_error(sizeof(handle) == handle_size, "Invalid handle size"); - test_assert_error(handle >= 0, "Invalid handle"); - - // Create semaphore from sync fd - cl_semaphore_properties_khr sema_2_props[] = { - static_cast(CL_SEMAPHORE_TYPE_KHR), - static_cast(CL_SEMAPHORE_TYPE_BINARY_KHR), - CL_SEMAPHORE_HANDLE_SYNC_FD_KHR, - static_cast(handle), 0 - }; - - cl_semaphore_khr sema_2 = - clCreateSemaphoreWithPropertiesKHR(context, sema_2_props, &err); - test_error(err, "Could not create semaphore"); - - // Wait semaphore - clEventWrapper wait_event; - err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_2, nullptr, 0, nullptr, - &wait_event); - test_error(err, "Could not wait semaphore"); - - // Finish - err = clFinish(queue); - test_error(err, "Could not finish queue"); - - // Check all events are completed - test_assert_event_complete(signal_event); - test_assert_event_complete(wait_event); - - // Release semaphore - err = clReleaseSemaphoreKHR(sema_1); - test_error(err, "Could not release semaphore"); - - err = clReleaseSemaphoreKHR(sema_2); - test_error(err, "Could not release semaphore"); - return TEST_PASS; + return MakeAndRunTest(deviceID, context, + defaultQueue); } From 9116bb7acbadfd79db5daf979dddea81b7fd4e75 Mon Sep 17 00:00:00 2001 From: Sven van Haastregt Date: Tue, 3 Sep 2024 19:30:13 +0200 Subject: [PATCH 7/7] [NFC] math_brute_force: move TestInfoBase to common.h (#2059) The various forms of `TestInfoBase` have many members in common, so avoid duplicating the struct definition and move it to `common.h`. Provide a description and initializer for every struct member, and drop initializations done with `memset`. Signed-off-by: Sven van Haastregt --- .../math_brute_force/binary_half.cpp | 23 +-------- .../math_brute_force/binary_i_half.cpp | 19 +------- test_conformance/math_brute_force/common.h | 47 +++++++++++++++++++ .../math_brute_force/macro_binary_half.cpp | 18 +------ .../math_brute_force/macro_unary_half.cpp | 17 +------ .../math_brute_force/unary_half.cpp | 22 +-------- 6 files changed, 52 insertions(+), 94 deletions(-) diff --git a/test_conformance/math_brute_force/binary_half.cpp b/test_conformance/math_brute_force/binary_half.cpp index 3a2395c705..897bd6e7c6 100644 --- a/test_conformance/math_brute_force/binary_half.cpp +++ b/test_conformance/math_brute_force/binary_half.cpp @@ -56,27 +56,8 @@ struct ThreadInfo tQueue; // per thread command queue to improve performance }; -struct TestInfoBase -{ - size_t subBufferSize; // Size of the sub-buffer in elements - const Func *f; // A pointer to the function info - - cl_uint threadCount; // Number of worker threads - cl_uint jobCount; // Number of jobs - cl_uint step; // step between each chunk and the next. - cl_uint scale; // stride between individual test values - float ulps; // max_allowed ulps - int ftz; // non-zero if running in flush to zero mode - - int isFDim; - int skipNanInf; - int isNextafter; -}; - struct TestInfo : public TestInfoBase { - TestInfo(const TestInfoBase &base): TestInfoBase(base) {} - // Array of thread specific information std::vector tinfo; @@ -646,7 +627,6 @@ cl_int TestHalf(cl_uint job_id, cl_uint thread_id, void *data) int TestFunc_Half_Half_Half_common(const Func *f, MTdata d, int isNextafter, bool relaxedMode) { - TestInfoBase test_info_base; cl_int error; float maxError = 0.0f; double maxErrorVal = 0.0; @@ -654,8 +634,7 @@ int TestFunc_Half_Half_Half_common(const Func *f, MTdata d, int isNextafter, logFunctionInfo(f->name, sizeof(cl_half), relaxedMode); // Init test_info - memset(&test_info_base, 0, sizeof(test_info_base)); - TestInfo test_info(test_info_base); + TestInfo test_info; test_info.threadCount = GetThreadCount(); test_info.subBufferSize = BUFFER_SIZE diff --git a/test_conformance/math_brute_force/binary_i_half.cpp b/test_conformance/math_brute_force/binary_i_half.cpp index c74a845a4f..2d6e827fdf 100644 --- a/test_conformance/math_brute_force/binary_i_half.cpp +++ b/test_conformance/math_brute_force/binary_i_half.cpp @@ -52,23 +52,8 @@ typedef struct ThreadInfo tQueue; // per thread command queue to improve performance } ThreadInfo; -struct TestInfoBase -{ - size_t subBufferSize; // Size of the sub-buffer in elements - const Func *f; // A pointer to the function info - - cl_uint threadCount; // Number of worker threads - cl_uint jobCount; // Number of jobs - cl_uint step; // step between each chunk and the next. - cl_uint scale; // stride between individual test values - float ulps; // max_allowed ulps - int ftz; // non-zero if running in flush to zero mode -}; - struct TestInfo : public TestInfoBase { - TestInfo(const TestInfoBase &base): TestInfoBase(base) {} - // Array of thread specific information std::vector tinfo; @@ -415,7 +400,6 @@ cl_int TestHalf(cl_uint job_id, cl_uint thread_id, void *data) int TestFunc_Half_Half_Int(const Func *f, MTdata d, bool relaxedMode) { - TestInfoBase test_info_base; cl_int error; size_t i, j; float maxError = 0.0f; @@ -425,8 +409,7 @@ int TestFunc_Half_Half_Int(const Func *f, MTdata d, bool relaxedMode) logFunctionInfo(f->name, sizeof(cl_half), relaxedMode); // Init test_info - memset(&test_info_base, 0, sizeof(test_info_base)); - TestInfo test_info(test_info_base); + TestInfo test_info; test_info.threadCount = GetThreadCount(); test_info.subBufferSize = BUFFER_SIZE diff --git a/test_conformance/math_brute_force/common.h b/test_conformance/math_brute_force/common.h index eb82c5f8f3..3f89ef6cb8 100644 --- a/test_conformance/math_brute_force/common.h +++ b/test_conformance/math_brute_force/common.h @@ -86,6 +86,53 @@ struct BuildKernelInfo bool relaxedMode; }; +// Data common to all math tests. +struct TestInfoBase +{ + TestInfoBase() = default; + ~TestInfoBase() = default; + + // Prevent accidental copy/move. + TestInfoBase(const TestInfoBase &) = delete; + TestInfoBase &operator=(const TestInfoBase &) = delete; + TestInfoBase(TestInfoBase &&h) = delete; + TestInfoBase &operator=(TestInfoBase &&h) = delete; + + // Size of the sub-buffer in elements. + size_t subBufferSize = 0; + // Function info. + const Func *f = nullptr; + + // Number of worker threads. + cl_uint threadCount = 0; + // Number of jobs. + cl_uint jobCount = 0; + // step between each chunk and the next. + cl_uint step = 0; + // stride between individual test values. + cl_uint scale = 0; + // max_allowed ulps. + float ulps = -1.f; + // non-zero if running in flush to zero mode. + int ftz = 0; + + // 1 if running the fdim test. + int isFDim = 0; + // 1 if input/output NaNs and INFs are skipped. + int skipNanInf = 0; + // 1 if running the nextafter test. + int isNextafter = 0; + + // 1 if the function is only to be evaluated over a range. + int isRangeLimited = 0; + + // Result limit for half_sin/half_cos/half_tan. + float half_sin_cos_tan_limit = -1.f; + + // Whether the test is being run in relaxed mode. + bool relaxedMode = false; +}; + using SourceGenerator = std::string (*)(const std::string &kernel_name, const char *builtin, cl_uint vector_size_index); diff --git a/test_conformance/math_brute_force/macro_binary_half.cpp b/test_conformance/math_brute_force/macro_binary_half.cpp index ea4ef8128f..a8f459a7af 100644 --- a/test_conformance/math_brute_force/macro_binary_half.cpp +++ b/test_conformance/math_brute_force/macro_binary_half.cpp @@ -45,22 +45,8 @@ struct ThreadInfo tQueue; // per thread command queue to improve performance }; -struct TestInfoBase -{ - size_t subBufferSize; // Size of the sub-buffer in elements - const Func *f; // A pointer to the function info - - cl_uint threadCount; // Number of worker threads - cl_uint jobCount; // Number of jobs - cl_uint step; // step between each chunk and the next. - cl_uint scale; // stride between individual test values - int ftz; // non-zero if running in flush to zero mode -}; - struct TestInfo : public TestInfoBase { - TestInfo(const TestInfoBase &base): TestInfoBase(base) {} - // Array of thread specific information std::vector tinfo; @@ -430,15 +416,13 @@ cl_int TestHalf(cl_uint job_id, cl_uint thread_id, void *data) int TestMacro_Int_Half_Half(const Func *f, MTdata d, bool relaxedMode) { - TestInfoBase test_info_base; cl_int error; size_t i, j; logFunctionInfo(f->name, sizeof(cl_half), relaxedMode); // Init test_info - memset(&test_info_base, 0, sizeof(test_info_base)); - TestInfo test_info(test_info_base); + TestInfo test_info; test_info.threadCount = GetThreadCount(); test_info.subBufferSize = BUFFER_SIZE diff --git a/test_conformance/math_brute_force/macro_unary_half.cpp b/test_conformance/math_brute_force/macro_unary_half.cpp index cb20205514..a1e9211876 100644 --- a/test_conformance/math_brute_force/macro_unary_half.cpp +++ b/test_conformance/math_brute_force/macro_unary_half.cpp @@ -43,21 +43,8 @@ struct ThreadInfo tQueue; // per thread command queue to improve performance }; -struct TestInfoBase -{ - size_t subBufferSize; // Size of the sub-buffer in elements - const Func *f; // A pointer to the function info - cl_uint threadCount; // Number of worker threads - cl_uint jobCount; // Number of jobs - cl_uint step; // step between each chunk and the next. - cl_uint scale; // stride between individual test values - int ftz; // non-zero if running in flush to zero mode -}; - struct TestInfo : public TestInfoBase { - TestInfo(const TestInfoBase &base): TestInfoBase(base) {} - // Array of thread specific information std::vector tinfo; @@ -328,14 +315,12 @@ cl_int TestHalf(cl_uint job_id, cl_uint thread_id, void *data) int TestMacro_Int_Half(const Func *f, MTdata d, bool relaxedMode) { - TestInfoBase test_info_base; cl_int error; size_t i, j; logFunctionInfo(f->name, sizeof(cl_half), relaxedMode); // Init test_info - memset(&test_info_base, 0, sizeof(test_info_base)); - TestInfo test_info(test_info_base); + TestInfo test_info; test_info.threadCount = GetThreadCount(); test_info.subBufferSize = BUFFER_SIZE diff --git a/test_conformance/math_brute_force/unary_half.cpp b/test_conformance/math_brute_force/unary_half.cpp index 0980fb16ca..7af16e7386 100644 --- a/test_conformance/math_brute_force/unary_half.cpp +++ b/test_conformance/math_brute_force/unary_half.cpp @@ -45,26 +45,8 @@ typedef struct ThreadInfo tQueue; // per thread command queue to improve performance } ThreadInfo; -struct TestInfoBase -{ - size_t subBufferSize; // Size of the sub-buffer in elements - const Func *f; // A pointer to the function info - cl_uint threadCount; // Number of worker threads - cl_uint jobCount; // Number of jobs - cl_uint step; // step between each chunk and the next. - cl_uint scale; // stride between individual test values - float ulps; // max_allowed ulps - int ftz; // non-zero if running in flush to zero mode - - int isRangeLimited; // 1 if the function is only to be evaluated over a - // range - float half_sin_cos_tan_limit; -}; - struct TestInfo : public TestInfoBase { - TestInfo(const TestInfoBase &base): TestInfoBase(base) {} - // Array of thread specific information std::vector tinfo; @@ -351,7 +333,6 @@ cl_int TestHalf(cl_uint job_id, cl_uint thread_id, void *data) int TestFunc_Half_Half(const Func *f, MTdata d, bool relaxedMode) { - TestInfoBase test_info_base; cl_int error; size_t i, j; float maxError = 0.0f; @@ -360,8 +341,7 @@ int TestFunc_Half_Half(const Func *f, MTdata d, bool relaxedMode) logFunctionInfo(f->name, sizeof(cl_half), relaxedMode); // Init test_info - memset(&test_info_base, 0, sizeof(test_info_base)); - TestInfo test_info(test_info_base); + TestInfo test_info; test_info.threadCount = GetThreadCount();