From 5ef065a74cabc91ae16c6115cb88fe5f58b68cd7 Mon Sep 17 00:00:00 2001 From: Nives Vukovic Date: Fri, 8 Dec 2023 10:11:51 +0000 Subject: [PATCH] EXSWHTEC-316 - Implement tests for Complex type functions (#356) Change-Id: I67d5e1fd02c7e40319f135bc9ae5bbbde85b5ee7 [ROCm/hip-tests commit: 2f783afe8c45a50cbb3c8429512273ffe7596160] --- .../catch/hipTestMain/config/config_amd_linux | 16 + .../hipTestMain/config/config_amd_windows | 13 + .../config/config_nvidia_linux.json | 15 +- .../config/config_nvidia_windows.json | 15 +- .../catch/include/hip_test_defgroups.hh | 7 + projects/hip-tests/catch/unit/CMakeLists.txt | 1 + .../catch/unit/complex/CMakeLists.txt | 72 ++ .../hip-tests/catch/unit/complex/complex.cc | 479 ++++++++++++++ .../unit/complex/complex_basic_common.hh | 106 +++ .../complex/complex_cast_negative_kernels.cc | 113 ++++ .../complex_cast_negative_kernels_rtc.hh | 121 ++++ .../unit/complex/complex_function_common.hh | 311 +++++++++ .../complex/complex_make_negative_kernels.cc | 121 ++++ .../complex_make_negative_kernels_rtc.hh | 198 ++++++ .../complex_negative_kernels_1Arg_double.cc | 75 +++ .../complex_negative_kernels_1Arg_float.cc | 75 +++ .../complex_negative_kernels_1Arg_rtc.hh | 207 ++++++ .../complex_negative_kernels_2Arg_double.cc | 105 +++ .../complex_negative_kernels_2Arg_float.cc | 105 +++ .../complex_negative_kernels_2Arg_rtc.hh | 620 ++++++++++++++++++ .../complex/complex_negative_kernels_3Arg.cc | 247 +++++++ .../complex_negative_kernels_3Arg_rtc.hh | 246 +++++++ 22 files changed, 3266 insertions(+), 2 deletions(-) create mode 100644 projects/hip-tests/catch/unit/complex/CMakeLists.txt create mode 100644 projects/hip-tests/catch/unit/complex/complex.cc create mode 100644 projects/hip-tests/catch/unit/complex/complex_basic_common.hh create mode 100644 projects/hip-tests/catch/unit/complex/complex_cast_negative_kernels.cc create mode 100644 projects/hip-tests/catch/unit/complex/complex_cast_negative_kernels_rtc.hh create mode 100644 projects/hip-tests/catch/unit/complex/complex_function_common.hh create mode 100644 projects/hip-tests/catch/unit/complex/complex_make_negative_kernels.cc create mode 100644 projects/hip-tests/catch/unit/complex/complex_make_negative_kernels_rtc.hh create mode 100644 projects/hip-tests/catch/unit/complex/complex_negative_kernels_1Arg_double.cc create mode 100644 projects/hip-tests/catch/unit/complex/complex_negative_kernels_1Arg_float.cc create mode 100644 projects/hip-tests/catch/unit/complex/complex_negative_kernels_1Arg_rtc.hh create mode 100644 projects/hip-tests/catch/unit/complex/complex_negative_kernels_2Arg_double.cc create mode 100644 projects/hip-tests/catch/unit/complex/complex_negative_kernels_2Arg_float.cc create mode 100644 projects/hip-tests/catch/unit/complex/complex_negative_kernels_2Arg_rtc.hh create mode 100644 projects/hip-tests/catch/unit/complex/complex_negative_kernels_3Arg.cc create mode 100644 projects/hip-tests/catch/unit/complex/complex_negative_kernels_3Arg_rtc.hh diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux index bba683b4aa..2b93de19f8 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux @@ -127,6 +127,7 @@ "Unit_deviceAllocation_InOneThread_AccessInAllThreads", "=== Patch which removes the typetraits implementation from std namespace in hiprtc is reverted ===", "Unit_hiprtc_stdheaders", +<<<<<<< HEAD "Unit_hipMemAddressFree_negative", "Unit_hipMemAddressReserve_AlignmentTest", "Unit_hipMemAddressReserve_Negative", @@ -249,6 +250,21 @@ "Unit_hipGraphicsResourceGetMappedPointer_Negative_Parameters", "Unit_hipGraphicsUnmapResources_Negative_Parameters", "Unit_hipGraphicsUnregisterResource_Negative_Parameters", +======= + "=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/356 ===", + "Unit_Device_Complex_Unary_Negative_Parameters_RTC", + "Unit_Device_Complex_Binary_Negative_Parameters_RTC", + "Unit_Device_Complex_hipCfma_Negative_Parameters_RTC", + "Unit_Device_make_Complex_Negative_Parameters_RTC", + "Unit_Device_Complex_Cast_Negative_Parameters_RTC", + "Unit_Device_Complex_make_Negative", + "Unit_Device_Complex_Cast_Negative", + "Unit_Device_Complex_Unary_float_Negative", + "Unit_Device_Complex_Unary_double_Negative", + "Unit_Device_Complex_Binary_float_Negative", + "Unit_Device_Complex_Binary_double_Negative", + "Unit_Device_Complex_hipCfma_Negative", +>>>>>>> ce6de407 (EXSWHTEC-316 - Implement tests for Complex type functions (#356)) #endif #if defined VEGA20 "=== SWDEV-419112 Below tests fail in stress test on 29/08/23 ===", diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows index f9080bed55..e60a130b00 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows @@ -335,6 +335,19 @@ "=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/279 ===", "Unit_Device_memcpy_Negative", "Unit_Device_memset_Negative", + "=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/356 ===", + "Unit_Device_Complex_Unary_Negative_Parameters_RTC", + "Unit_Device_Complex_Binary_Negative_Parameters_RTC", + "Unit_Device_Complex_hipCfma_Negative_Parameters_RTC", + "Unit_Device_make_Complex_Negative_Parameters_RTC", + "Unit_Device_Complex_Cast_Negative_Parameters_RTC", + "Unit_Device_Complex_make_Negative", + "Unit_Device_Complex_Cast_Negative", + "Unit_Device_Complex_Unary_float_Negative", + "Unit_Device_Complex_Unary_double_Negative", + "Unit_Device_Complex_Binary_float_Negative", + "Unit_Device_Complex_Binary_double_Negative", + "Unit_Device_Complex_hipCfma_Negative", #endif "End of json" ] diff --git a/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux.json b/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux.json index 5daf1caa1b..cdf3eccc68 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux.json @@ -57,6 +57,19 @@ "Unit_Printf_Negative", "=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/279 ===", "Unit_Device_memcpy_Negative", - "Unit_Device_memset_Negative" + "Unit_Device_memset_Negative", + "=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/356 ===", + "Unit_Device_Complex_Unary_Negative_Parameters_RTC", + "Unit_Device_Complex_Binary_Negative_Parameters_RTC", + "Unit_Device_Complex_hipCfma_Negative_Parameters_RTC", + "Unit_Device_make_Complex_Negative_Parameters_RTC", + "Unit_Device_Complex_Cast_Negative_Parameters_RTC", + "Unit_Device_Complex_make_Negative", + "Unit_Device_Complex_Cast_Negative", + "Unit_Device_Complex_Unary_float_Negative", + "Unit_Device_Complex_Unary_double_Negative", + "Unit_Device_Complex_Binary_float_Negative", + "Unit_Device_Complex_Binary_double_Negative", + "Unit_Device_Complex_hipCfma_Negative" ] } diff --git a/projects/hip-tests/catch/hipTestMain/config/config_nvidia_windows.json b/projects/hip-tests/catch/hipTestMain/config/config_nvidia_windows.json index 5f35a92780..94aaeb1201 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_nvidia_windows.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_nvidia_windows.json @@ -15,6 +15,19 @@ "Unit_ChannelDescriptor_Positive_Basic_3D - ulong3", "Unit_ChannelDescriptor_Positive_Basic_3D - long3", "Unit_ChannelDescriptor_Positive_Basic_4D - ulong4", - "Unit_ChannelDescriptor_Positive_Basic_4D - long4" + "Unit_ChannelDescriptor_Positive_Basic_4D - long4", + "=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/356 ===", + "Unit_Device_Complex_Unary_Negative_Parameters_RTC", + "Unit_Device_Complex_Binary_Negative_Parameters_RTC", + "Unit_Device_Complex_hipCfma_Negative_Parameters_RTC", + "Unit_Device_make_Complex_Negative_Parameters_RTC", + "Unit_Device_Complex_Cast_Negative_Parameters_RTC", + "Unit_Device_Complex_make_Negative", + "Unit_Device_Complex_Cast_Negative", + "Unit_Device_Complex_Unary_float_Negative", + "Unit_Device_Complex_Unary_double_Negative", + "Unit_Device_Complex_Binary_float_Negative", + "Unit_Device_Complex_Binary_double_Negative", + "Unit_Device_Complex_hipCfma_Negative" ] } diff --git a/projects/hip-tests/catch/include/hip_test_defgroups.hh b/projects/hip-tests/catch/include/hip_test_defgroups.hh index 44a0b406f5..cb95778f3e 100644 --- a/projects/hip-tests/catch/include/hip_test_defgroups.hh +++ b/projects/hip-tests/catch/include/hip_test_defgroups.hh @@ -193,3 +193,10 @@ THE SOFTWARE. * This section describes tests for the surface management functions of HIP runtime API. * @} */ + + /** + * @defgroup ComplexTest Complex type + * @{ + * This section describes tests for the Complex type functions. + * @} + */ diff --git a/projects/hip-tests/catch/unit/CMakeLists.txt b/projects/hip-tests/catch/unit/CMakeLists.txt index 5557f5c3cc..8fac51f5c6 100644 --- a/projects/hip-tests/catch/unit/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/CMakeLists.txt @@ -45,6 +45,7 @@ add_subdirectory(channelDescriptor) add_subdirectory(executionControl) add_subdirectory(vector_types) add_subdirectory(atomics) +add_subdirectory(complex) add_subdirectory(p2p) add_subdirectory(gcc) diff --git a/projects/hip-tests/catch/unit/complex/CMakeLists.txt b/projects/hip-tests/catch/unit/complex/CMakeLists.txt new file mode 100644 index 0000000000..971369e2b1 --- /dev/null +++ b/projects/hip-tests/catch/unit/complex/CMakeLists.txt @@ -0,0 +1,72 @@ +# Copyright (c) 2023 Advanced Micro Devices, Inc. All Rights Reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +# Common Tests - Test independent of all platforms + +set(TEST_SRC + complex.cc +) + +if(HIP_PLATFORM MATCHES "nvidia") + set(LINKER_LIBS nvrtc) +elseif(HIP_PLATFORM MATCHES "amd") + set(LINKER_LIBS hiprtc) +endif() + +hip_add_exe_to_target(NAME ComplexTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests + LINKER_LIBS ${LINKER_LIBS}) + +add_test(NAME Unit_Device_Complex_make_Negative + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + complex_make_negative_kernels.cc 54) + +add_test(NAME Unit_Device_Complex_Cast_Negative + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + complex_cast_negative_kernels.cc 28) + +add_test(NAME Unit_Device_Complex_Unary_float_Negative + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + complex_negative_kernels_1Arg_float.cc 62) + +add_test(NAME Unit_Device_Complex_Unary_double_Negative + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + complex_negative_kernels_1Arg_double.cc 62) + +add_test(NAME Unit_Device_Complex_Binary_float_Negative + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + complex_negative_kernels_2Arg_float.cc 88) + +add_test(NAME Unit_Device_Complex_Binary_double_Negative + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + complex_negative_kernels_2Arg_double.cc 88) + +add_test(NAME Unit_Device_Complex_hipCfma_Negative + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + complex_negative_kernels_3Arg.cc 60) + diff --git a/projects/hip-tests/catch/unit/complex/complex.cc b/projects/hip-tests/catch/unit/complex/complex.cc new file mode 100644 index 0000000000..11b42eb0e5 --- /dev/null +++ b/projects/hip-tests/catch/unit/complex/complex.cc @@ -0,0 +1,479 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "complex_function_common.hh" +#include "complex_cast_negative_kernels_rtc.hh" +#include "complex_make_negative_kernels_rtc.hh" +#include "complex_negative_kernels_1Arg_rtc.hh" +#include "complex_negative_kernels_2Arg_rtc.hh" +#include "complex_negative_kernels_3Arg_rtc.hh" + +/** + * @addtogroup complex complex + * @{ + * @ingroup ComplexTest + * Contains unit tests for complex type functions + */ + + +/** + * Test Description + * ------------------------ + * - Test that checks unary complex functions on device for reduced set of input values. The + * results are compared against manually calculated ones: + * -# hipConj, hipConjf + * -# hipCreal, hipCrealf + * -# hipCimag, hipCimagf + * -# hipCabs, hipCabsf + * -# hipCsqabs, hipCsqabsf + * + * Test source + * ------------------------ + * - unit/complex/complex.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_Device_Complex_Unary_Device_Sanity_Positive", "", hipFloatComplex, + hipDoubleComplex) { + decltype(TestType().x) input_r = GENERATE(-4.75, 0, 1.75); + decltype(TestType().x) input_i = GENERATE(-4.75, 0, 1.75); + + TestType input_val = MakeComplexType(input_r, input_i); + for (const auto function : + {ComplexFunction::kConj, ComplexFunction::kReal, ComplexFunction::kImag, + ComplexFunction::kAbs, ComplexFunction::kSqabs}) { + DYNAMIC_SECTION("function: " << to_string(function)) { + ComplexFunctionUnaryDeviceTest(function, input_val); + } + } +} + +/** + * Test Description + * ------------------------ + * - Test that checks unary complex functions on host for reduced set of input values. The + * results are compared against manually calculated ones: + * -# hipConj, hipConjf + * -# hipCreal, hipCrealf + * -# hipCimag, hipCimagf + * -# hipCabs, hipCabsf + * -# hipCsqabs, hipCsqabsf + * + * Test source + * ------------------------ + * - unit/complex/complex.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_Device_Complex_Unary_Host_Sanity_Positive", "", hipFloatComplex, + hipDoubleComplex) { + decltype(TestType().x) input_r = GENERATE(-4.75, 0, 1.75); + decltype(TestType().x) input_i = GENERATE(-4.75, 0, 1.75); + + TestType input_val = MakeComplexType(input_r, input_i); + for (const auto function : + {ComplexFunction::kConj, ComplexFunction::kReal, ComplexFunction::kImag, + ComplexFunction::kAbs, ComplexFunction::kSqabs}) { + DYNAMIC_SECTION("function: " << to_string(function)) { + ComplexFunctionUnaryHostTest(function, input_val); + } + } +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for unary complex functions. + * + * Test source + * ------------------------ + * - unit/complex/complex.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_Complex_Unary_Negative_Parameters_RTC") { + ComplexTypeRTCWrapper<28>(kComplexConj); + ComplexTypeRTCWrapper<24>(kComplexReal); + ComplexTypeRTCWrapper<24>(kComplexImag); + ComplexTypeRTCWrapper<24>(kComplexAbs); + ComplexTypeRTCWrapper<24>(kComplexSqabs); +} + +/** + * Test Description + * ------------------------ + * - Test that checks binary complex functions on device for reduced set of input values. The + * results are compared against manually calculated ones: + * -# hipCadd, hipCaddf + * -# hipCsub, hipCsubf + * -# hipCmul, hipCmulf + * -# hipCdiv, hipCdivf + * + * Test source + * ------------------------ + * - unit/complex/complex.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_Device_Complex_Binary_Device_Sanity_Positive", "", hipFloatComplex, + hipDoubleComplex) { + decltype(TestType().x) input1_r = GENERATE(-4.75, 0, 1.75); + decltype(TestType().x) input1_i = GENERATE(-4.75, 0, 1.75); + decltype(TestType().x) input2_r = GENERATE(-4.75, 0, 1.75); + decltype(TestType().x) input2_i = GENERATE(-4.75, 0, 1.75); + + TestType input_val1 = MakeComplexType(input1_r, input1_i); + TestType input_val2 = MakeComplexType(input2_r, input2_i); + for (const auto function : {ComplexFunction::kAdd, ComplexFunction::kSub, ComplexFunction::kMul, + ComplexFunction::kDiv}) { + DYNAMIC_SECTION("function: " << to_string(function)) { + ComplexFunctionBinaryDeviceTest(function, input_val1, input_val2); + } + } +} + +/** + * Test Description + * ------------------------ + * - Test that checks binary complex functions on host for reduced set of input values. The + * results are compared against manually calculated ones: + * -# hipCadd, hipCaddf + * -# hipCsub, hipCsubf + * -# hipCmul, hipCmulf + * -# hipCdiv, hipCdivf + * + * Test source + * ------------------------ + * - unit/complex/complex.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_Device_Complex_Binary_Host_Sanity_Positive", "", hipFloatComplex, + hipDoubleComplex) { + decltype(TestType().x) input1_r = GENERATE(-4.75, 0, 1.75); + decltype(TestType().x) input1_i = GENERATE(-4.75, 0, 1.75); + decltype(TestType().x) input2_r = GENERATE(-4.75, 0, 1.75); + decltype(TestType().x) input2_i = GENERATE(-4.75, 0, 1.75); + + TestType input_val1 = MakeComplexType(input1_r, input1_i); + TestType input_val2 = MakeComplexType(input2_r, input2_i); + for (const auto function : {ComplexFunction::kAdd, ComplexFunction::kSub, ComplexFunction::kMul, + ComplexFunction::kDiv}) { + DYNAMIC_SECTION("function: " << to_string(function)) { + ComplexFunctionBinaryHostTest(function, input_val1, input_val2); + } + } +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for binary complex functions. + * + * Test source + * ------------------------ + * - unit/complex/complex.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_Complex_Binary_Negative_Parameters_RTC") { + ComplexTypeRTCWrapper<44>(kComplexAdd); + ComplexTypeRTCWrapper<44>(kComplexSub); + ComplexTypeRTCWrapper<44>(kComplexMul); + ComplexTypeRTCWrapper<44>(kComplexDiv); +} + +/** + * Test Description + * ------------------------ + * - Test that checks hipCfma/hipCfmaf complex functions on device for reduced set of input + * values. The results are compared against manually calculated ones. + * + * Test source + * ------------------------ + * - unit/complex/complex.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_Device_Complex_hipCfma_Device_Sanity_Positive", "", hipFloatComplex, + hipDoubleComplex) { + decltype(TestType().x) input1_r = GENERATE(-4.75, 0, 1.75); + decltype(TestType().x) input1_i = GENERATE(-4.75, 0, 1.75); + decltype(TestType().x) input2_r = GENERATE(-4.75, 0, 1.75); + decltype(TestType().x) input2_i = GENERATE(-4.75, 0, 1.75); + decltype(TestType().x) input3_r = GENERATE(-4.75, 0, 1.75); + decltype(TestType().x) input3_i = GENERATE(-4.75, 0, 1.75); + + TestType input_val1 = MakeComplexType(input1_r, input1_i); + TestType input_val2 = MakeComplexType(input2_r, input2_i); + TestType input_val3 = MakeComplexType(input3_r, input3_i); + + ComplexFunctionTernaryDeviceTest(ComplexFunction::kFma, input_val1, input_val2, input_val3); +} + +/** + * Test Description + * ------------------------ + * - Test that checks hipCfma/hipCfmaf complex functions on host for reduced set of input + * values. The results are compared against manually calculated ones. + * + * Test source + * ------------------------ + * - unit/complex/complex.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_Device_Complex_hipCfma_Host_Sanity_Positive", "", hipFloatComplex, + hipDoubleComplex) { + decltype(TestType().x) input1_r = GENERATE(-4.75, 0, 1.75); + decltype(TestType().x) input1_i = GENERATE(-4.75, 0, 1.75); + decltype(TestType().x) input2_r = GENERATE(-4.75, 0, 1.75); + decltype(TestType().x) input2_i = GENERATE(-4.75, 0, 1.75); + decltype(TestType().x) input3_r = GENERATE(-4.75, 0, 1.75); + decltype(TestType().x) input3_i = GENERATE(-4.75, 0, 1.75); + + TestType input_val1 = MakeComplexType(input1_r, input1_i); + TestType input_val2 = MakeComplexType(input2_r, input2_i); + TestType input_val3 = MakeComplexType(input3_r, input3_i); + + ComplexFunctionTernaryHostTest(ComplexFunction::kFma, input_val1, input_val2, input_val3); +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for hipCfma/hipCfmaf complex function. + * + * Test source + * ------------------------ + * - unit/complex/complex.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_Complex_hipCfma_Negative_Parameters_RTC") { + ComplexTypeRTCWrapper<60>(kComplexFma); +} + +/** + * Test Description + * ------------------------ + * - Test that checks make_hipFloatComplex/make_hipDoubleComplex functions on device for reduced + * set of input values. The results are compared against manually calculated ones. + * + * Test source + * ------------------------ + * - unit/complex/complex.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_Device_make_Complex_Device_Positive", "", hipFloatComplex, + hipDoubleComplex) { + decltype(TestType().x) input_r = GENERATE(-0.25, 0, 0.25); + decltype(TestType().x) input_i = GENERATE(-1.75, 0, 1.75); + + LinearAllocGuard result_d(LinearAllocs::hipMalloc, sizeof(TestType)); + LinearAllocGuard result_h(LinearAllocs::hipHostMalloc, sizeof(TestType)); + + MakeComplexTypeKernel<<<1, 1>>>(result_d.ptr(), input_r, input_i); + HIP_CHECK(hipMemcpy(result_h.ptr(), result_d.ptr(), sizeof(TestType), hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + + REQUIRE(result_h.ptr()[0].x == input_r); + REQUIRE(result_h.ptr()[0].y == input_i); +} + +/** + * Test Description + * ------------------------ + * - Test that checks make_hipFloatComplex/make_hipDoubleComplex functions on device for reduced + * set of input values. The results are compared against manually calculated ones. + * + * Test source + * ------------------------ + * - unit/complex/complex.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_Device_make_Complex_Host_Positive", "", hipFloatComplex, + hipDoubleComplex) { + decltype(TestType().x) input_r = GENERATE(-0.25, 0, 0.25); + decltype(TestType().x) input_i = GENERATE(-1.75, 0, 1.75); + + TestType result = MakeComplexType(input_r, input_i); + + REQUIRE(result.x == input_r); + REQUIRE(result.y == input_i); +} + +#if HT_AMD // EXSWHTEC-321 +/** + * Test Description + * ------------------------ + * - Test that checks make_hipComplex functions on device for reduced set of input values. The + * results are compared against manually calculated ones. + * + * Test source + * ------------------------ + * - unit/complex/complex.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_make_hipComplex_Device_Positive") { + float input_r = GENERATE(-0.25, 0, 0.25); + float input_i = GENERATE(-1.75, 0, 1.75); + + LinearAllocGuard result_d(LinearAllocs::hipMalloc, sizeof(hipComplex)); + LinearAllocGuard result_h(LinearAllocs::hipHostMalloc, sizeof(hipComplex)); + + MakeHipComplexTypeKernel<<<1, 1>>>(result_d.ptr(), input_r, input_i); + HIP_CHECK(hipMemcpy(result_h.ptr(), result_d.ptr(), sizeof(hipComplex), hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + + REQUIRE(result_h.ptr()[0].x == input_r); + REQUIRE(result_h.ptr()[0].y == input_i); +} + +/** + * Test Description + * ------------------------ + * - Test that checks make_hipComplex functions on host for reduced set of input values. The + * results are compared against manually calculated ones. + * + * Test source + * ------------------------ + * - unit/complex/complex.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_make_hipComplex_Host_Positive") { + float input_r = GENERATE(-0.25, 0, 0.25); + float input_i = GENERATE(-1.75, 0, 1.75); + + hipComplex result = make_hipComplex(input_r, input_i); + + REQUIRE(result.x == input_r); + REQUIRE(result.y == input_i); +} +#endif + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for make complex functions. + * + * Test source + * ------------------------ + * - unit/complex/complex.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_make_Complex_Negative_Parameters_RTC") { + ComplexTypeRTCWrapper<18>(kMakeHipComplex); + ComplexTypeRTCWrapper<18>(kMakeHipFloatComplex); + ComplexTypeRTCWrapper<18>(kMakeHipDoubleComplex); +} + +/** + * Test Description + * ------------------------ + * - Test that checks hipComplexDoubleToFloat/hipComplexFloatToDouble functions on device for + * reduced set of input values. The results are compared against manually calculated ones. + * + * Test source + * ------------------------ + * - unit/complex/complex.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_Device_Complex_Cast_Device_Sanity_Positive", "", hipFloatComplex, + hipDoubleComplex) { + decltype(TestType().x) input_r = GENERATE(-0.25, 0, 0.25); + decltype(TestType().x) input_i = GENERATE(-1.75, 0, 1.75); + TestType input = MakeComplexType(input_r, input_i); + + LinearAllocGuard> result_d{LinearAllocs::hipMalloc, + sizeof(CastType_t)}; + LinearAllocGuard> result_h{LinearAllocs::hipHostMalloc, + sizeof(CastType_t)}; + + CastComplexTypeKernel<<<1, 1>>>(result_d.ptr(), input); + HIP_CHECK(hipMemcpy(result_h.ptr(), result_d.ptr(), sizeof(CastType_t), + hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + + REQUIRE(result_h.ptr()[0].x == static_cast().x)>(input_r)); + REQUIRE(result_h.ptr()[0].y == static_cast().x)>(input_i)); +} + +/** + * Test Description + * ------------------------ + * - Test that checks hipComplexDoubleToFloat/hipComplexFloatToDouble functions on host for + * reduced set of input values. The results are compared against manually calculated ones. + * + * Test source + * ------------------------ + * - unit/complex/complex.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_Device_Complex_Cast_Host_Sanity_Positive", "", hipFloatComplex, + hipDoubleComplex) { + decltype(TestType().x) input_r = GENERATE(-0.25, 0, 0.25); + decltype(TestType().x) input_i = GENERATE(-1.75, 0, 1.75); + TestType input = MakeComplexType(input_r, input_i); + + CastType_t result = CastComplexType>(input); + + REQUIRE(result.x == static_cast().x)>(input_r)); + REQUIRE(result.y == static_cast().x)>(input_i)); +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for complex cast functions. + * + * Test source + * ------------------------ + * - unit/complex/complex.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_Complex_Cast_Negative_Parameters_RTC") { + ComplexTypeRTCWrapper<14>(kComplexDoubleToFloat); + ComplexTypeRTCWrapper<14>(kComplexFloatToDouble); +} diff --git a/projects/hip-tests/catch/unit/complex/complex_basic_common.hh b/projects/hip-tests/catch/unit/complex/complex_basic_common.hh new file mode 100644 index 0000000000..f0e07d6cc4 --- /dev/null +++ b/projects/hip-tests/catch/unit/complex/complex_basic_common.hh @@ -0,0 +1,106 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#pragma once + +#include +#include +#include + +template +__host__ __device__ T MakeComplexType(decltype(T().x) input_val1, decltype(T().x) input_val2) { + if constexpr (std::is_same_v) { + return make_hipFloatComplex(input_val1, input_val2); + } else { + return make_hipDoubleComplex(input_val1, input_val2); + } +} + +template +__global__ void MakeComplexTypeKernel(T* const output_val, decltype(T().x) const input_val1, + decltype(T().x) const input_val2) { + *output_val = MakeComplexType(input_val1, input_val2); +} +#if HT_AMD // EXSWHTEC-321 +__global__ void MakeHipComplexTypeKernel(hipComplex* const output_val, float const input_val1, + float const input_val2) { + *output_val = make_hipComplex(input_val1, input_val2); +} +#endif +template struct CastType {}; + +template <> struct CastType { + using type = hipDoubleComplex; +}; + +template <> struct CastType { + using type = hipFloatComplex; +}; + +template using CastType_t = typename CastType::type; + +template __device__ __host__ T1 CastComplexType(T2 const input_val) { + if constexpr (std::is_same_v) { + return hipComplexDoubleToFloat(input_val); + } else if constexpr (std::is_same_v) { + return hipComplexFloatToDouble(input_val); + } +} + +template +__global__ void CastComplexTypeKernel(T1* const output_val, T2 const input_val) { + *output_val = CastComplexType(input_val); +} + +template void CompareValues(T actual_val, T ref_val, double margin) { + if (!std::isnan(ref_val)) { + REQUIRE_THAT(actual_val, Catch::WithinAbs(ref_val, margin)); + } +} + +template void ComplexTypeRTCWrapper(const char* program_source) { + hiprtcProgram program{}; + HIPRTC_CHECK(hiprtcCreateProgram(&program, program_source, "complex_type_kernels.cc", 0, nullptr, + nullptr)); + +#if HT_AMD + std::string args = std::string("-ferror-limit=100"); + const char* options[] = {args.c_str()}; + hiprtcResult result{hiprtcCompileProgram(program, 1, options)}; +#else + hiprtcResult result{hiprtcCompileProgram(program, 0, nullptr)}; +#endif + + size_t log_size{}; + HIPRTC_CHECK(hiprtcGetProgramLogSize(program, &log_size)); + std::string log(log_size, ' '); + HIPRTC_CHECK(hiprtcGetProgramLog(program, log.data())); + int error_count{0}; + + std::string error_message{"error:"}; + + size_t npos_e = log.find(error_message, 0); + while (npos_e != std::string::npos) { + ++error_count; + npos_e = log.find(error_message, npos_e + 1); + } + + HIPRTC_CHECK(hiprtcDestroyProgram(&program)); + HIPRTC_CHECK_ERROR(result, HIPRTC_ERROR_COMPILATION); + REQUIRE(error_count == expected_errors_num); +} diff --git a/projects/hip-tests/catch/unit/complex/complex_cast_negative_kernels.cc b/projects/hip-tests/catch/unit/complex/complex_cast_negative_kernels.cc new file mode 100644 index 0000000000..23b0700205 --- /dev/null +++ b/projects/hip-tests/catch/unit/complex/complex_cast_negative_kernels.cc @@ -0,0 +1,113 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include + +class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} +}; + +__global__ void hipComplexDoubleToFloat_kernel_v1(hipFloatComplex* result, hipDoubleComplex* x) { + *result = hipComplexDoubleToFloat(x); +} +__global__ void hipComplexDoubleToFloat_kernel_v2(hipFloatComplex* result, hipFloatComplex x) { + *result = hipComplexDoubleToFloat(x); +} +__global__ void hipComplexDoubleToFloat_kernel_v3(hipFloatComplex* result, double x) { + *result = hipComplexDoubleToFloat(x); +} +__global__ void hipComplexDoubleToFloat_kernel_v4(hipFloatComplex* result, Dummy x) { + *result = hipComplexDoubleToFloat(x); +} +__global__ void hipComplexDoubleToFloat_kernel_v5(float* result, hipDoubleComplex x) { + *result = hipComplexDoubleToFloat(x); +} +__global__ void hipComplexDoubleToFloat_kernel_v6(hipDoubleComplex* result, hipDoubleComplex x) { + *result = hipComplexDoubleToFloat(x); +} +__global__ void hipComplexDoubleToFloat_kernel_v7(Dummy* result, hipDoubleComplex x) { + *result = hipComplexDoubleToFloat(x); +} +void hipComplexDoubleToFloat_v1(hipFloatComplex* result, hipDoubleComplex* x) { + *result = hipComplexDoubleToFloat(x); +} +void hipComplexDoubleToFloat_v2(hipFloatComplex* result, hipFloatComplex x) { + *result = hipComplexDoubleToFloat(x); +} +void hipComplexDoubleToFloat_v3(hipFloatComplex* result, double x) { + *result = hipComplexDoubleToFloat(x); +} +void hipComplexDoubleToFloat_v4(hipFloatComplex* result, Dummy x) { + *result = hipComplexDoubleToFloat(x); +} +void hipComplexDoubleToFloat_v5(float* result, hipDoubleComplex x) { + *result = hipComplexDoubleToFloat(x); +} +void hipComplexDoubleToFloat_v6(hipDoubleComplex* result, hipDoubleComplex x) { + *result = hipComplexDoubleToFloat(x); +} +void hipComplexDoubleToFloat_v7(Dummy* result, hipDoubleComplex x) { + *result = hipComplexDoubleToFloat(x); +} + +__global__ void hipComplexFloatToDouble_kernel_v1(hipDoubleComplex* result, hipFloatComplex* x) { + *result = hipComplexFloatToDouble(x); +} +__global__ void hipComplexFloatToDouble_kernel_v2(hipDoubleComplex* result, hipDoubleComplex x) { + *result = hipComplexFloatToDouble(x); +} +__global__ void hipComplexFloatToDouble_kernel_v3(hipDoubleComplex* result, float x) { + *result = hipComplexFloatToDouble(x); +} +__global__ void hipComplexFloatToDouble_kernel_v4(hipDoubleComplex* result, Dummy x) { + *result = hipComplexFloatToDouble(x); +} +__global__ void hipComplexFloatToDouble_kernel_v5(double* result, hipFloatComplex x) { + *result = hipComplexFloatToDouble(x); +} +__global__ void hipComplexFloatToDouble_kernel_v6(hipFloatComplex* result, hipFloatComplex x) { + *result = hipComplexFloatToDouble(x); +} +__global__ void hipComplexFloatToDouble_kernel_v7(Dummy* result, hipFloatComplex x) { + *result = hipComplexFloatToDouble(x); +} +void hipComplexFloatToDouble_v1(hipDoubleComplex* result, hipFloatComplex* x) { + *result = hipComplexFloatToDouble(x); +} +void hipComplexFloatToDouble_v2(hipDoubleComplex* result, hipDoubleComplex x) { + *result = hipComplexFloatToDouble(x); +} +void hipComplexFloatToDouble_v3(hipDoubleComplex* result, float x) { + *result = hipComplexFloatToDouble(x); +} +void hipComplexFloatToDouble_v4(hipDoubleComplex* result, Dummy x) { + *result = hipComplexFloatToDouble(x); +} +void hipComplexFloatToDouble_v5(double* result, hipFloatComplex x) { + *result = hipComplexFloatToDouble(x); +} +void hipComplexFloatToDouble_v6(hipFloatComplex* result, hipFloatComplex x) { + *result = hipComplexFloatToDouble(x); +} +void hipComplexFloatToDouble_v7(Dummy* result, hipFloatComplex x) { + *result = hipComplexFloatToDouble(x); +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/complex/complex_cast_negative_kernels_rtc.hh b/projects/hip-tests/catch/unit/complex/complex_cast_negative_kernels_rtc.hh new file mode 100644 index 0000000000..c2ef7bef8d --- /dev/null +++ b/projects/hip-tests/catch/unit/complex/complex_cast_negative_kernels_rtc.hh @@ -0,0 +1,121 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +static constexpr auto kComplexDoubleToFloat{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void hipComplexDoubleToFloat_kernel_v1(hipFloatComplex* result, hipDoubleComplex* x) { + *result = hipComplexDoubleToFloat(x); + } + __global__ void hipComplexDoubleToFloat_kernel_v2(hipFloatComplex* result, hipFloatComplex x) { + *result = hipComplexDoubleToFloat(x); + } + __global__ void hipComplexDoubleToFloat_kernel_v3(hipFloatComplex* result, double x) { + *result = hipComplexDoubleToFloat(x); + } + __global__ void hipComplexDoubleToFloat_kernel_v4(hipFloatComplex* result, Dummy x) { + *result = hipComplexDoubleToFloat(x); + } + __global__ void hipComplexDoubleToFloat_kernel_v5(float* result, hipDoubleComplex x) { + *result = hipComplexDoubleToFloat(x); + } + __global__ void hipComplexDoubleToFloat_kernel_v6(hipDoubleComplex* result, hipDoubleComplex x) { + *result = hipComplexDoubleToFloat(x); + } + __global__ void hipComplexDoubleToFloat_kernel_v7(Dummy* result, hipDoubleComplex x) { + *result = hipComplexDoubleToFloat(x); + } + void hipComplexDoubleToFloat_v1(hipFloatComplex* result, hipDoubleComplex* x) { + *result = hipComplexDoubleToFloat(x); + } + void hipComplexDoubleToFloat_v2(hipFloatComplex* result, hipFloatComplex x) { + *result = hipComplexDoubleToFloat(x); + } + void hipComplexDoubleToFloat_v3(hipFloatComplex* result, double x) { + *result = hipComplexDoubleToFloat(x); + } + void hipComplexDoubleToFloat_v4(hipFloatComplex* result, Dummy x) { + *result = hipComplexDoubleToFloat(x); + } + void hipComplexDoubleToFloat_v5(float* result, hipDoubleComplex x) { + *result = hipComplexDoubleToFloat(x); + } + void hipComplexDoubleToFloat_v6(hipDoubleComplex* result, hipDoubleComplex x) { + *result = hipComplexDoubleToFloat(x); + } + void hipComplexDoubleToFloat_v7(Dummy* result, hipDoubleComplex x) { + *result = hipComplexDoubleToFloat(x); + } + +)"}; + +static constexpr auto kComplexFloatToDouble{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void hipComplexFloatToDouble_kernel_v1(hipDoubleComplex* result, hipFloatComplex* x) { + *result = hipComplexFloatToDouble(x); + } + __global__ void hipComplexFloatToDouble_kernel_v2(hipDoubleComplex* result, hipDoubleComplex x) { + *result = hipComplexFloatToDouble(x); + } + __global__ void hipComplexFloatToDouble_kernel_v3(hipDoubleComplex* result, float x) { + *result = hipComplexFloatToDouble(x); + } + __global__ void hipComplexFloatToDouble_kernel_v4(hipDoubleComplex* result, Dummy x) { + *result = hipComplexFloatToDouble(x); + } +__global__ void hipComplexFloatToDouble_kernel_v5(double* result, hipFloatComplex x) { + *result = hipComplexFloatToDouble(x); +} + __global__ void hipComplexFloatToDouble_kernel_v6(hipFloatComplex* result, hipFloatComplex x) { + *result = hipComplexFloatToDouble(x); + } + __global__ void hipComplexFloatToDouble_kernel_v7(Dummy* result, hipFloatComplex x) { + *result = hipComplexFloatToDouble(x); + } + void hipComplexFloatToDouble_v1(hipDoubleComplex* result, hipFloatComplex* x) { + *result = hipComplexFloatToDouble(x); + } + void hipComplexFloatToDouble_v2(hipDoubleComplex* result, hipDoubleComplex x) { + *result = hipComplexFloatToDouble(x); + } + void hipComplexFloatToDouble_v3(hipDoubleComplex* result, float x) { + *result = hipComplexFloatToDouble(x); + } + void hipComplexFloatToDouble_v4(hipDoubleComplex* result, Dummy x) { + *result = hipComplexFloatToDouble(x); + } + void hipComplexFloatToDouble_v5(double* result, hipFloatComplex x) { + *result = hipComplexFloatToDouble(x); + } + void hipComplexFloatToDouble_v6(hipFloatComplex* result, hipFloatComplex x) { + *result = hipComplexFloatToDouble(x); + } + void hipComplexFloatToDouble_v7(Dummy* result, hipFloatComplex x) { + *result = hipComplexFloatToDouble(x); + } +)"}; \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/complex/complex_function_common.hh b/projects/hip-tests/catch/unit/complex/complex_function_common.hh new file mode 100644 index 0000000000..2cecfc9b2d --- /dev/null +++ b/projects/hip-tests/catch/unit/complex/complex_function_common.hh @@ -0,0 +1,311 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#pragma once + +#include "complex_basic_common.hh" + +enum class ComplexFunction { kReal, kImag, kConj, kAdd, kSub, kMul, kDiv, kAbs, kSqabs, kFma }; + +inline std::string to_string(ComplexFunction function) { + switch (function) { + case ComplexFunction::kReal: + return "real"; + case ComplexFunction::kImag: + return "imaginary"; + case ComplexFunction::kConj: + return "conjugate "; + case ComplexFunction::kAdd: + return "addition"; + case ComplexFunction::kSub: + return "subtract"; + case ComplexFunction::kMul: + return "multiply"; + case ComplexFunction::kDiv: + return "divide"; + case ComplexFunction::kAbs: + return "absolute"; + case ComplexFunction::kSqabs: + return "square absolute"; + case ComplexFunction::kFma: + return "fused multiply"; + default: + return "Unknown"; + } +} + +// Function that validates complex functions with complex type result +template +void ValidateComplexResultFunction(ComplexFunction function, T input_val1, T input_val2, + T input_val3, T actual_val) { + decltype(T().x) ref_val_r; + decltype(T().x) ref_val_i; + double margin = 0; + + switch (function) { + case ComplexFunction::kAdd: { + ref_val_r = input_val1.x + input_val2.x; + ref_val_i = input_val1.y + input_val2.y; + break; + } + case ComplexFunction::kSub: { + ref_val_r = input_val1.x - input_val2.x; + ref_val_i = input_val1.y - input_val2.y; + break; + } + case ComplexFunction::kMul: { + ref_val_r = input_val1.x * input_val2.x - input_val1.y * input_val2.y; + ref_val_i = input_val1.y * input_val2.x + input_val1.x * input_val2.y; + break; + } + case ComplexFunction::kDiv: { + decltype(T().x) sqabs = input_val2.x * input_val2.x + input_val2.y * input_val2.y; + ref_val_r = (input_val1.x * input_val2.x + input_val1.y * input_val2.y) / sqabs; + ref_val_i = (input_val1.y * input_val2.x - input_val1.x * input_val2.y) / sqabs; +#if HT_NVIDIA + // Nvidia implementation uses scaling to guard against intermediate underflow and overflow + margin = 0.000001; +#endif + break; + } + case ComplexFunction::kConj: { + ref_val_r = input_val1.x; + ref_val_i = -input_val1.y; + break; + } + case ComplexFunction::kFma: { + ref_val_r = (input_val1.x * input_val2.x) + input_val3.x; + ref_val_i = (input_val2.x * input_val1.y) + input_val3.y; + + ref_val_r = -(input_val1.y * input_val2.y) + ref_val_r; + ref_val_i = (input_val1.x * input_val2.y) + ref_val_i; + break; + } + default: { + ref_val_r = input_val1.x; + ref_val_i = input_val1.y; + break; + } + } + + CompareValues(actual_val.x, ref_val_r, margin); + CompareValues(actual_val.y, ref_val_i, margin); +} + +// Function that validates complex functions with scalar type result +template +void ValidateScalarResultFunction(ComplexFunction function, T input_val, + decltype(T().x) actual_val) { + decltype(T().x) ref_val; + + switch (function) { + case ComplexFunction::kReal: { + ref_val = input_val.x; + break; + } + case ComplexFunction::kImag: { + ref_val = input_val.y; + break; + } + case ComplexFunction::kAbs: { + decltype(T().x) sqabs = input_val.x * input_val.x + input_val.y * input_val.y; + ref_val = std::sqrt(sqabs); + break; + } + case ComplexFunction::kSqabs: { + ref_val = input_val.x * input_val.x + input_val.y * input_val.y; + break; + } + default: { + ref_val = input_val.x; + break; + } + } + + CompareValues(actual_val, ref_val, 0); +} + +// Function that performs complex functions with complex type result on host/device +template +__device__ __host__ void PerformComplexResultFunction(ComplexFunction function, T* output_val, + T input_val1, T input_val2, T input_val3) { + if (function == ComplexFunction::kAdd) { + if constexpr (std::is_same_v) { + *output_val = hipCaddf(input_val1, input_val2); + } else if constexpr (std::is_same_v) { + *output_val = hipCadd(input_val1, input_val2); + } + } else if (function == ComplexFunction::kSub) { + if constexpr (std::is_same_v) { + *output_val = hipCsubf(input_val1, input_val2); + } else if constexpr (std::is_same_v) { + *output_val = hipCsub(input_val1, input_val2); + } + } else if (function == ComplexFunction::kMul) { + if constexpr (std::is_same_v) { + *output_val = hipCmulf(input_val1, input_val2); + } else if constexpr (std::is_same_v) { + *output_val = hipCmul(input_val1, input_val2); + } + } else if (function == ComplexFunction::kDiv) { + if constexpr (std::is_same_v) { + *output_val = hipCdivf(input_val1, input_val2); + } else if constexpr (std::is_same_v) { + *output_val = hipCdiv(input_val1, input_val2); + } + } else if (function == ComplexFunction::kConj) { + if constexpr (std::is_same_v) { + *output_val = hipConjf(input_val1); + } else if constexpr (std::is_same_v) { + *output_val = hipConj(input_val1); + } + } else if (function == ComplexFunction::kFma) { + if constexpr (std::is_same_v) { + *output_val = hipCfmaf(input_val1, input_val2, input_val3); + } else if constexpr (std::is_same_v) { + *output_val = hipCfma(input_val1, input_val2, input_val3); + } + } else { + *output_val = input_val1; + } +} + +// Function that performs complex functions with scalar type result on host/device +template +__device__ __host__ void PerformScalarResultFunction(ComplexFunction function, + decltype(T().x)* output_val, T input_val) { + if (function == ComplexFunction::kReal) { + if constexpr (std::is_same_v) { + *output_val = hipCrealf(input_val); + } else if constexpr (std::is_same_v) { + *output_val = hipCreal(input_val); + } + } else if (function == ComplexFunction::kImag) { + if constexpr (std::is_same_v) { + *output_val = hipCimagf(input_val); + } else if constexpr (std::is_same_v) { + *output_val = hipCimag(input_val); + } + } else if (function == ComplexFunction::kAbs) { + if constexpr (std::is_same_v) { + *output_val = hipCabsf(input_val); + } else if constexpr (std::is_same_v) { + *output_val = hipCabs(input_val); + } + } else if (function == ComplexFunction::kSqabs) { + if constexpr (std::is_same_v) { + *output_val = hipCsqabsf(input_val); + } else if constexpr (std::is_same_v) { + *output_val = hipCsqabs(input_val); + } + } else { + *output_val = input_val.x; + } +} + +// Kernel that calls device function which performs complex functions with complex type result +template +__global__ void ComplexResultKernel(ComplexFunction function, T* output_val, T input_val1, + T input_val2, T input_val3) { + PerformComplexResultFunction(function, output_val, input_val1, input_val2, input_val3); +} + +// Kernel that calls device function which performs complex functions with scalar type result +template +__global__ void ScalarResultKernel(ComplexFunction function, decltype(T().x)* output_val, + T input_val) { + PerformScalarResultFunction(function, output_val, input_val); +} + +// Wrapper function for testing complex functions with one input parameter on device +template void ComplexFunctionUnaryDeviceTest(ComplexFunction function, T input_val) { + if (function == ComplexFunction::kConj) { + LinearAllocGuard result_d{LinearAllocs::hipMalloc, sizeof(T)}; + LinearAllocGuard result_h{LinearAllocs::hipHostMalloc, sizeof(T)}; + + ComplexResultKernel<<<1, 1>>>(function, result_d.ptr(), input_val, input_val, input_val); + HIP_CHECK(hipMemcpy(result_h.ptr(), result_d.ptr(), sizeof(T), hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + ValidateComplexResultFunction(function, input_val, input_val, input_val, result_h.ptr()[0]); + } else { + LinearAllocGuard result_d{LinearAllocs::hipMalloc, sizeof(decltype(T().x))}; + LinearAllocGuard result_h{LinearAllocs::hipHostMalloc, + sizeof(decltype(T().x))}; + + ScalarResultKernel<<<1, 1>>>(function, result_d.ptr(), input_val); + HIP_CHECK( + hipMemcpy(result_h.ptr(), result_d.ptr(), sizeof(decltype(T().x)), hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + ValidateScalarResultFunction(function, input_val, result_h.ptr()[0]); + } +} + +// Wrapper function for testing complex functions with one input parameter on host +template void ComplexFunctionUnaryHostTest(ComplexFunction function, T input_val) { + if (function == ComplexFunction::kConj) { + T result; + PerformComplexResultFunction(function, &result, input_val, input_val, input_val); + ValidateComplexResultFunction(function, input_val, input_val, input_val, result); + } else { + decltype(T().x) result; + PerformScalarResultFunction(function, &result, input_val); + ValidateScalarResultFunction(function, input_val, result); + } +} + +// Wrapper function for testing complex functions with two input parameters on device +template +void ComplexFunctionBinaryDeviceTest(ComplexFunction function, T input_val1, T input_val2) { + LinearAllocGuard result_d{LinearAllocs::hipMalloc, sizeof(T)}; + LinearAllocGuard result_h{LinearAllocs::hipHostMalloc, sizeof(T)}; + + ComplexResultKernel<<<1, 1>>>(function, result_d.ptr(), input_val1, input_val2, input_val2); + HIP_CHECK(hipMemcpy(result_h.ptr(), result_d.ptr(), sizeof(T), hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + ValidateComplexResultFunction(function, input_val1, input_val2, input_val2, result_h.ptr()[0]); +} + +// Wrapper function for testing complex functions with two input parameters on host +template +void ComplexFunctionBinaryHostTest(ComplexFunction function, T input_val1, T input_val2) { + T result; + PerformComplexResultFunction(function, &result, input_val1, input_val2, input_val2); + ValidateComplexResultFunction(function, input_val1, input_val2, input_val2, result); +} + +// Wrapper function for testing complex functions with three input parameters on device +template +void ComplexFunctionTernaryDeviceTest(ComplexFunction function, T input_val1, T input_val2, + T input_val3) { + LinearAllocGuard result_d{LinearAllocs::hipMalloc, sizeof(T)}; + LinearAllocGuard result_h{LinearAllocs::hipHostMalloc, sizeof(T)}; + + ComplexResultKernel<<<1, 1>>>(function, result_d.ptr(), input_val1, input_val2, input_val3); + HIP_CHECK(hipMemcpy(result_h.ptr(), result_d.ptr(), sizeof(T), hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + ValidateComplexResultFunction(function, input_val1, input_val2, input_val3, result_h.ptr()[0]); +} + +// Wrapper function for testing complex functions with three input parameters on host +template +void ComplexFunctionTernaryHostTest(ComplexFunction function, T input_val1, T input_val2, + T input_val3) { + T result; + PerformComplexResultFunction(function, &result, input_val1, input_val2, input_val3); + ValidateComplexResultFunction(function, input_val1, input_val2, input_val3, result); +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/complex/complex_make_negative_kernels.cc b/projects/hip-tests/catch/unit/complex/complex_make_negative_kernels.cc new file mode 100644 index 0000000000..e18943f6f3 --- /dev/null +++ b/projects/hip-tests/catch/unit/complex/complex_make_negative_kernels.cc @@ -0,0 +1,121 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include + +class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} +}; + +#define NEGATIVE_SHELL_MAKE_FLOAT(T, func_name) \ + __global__ void func_name##_kernel_v1(T* result, float* x, float y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v2(T* result, float x, float* y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v3(T* result, T x, float y) { *result = func_name(x, y); } \ + __global__ void func_name##_kernel_v4(T* result, float x, T y) { *result = func_name(x, y); } \ + __global__ void func_name##_kernel_v5(T* result, Dummy x, float y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v6(T* result, float x, Dummy y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v7(float* result, float x, float y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v8(hipDoubleComplex* result, float x, float y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v9(Dummy* result, float x, float y) { \ + *result = func_name(x, y); \ + } \ + void func_name##_v1(T* result, float* x, float y) { *result = func_name(x, y); } \ + void func_name##_v2(T* result, float x, float* y) { *result = func_name(x, y); } \ + void func_name##_v3(T* result, T x, float y) { *result = func_name(x, y); } \ + void func_name##_v4(T* result, float x, T y) { *result = func_name(x, y); } \ + void func_name##_v5(T* result, Dummy x, float y) { *result = func_name(x, y); } \ + void func_name##_v6(T* result, float x, Dummy y) { *result = func_name(x, y); } \ + void func_name##_v7(float* result, float x, float y) { *result = func_name(x, y); } \ + void func_name##_v8(hipDoubleComplex* result, float x, float y) { *result = func_name(x, y); } \ + void func_name##_v9(Dummy* result, float x, float y) { *result = func_name(x, y); } + +__global__ void make_hipDoubleComplex_kernel_v1(hipDoubleComplex* result, double* x, double y) { + *result = make_hipDoubleComplex(x, y); +} +__global__ void make_hipDoubleComplex_kernel_v2(hipDoubleComplex* result, double x, double* y) { + *result = make_hipDoubleComplex(x, y); +} +__global__ void make_hipDoubleComplex_kernel_v3(hipDoubleComplex* result, hipDoubleComplex x, + double y) { + *result = make_hipDoubleComplex(x, y); +} +__global__ void make_hipDoubleComplex_kernel_v4(hipDoubleComplex* result, double x, + hipDoubleComplex y) { + *result = make_hipDoubleComplex(x, y); +} +__global__ void make_hipDoubleComplex_kernel_v5(hipDoubleComplex* result, Dummy x, double y) { + *result = make_hipDoubleComplex(x, y); +} +__global__ void make_hipDoubleComplex_kernel_v6(hipDoubleComplex* result, double x, Dummy y) { + *result = make_hipDoubleComplex(x, y); +} +__global__ void make_hipDoubleComplex_kernel_v7(double* result, double x, double y) { + *result = make_hipDoubleComplex(x, y); +} +__global__ void make_hipDoubleComplex_kernel_v8(hipFloatComplex* result, double x, double y) { + *result = make_hipDoubleComplex(x, y); +} +__global__ void make_hipDoubleComplex_kernel_v9(Dummy* result, double x, double y) { + *result = make_hipDoubleComplex(x, y); +} +void make_hipDoubleComplex_v1(hipDoubleComplex* result, double* x, double y) { + *result = make_hipDoubleComplex(x, y); +} +void make_hipDoubleComplex_v2(hipDoubleComplex* result, double x, double* y) { + *result = make_hipDoubleComplex(x, y); +} +void make_hipDoubleComplex_v3(hipDoubleComplex* result, hipDoubleComplex x, double y) { + *result = make_hipDoubleComplex(x, y); +} +void make_hipDoubleComplex_v4(hipDoubleComplex* result, double x, hipDoubleComplex y) { + *result = make_hipDoubleComplex(x, y); +} +void make_hipDoubleComplex_v5(hipDoubleComplex* result, Dummy x, double y) { + *result = make_hipDoubleComplex(x, y); +} +void make_hipDoubleComplex_v6(hipDoubleComplex* result, double x, Dummy y) { + *result = make_hipDoubleComplex(x, y); +} +void make_hipDoubleComplex_v7(float* result, double x, double y) { + *result = make_hipDoubleComplex(x, y); +} +void make_hipDoubleComplex_v8(hipFloatComplex* result, double x, double y) { + *result = make_hipDoubleComplex(x, y); +} +void make_hipDoubleComplex_v9(Dummy* result, double x, double y) { + *result = make_hipDoubleComplex(x, y); +} + +NEGATIVE_SHELL_MAKE_FLOAT(hipFloatComplex, make_hipFloatComplex) +NEGATIVE_SHELL_MAKE_FLOAT(hipComplex, make_hipComplex) \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/complex/complex_make_negative_kernels_rtc.hh b/projects/hip-tests/catch/unit/complex/complex_make_negative_kernels_rtc.hh new file mode 100644 index 0000000000..c643d02c09 --- /dev/null +++ b/projects/hip-tests/catch/unit/complex/complex_make_negative_kernels_rtc.hh @@ -0,0 +1,198 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +static constexpr auto kMakeHipComplex{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void make_hipComplex_kernel_v1(hipComplex* result, float* x, float y) { + *result = make_hipComplex(x, y); + } + __global__ void make_hipComplex_kernel_v2(hipComplex* result, float x, float* y) { + *result = make_hipComplex(x, y); + } + __global__ void make_hipComplex_kernel_v3(hipComplex* result, hipComplex x, float y) { + *result = make_hipComplex(x, y); + } + __global__ void make_hipComplex_kernel_v4(hipComplex* result, float x, hipComplex y) { + *result = make_hipComplex(x, y); + } + __global__ void make_hipComplex_kernel_v5(hipComplex* result, Dummy x, float y) { + *result = make_hipComplex(x, y); + } + __global__ void make_hipComplex_kernel_v6(hipComplex* result, float x, Dummy y) { + *result = make_hipComplex(x, y); + } + __global__ void make_hipComplex_kernel_v7(float* result, float x, float y) { + *result = make_hipComplex(x, y); + } + __global__ void make_hipComplex_kernel_v8(hipDoubleComplex* result, float x, float y) { + *result = make_hipComplex(x, y); + } + __global__ void make_hipComplex_kernel_v9(Dummy* result, float x, float y) { + *result = make_hipComplex(x, y); + } + void make_hipComplex_v1(hipComplex* result, float* x, float y) { *result = make_hipComplex(x, y); } + void make_hipComplex_v2(hipComplex* result, float x, float* y) { *result = make_hipComplex(x, y); } + void make_hipComplex_v3(hipComplex* result, hipComplex x, float y) { + *result = make_hipComplex(x, y); + } + void make_hipComplex_v4(hipComplex* result, float x, hipComplex y) { + *result = make_hipComplex(x, y); + } + void make_hipComplex_v5(hipComplex* result, Dummy x, float y) { *result = make_hipComplex(x, y); } + void make_hipComplex_v6(hipComplex* result, float x, Dummy y) { *result = make_hipComplex(x, y); } + void make_hipComplex_v7(float* result, float x, float y) { *result = make_hipComplex(x, y); } + void make_hipComplex_v8(hipDoubleComplex* result, float x, float y) { + *result = make_hipComplex(x, y); + } + void make_hipComplex_v9(Dummy* result, float x, float y) { *result = make_hipComplex(x, y); } +)"}; + +static constexpr auto kMakeHipFloatComplex{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void make_hipFloatComplex_kernel_v1(hipFloatComplex* result, float* x, float y) { + *result = make_hipFloatComplex(x, y); + } + __global__ void make_hipFloatComplex_kernel_v2(hipFloatComplex* result, float x, float* y) { + *result = make_hipFloatComplex(x, y); + } + __global__ void make_hipFloatComplex_kernel_v3(hipFloatComplex* result, hipFloatComplex x, + float y) { + *result = make_hipFloatComplex(x, y); + } + __global__ void make_hipFloatComplex_kernel_v4(hipFloatComplex* result, float x, + hipFloatComplex y) { + *result = make_hipFloatComplex(x, y); + } + __global__ void make_hipFloatComplex_kernel_v5(hipFloatComplex* result, Dummy x, float y) { + *result = make_hipFloatComplex(x, y); + } + __global__ void make_hipFloatComplex_kernel_v6(hipFloatComplex* result, float x, Dummy y) { + *result = make_hipFloatComplex(x, y); + } + __global__ void make_hipFloatComplex_kernel_v7(float* result, float x, float y) { + *result = make_hipFloatComplex(x, y); + } + __global__ void make_hipFloatComplex_kernel_v8(hipDoubleComplex* result, float x, float y) { + *result = make_hipFloatComplex(x, y); + } + __global__ void make_hipFloatComplex_kernel_v9(Dummy* result, float x, float y) { + *result = make_hipFloatComplex(x, y); + } + void make_hipFloatComplex_v1(hipFloatComplex* result, float* x, float y) { + *result = make_hipFloatComplex(x, y); + } + void make_hipFloatComplex_v2(hipFloatComplex* result, float x, float* y) { + *result = make_hipFloatComplex(x, y); + } + void make_hipFloatComplex_v3(hipFloatComplex* result, hipFloatComplex x, float y) { + *result = make_hipFloatComplex(x, y); + } + void make_hipFloatComplex_v4(hipFloatComplex* result, float x, hipFloatComplex y) { + *result = make_hipFloatComplex(x, y); + } + void make_hipFloatComplex_v5(hipFloatComplex* result, Dummy x, float y) { + *result = make_hipFloatComplex(x, y); + } + void make_hipFloatComplex_v6(hipFloatComplex* result, float x, Dummy y) { + *result = make_hipFloatComplex(x, y); + } + void make_hipFloatComplex_v7(float* result, float x, float y) { + *result = make_hipFloatComplex(x, y); + } + void make_hipFloatComplex_v8(hipDoubleComplex* result, float x, float y) { + *result = make_hipFloatComplex(x, y); + } + void make_hipFloatComplex_v9(Dummy* result, float x, float y) { + *result = make_hipFloatComplex(x, y); + } +)"}; + +static constexpr auto kMakeHipDoubleComplex{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void make_hipDoubleComplex_kernel_v1(hipDoubleComplex* result, double* x, double y) { + *result = make_hipDoubleComplex(x, y); + } + __global__ void make_hipDoubleComplex_kernel_v2(hipDoubleComplex* result, double x, double* y) { + *result = make_hipDoubleComplex(x, y); + } + __global__ void make_hipDoubleComplex_kernel_v3(hipDoubleComplex* result, hipDoubleComplex x, + double y) { + *result = make_hipDoubleComplex(x, y); + } + __global__ void make_hipDoubleComplex_kernel_v4(hipDoubleComplex* result, double x, + hipDoubleComplex y) { + *result = make_hipDoubleComplex(x, y); + } + __global__ void make_hipDoubleComplex_kernel_v5(hipDoubleComplex* result, Dummy x, double y) { + *result = make_hipDoubleComplex(x, y); + } + __global__ void make_hipDoubleComplex_kernel_v6(hipDoubleComplex* result, double x, Dummy y) { + *result = make_hipDoubleComplex(x, y); + } + __global__ void make_hipDoubleComplex_kernel_v7(double* result, double x, double y) { + *result = make_hipDoubleComplex(x, y); + } + __global__ void make_hipDoubleComplex_kernel_v8(hipFloatComplex* result, double x, double y) { + *result = make_hipDoubleComplex(x, y); + } + __global__ void make_hipDoubleComplex_kernel_v9(Dummy* result, double x, double y) { + *result = make_hipDoubleComplex(x, y); + } + void make_hipDoubleComplex_v1(hipDoubleComplex* result, double* x, double y) { + *result = make_hipDoubleComplex(x, y); + } + void make_hipDoubleComplex_v2(hipDoubleComplex* result, double x, double* y) { + *result = make_hipDoubleComplex(x, y); + } + void make_hipDoubleComplex_v3(hipDoubleComplex* result, hipDoubleComplex x, double y) { + *result = make_hipDoubleComplex(x, y); + } + void make_hipDoubleComplex_v4(hipDoubleComplex* result, double x, hipDoubleComplex y) { + *result = make_hipDoubleComplex(x, y); + } + void make_hipDoubleComplex_v5(hipDoubleComplex* result, Dummy x, double y) { + *result = make_hipDoubleComplex(x, y); + } + void make_hipDoubleComplex_v6(hipDoubleComplex* result, double x, Dummy y) { + *result = make_hipDoubleComplex(x, y); + } + void make_hipDoubleComplex_v7(float* result, double x, double y) { + *result = make_hipDoubleComplex(x, y); + } + void make_hipDoubleComplex_v8(hipFloatComplex* result, double x, double y) { + *result = make_hipDoubleComplex(x, y); + } + void make_hipDoubleComplex_v9(Dummy* result, double x, double y) { + *result = make_hipDoubleComplex(x, y); + } +)"}; diff --git a/projects/hip-tests/catch/unit/complex/complex_negative_kernels_1Arg_double.cc b/projects/hip-tests/catch/unit/complex/complex_negative_kernels_1Arg_double.cc new file mode 100644 index 0000000000..31ff6ce5b8 --- /dev/null +++ b/projects/hip-tests/catch/unit/complex/complex_negative_kernels_1Arg_double.cc @@ -0,0 +1,75 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include + +class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} +}; + +#define NEGATIVE_SHELL_ONE_ARG_DOUBLE(func_name) \ + __global__ void func_name##_kernel_v1(double* result, hipDoubleComplex* x) { \ + *result = func_name(x); \ + } \ + __global__ void func_name##_kernel_v2(double* result, hipFloatComplex x) { \ + *result = func_name(x); \ + } \ + __global__ void func_name##_kernel_v3(double* result, double x) { *result = func_name(x); } \ + __global__ void func_name##_kernel_v4(double* result, Dummy x) { *result = func_name(x); } \ + __global__ void func_name##_kernel_v5(hipDoubleComplex* result, hipDoubleComplex x) { \ + *result = func_name(x); \ + } \ + __global__ void func_name##_kernel_v6(Dummy* result, hipDoubleComplex x) { \ + *result = func_name(x); \ + } \ + void func_name##_v1(double* result, hipDoubleComplex* x) { *result = func_name(x); } \ + void func_name##_v2(double* result, hipFloatComplex x) { *result = func_name(x); } \ + void func_name##_v3(double* result, double x) { *result = func_name(x); } \ + void func_name##_v4(double* result, Dummy x) { *result = func_name(x); } \ + void func_name##_v5(hipDoubleComplex* result, hipDoubleComplex x) { *result = func_name(x); } \ + void func_name##_v6(Dummy* result, hipDoubleComplex x) { *result = func_name(x); } + +__global__ void hipConj_kernel_v1(hipDoubleComplex* result, hipDoubleComplex* x) { + *result = hipConj(x); +} +__global__ void hipConj_kernel_v2(hipDoubleComplex* result, hipFloatComplex x) { + *result = hipConj(x); +} +__global__ void hipConj_kernel_v3(hipDoubleComplex* result, double x) { *result = hipConj(x); } +__global__ void hipConj_kernel_v4(hipDoubleComplex* result, Dummy x) { *result = hipConj(x); } +__global__ void hipConj_kernel_v5(double* result, hipDoubleComplex x) { *result = hipConj(x); } +__global__ void hipConj_kernel_v6(hipFloatComplex* result, hipDoubleComplex x) { + *result = hipConj(x); +} +__global__ void hipConj_kernel_v7(Dummy* result, hipDoubleComplex x) { *result = hipConj(x); } +void hipConj_v1(hipDoubleComplex* result, hipDoubleComplex* x) { *result = hipConj(x); } +void hipConj_v2(hipDoubleComplex* result, hipFloatComplex x) { *result = hipConj(x); } +void hipConj_v3(hipDoubleComplex* result, double x) { *result = hipConj(x); } +void hipConj_v4(hipDoubleComplex* result, Dummy x) { *result = hipConj(x); } +void hipConj_v5(double* result, hipDoubleComplex x) { *result = hipConj(x); } +void hipConj_v6(hipFloatComplex* result, hipDoubleComplex x) { *result = hipConj(x); } +void hipConj_v7(Dummy* result, hipDoubleComplex x) { *result = hipConj(x); } + +NEGATIVE_SHELL_ONE_ARG_DOUBLE(hipCreal) +NEGATIVE_SHELL_ONE_ARG_DOUBLE(hipCimag) +NEGATIVE_SHELL_ONE_ARG_DOUBLE(hipCabs) +NEGATIVE_SHELL_ONE_ARG_DOUBLE(hipCsqabs) diff --git a/projects/hip-tests/catch/unit/complex/complex_negative_kernels_1Arg_float.cc b/projects/hip-tests/catch/unit/complex/complex_negative_kernels_1Arg_float.cc new file mode 100644 index 0000000000..381ea35c0b --- /dev/null +++ b/projects/hip-tests/catch/unit/complex/complex_negative_kernels_1Arg_float.cc @@ -0,0 +1,75 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include + +class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} +}; + +#define NEGATIVE_SHELL_ONE_ARG_FLOAT(func_name) \ + __global__ void func_name##_kernel_v1(float* result, hipFloatComplex* x) { \ + *result = func_name(x); \ + } \ + __global__ void func_name##_kernel_v2(float* result, hipDoubleComplex x) { \ + *result = func_name(x); \ + } \ + __global__ void func_name##_kernel_v3(float* result, float x) { *result = func_name(x); } \ + __global__ void func_name##_kernel_v4(float* result, Dummy x) { *result = func_name(x); } \ + __global__ void func_name##_kernel_v5(hipFloatComplex* result, hipFloatComplex x) { \ + *result = func_name(x); \ + } \ + __global__ void func_name##_kernel_v6(Dummy* result, hipFloatComplex x) { \ + *result = func_name(x); \ + } \ + void func_name##_v1(float* result, hipFloatComplex* x) { *result = func_name(x); } \ + void func_name##_v2(float* result, hipDoubleComplex x) { *result = func_name(x); } \ + void func_name##_v3(float* result, float x) { *result = func_name(x); } \ + void func_name##_v4(float* result, Dummy x) { *result = func_name(x); } \ + void func_name##_v5(hipFloatComplex* result, hipFloatComplex x) { *result = func_name(x); } \ + void func_name##_v6(Dummy* result, hipFloatComplex x) { *result = func_name(x); } + +__global__ void hipConjf_kernel_v1(hipFloatComplex* result, hipFloatComplex* x) { + *result = hipConjf(x); +} +__global__ void hipConjf_kernel_v2(hipFloatComplex* result, hipDoubleComplex x) { + *result = hipConjf(x); +} +__global__ void hipConjf_kernel_v3(hipFloatComplex* result, float x) { *result = hipConjf(x); } +__global__ void hipConjf_kernel_v4(hipFloatComplex* result, Dummy x) { *result = hipConjf(x); } +__global__ void hipConjf_kernel_v5(float* result, hipFloatComplex x) { *result = hipConjf(x); } +__global__ void hipConjf_kernel_v6(hipDoubleComplex* result, hipFloatComplex x) { + *result = hipConjf(x); +} +__global__ void hipConjf_kernel_v7(Dummy* result, hipFloatComplex x) { *result = hipConjf(x); } +void hipConjf_v1(hipFloatComplex* result, hipFloatComplex* x) { *result = hipConjf(x); } +void hipConjf_v2(hipFloatComplex* result, hipDoubleComplex x) { *result = hipConjf(x); } +void hipConjf_v3(hipFloatComplex* result, float x) { *result = hipConjf(x); } +void hipConjf_v4(hipFloatComplex* result, Dummy x) { *result = hipConjf(x); } +void hipConjf_v5(float* result, hipFloatComplex x) { *result = hipConjf(x); } +void hipConjf_v6(hipDoubleComplex* result, hipFloatComplex x) { *result = hipConjf(x); } +void hipConjf_v7(Dummy* result, hipFloatComplex x) { *result = hipConjf(x); } + +NEGATIVE_SHELL_ONE_ARG_FLOAT(hipCrealf) +NEGATIVE_SHELL_ONE_ARG_FLOAT(hipCimagf) +NEGATIVE_SHELL_ONE_ARG_FLOAT(hipCabsf) +NEGATIVE_SHELL_ONE_ARG_FLOAT(hipCsqabsf) diff --git a/projects/hip-tests/catch/unit/complex/complex_negative_kernels_1Arg_rtc.hh b/projects/hip-tests/catch/unit/complex/complex_negative_kernels_1Arg_rtc.hh new file mode 100644 index 0000000000..ad631a3701 --- /dev/null +++ b/projects/hip-tests/catch/unit/complex/complex_negative_kernels_1Arg_rtc.hh @@ -0,0 +1,207 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +static constexpr auto kComplexConj{R"( + __global__ void hipConjf_kernel_v1(hipFloatComplex* result, hipFloatComplex* x) { + *result = hipConjf(x); + } + __global__ void hipConjf_kernel_v2(hipFloatComplex* result, hipDoubleComplex x) { + *result = hipConjf(x); + } + __global__ void hipConjf_kernel_v3(hipFloatComplex* result, float x) { *result = hipConjf(x); } + __global__ void hipConjf_kernel_v4(hipFloatComplex* result, Dummy x) { *result = hipConjf(x); } + __global__ void hipConjf_kernel_v5(float* result, hipFloatComplex x) { *result = hipConjf(x); } + __global__ void hipConjf_kernel_v6(hipDoubleComplex* result, hipFloatComplex x) { + *result = hipConjf(x); + } + __global__ void hipConjf_kernel_v7(Dummy* result, hipFloatComplex x) { *result = hipConjf(x); } + __global__ void hipConj_kernel_v1(hipDoubleComplex* result, hipDoubleComplex* x) { + *result = hipConj(x); + } + __global__ void hipConj_kernel_v2(hipDoubleComplex* result, hipFloatComplex x) { + *result = hipConj(x); + } + __global__ void hipConj_kernel_v3(hipDoubleComplex* result, double x) { *result = hipConj(x); } + __global__ void hipConj_kernel_v4(hipDoubleComplex* result, Dummy x) { *result = hipConj(x); } + __global__ void hipConj_kernel_v5(double* result, hipDoubleComplex x) { *result = hipConj(x); } + __global__ void hipConj_kernel_v6(hipFloatComplex* result, hipDoubleComplex x) { + *result = hipConj(x); + } + __global__ void hipConj_kernel_v7(Dummy* result, hipDoubleComplex x) { *result = hipConj(x); } + void hipConjf_v1(hipFloatComplex* result, hipFloatComplex* x) { *result = hipConjf(x); } + void hipConjf_v2(hipFloatComplex* result, hipDoubleComplex x) { *result = hipConjf(x); } + void hipConjf_v3(hipFloatComplex* result, float x) { *result = hipConjf(x); } + void hipConjf_v4(hipFloatComplex* result, Dummy x) { *result = hipConjf(x); } + void hipConjf_v5(float* result, hipFloatComplex x) { *result = hipConjf(x); } + void hipConjf_v6(hipDoubleComplex* result, hipFloatComplex x) { *result = hipConjf(x); } + void hipConjf_v7(Dummy* result, hipFloatComplex x) { *result = hipConjf(x); } + void hipConj_v1(hipDoubleComplex* result, hipDoubleComplex* x) { *result = hipConj(x); } + void hipConj_v2(hipDoubleComplex* result, hipFloatComplex x) { *result = hipConj(x); } + void hipConj_v3(hipDoubleComplex* result, double x) { *result = hipConj(x); } + void hipConj_v4(hipDoubleComplex* result, Dummy x) { *result = hipConj(x); } + void hipConj_v5(double* result, hipDoubleComplex x) { *result = hipConj(x); } + void hipConj_v6(hipFloatComplex* result, hipDoubleComplex x) { *result = hipConj(x); } + void hipConj_v7(Dummy* result, hipDoubleComplex x) { *result = hipConj(x); } +)"}; + +static constexpr auto kComplexReal{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void hipCrealf_kernel_v1(float* result, hipFloatComplex* x) { *result = hipCrealf(x); } + __global__ void hipCrealf_kernel_v2(float* result, hipDoubleComplex x) { *result = hipCrealf(x); } + __global__ void hipCrealf_kernel_v3(float* result, float x) { *result = hipCrealf(x); } + __global__ void hipCrealf_kernel_v4(float* result, Dummy x) { *result = hipCrealf(x); } + __global__ void hipCrealf_kernel_v5(hipFloatComplex* result, hipFloatComplex x) { + *result = hipCrealf(x); + } + __global__ void hipCrealf_kernel_v6(Dummy* result, hipFloatComplex x) { *result = hipCrealf(x); } + __global__ void hipCreal_kernel_v1(double* result, hipDoubleComplex* x) { *result = hipCreal(x); } + __global__ void hipCreal_kernel_v2(double* result, hipFloatComplex x) { *result = hipCreal(x); } + __global__ void hipCreal_kernel_v3(double* result, double x) { *result = hipCreal(x); } + __global__ void hipCreal_kernel_v4(double* result, Dummy x) { *result = hipCreal(x); } + __global__ void hipCreal_kernel_v5(hipDoubleComplex* result, hipDoubleComplex x) { + *result = hipCreal(x); + } + __global__ void hipCreal_kernel_v6(Dummy* result, hipDoubleComplex x) { *result = hipCreal(x); } + void hipCrealf_v1(float* result, hipFloatComplex* x) { *result = hipCrealf(x); } + void hipCrealf_v2(float* result, hipDoubleComplex x) { *result = hipCrealf(x); } + void hipCrealf_v3(float* result, float x) { *result = hipCrealf(x); } + void hipCrealf_v4(float* result, Dummy x) { *result = hipCrealf(x); } + void hipCrealf_v5(hipFloatComplex* result, hipFloatComplex x) { *result = hipCrealf(x); } + void hipCrealf_v6(Dummy* result, hipFloatComplex x) { *result = hipCrealf(x); } + void hipCreal_v1(double* result, hipDoubleComplex* x) { *result = hipCreal(x); } + void hipCreal_v2(double* result, hipFloatComplex x) { *result = hipCreal(x); } + void hipCreal_v3(double* result, double x) { *result = hipCreal(x); } + void hipCreal_v4(double* result, Dummy x) { *result = hipCreal(x); } + void hipCreal_v5(hipDoubleComplex* result, hipDoubleComplex x) { *result = hipCreal(x); } + void hipCreal_v6(Dummy* result, hipDoubleComplex x) { *result = hipCreal(x); } +)"}; + +static constexpr auto kComplexImag{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void hipCimagf_kernel_v1(float* result, hipFloatComplex* x) { *result = hipCimagf(x); } + __global__ void hipCimagf_kernel_v2(float* result, hipDoubleComplex x) { *result = hipCimagf(x); } + __global__ void hipCimagf_kernel_v3(float* result, float x) { *result = hipCimagf(x); } + __global__ void hipCimagf_kernel_v4(float* result, Dummy x) { *result = hipCimagf(x); } + __global__ void hipCimagf_kernel_v5(hipFloatComplex* result, hipFloatComplex x) { + *result = hipCimagf(x); + } + __global__ void hipCimagf_kernel_v6(Dummy* result, hipFloatComplex x) { *result = hipCimagf(x); } + __global__ void hipCimag_kernel_v1(double* result, hipDoubleComplex* x) { *result = hipCimag(x); } + __global__ void hipCimag_kernel_v2(double* result, hipFloatComplex x) { *result = hipCimag(x); } + __global__ void hipCimag_kernel_v3(double* result, double x) { *result = hipCimag(x); } + __global__ void hipCimag_kernel_v4(double* result, Dummy x) { *result = hipCimag(x); } + __global__ void hipCimag_kernel_v5(hipDoubleComplex* result, hipDoubleComplex x) { + *result = hipCimag(x); + } + __global__ void hipCimag_kernel_v6(Dummy* result, hipDoubleComplex x) { *result = hipCimag(x); } + void hipCimagf_v1(float* result, hipFloatComplex* x) { *result = hipCimagf(x); } + void hipCimagf_v2(float* result, hipDoubleComplex x) { *result = hipCimagf(x); } + void hipCimagf_v3(float* result, float x) { *result = hipCimagf(x); } + void hipCimagf_v4(float* result, Dummy x) { *result = hipCimagf(x); } + void hipCimagf_v5(hipFloatComplex* result, hipFloatComplex x) { *result = hipCimagf(x); } + void hipCimagf_v6(Dummy* result, hipFloatComplex x) { *result = hipCimagf(x); } + void hipCimag_v1(double* result, hipDoubleComplex* x) { *result = hipCimag(x); } + void hipCimag_v2(double* result, hipFloatComplex x) { *result = hipCimag(x); } + void hipCimag_v3(double* result, double x) { *result = hipCimag(x); } + void hipCimag_v4(double* result, Dummy x) { *result = hipCimag(x); } + void hipCimag_v5(hipDoubleComplex* result, hipDoubleComplex x) { *result = hipCimag(x); } + void hipCimag_v6(Dummy* result, hipDoubleComplex x) { *result = hipCimag(x); } +)"}; + +static constexpr auto kComplexAbs{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void hipCabsf_kernel_v1(float* result, hipFloatComplex* x) { *result = hipCabsf(x); } + __global__ void hipCabsf_kernel_v2(float* result, hipDoubleComplex x) { *result = hipCabsf(x); } + __global__ void hipCabsf_kernel_v3(float* result, float x) { *result = hipCabsf(x); } + __global__ void hipCabsf_kernel_v4(float* result, Dummy x) { *result = hipCabsf(x); } + __global__ void hipCabsf_kernel_v5(hipFloatComplex* result, hipFloatComplex x) { + *result = hipCabsf(x); + } + __global__ void hipCabsf_kernel_v6(Dummy* result, hipFloatComplex x) { *result = hipCabsf(x); } + __global__ void hipCabs_kernel_v1(double* result, hipDoubleComplex* x) { *result = hipCabs(x); } + __global__ void hipCabs_kernel_v2(double* result, hipFloatComplex x) { *result = hipCabs(x); } + __global__ void hipCabs_kernel_v3(double* result, double x) { *result = v(x); } + __global__ void hipCabs_kernel_v4(double* result, Dummy x) { *result = hipCabs(x); } + __global__ void hipCabs_kernel_v5(hipDoubleComplex* result, hipDoubleComplex x) { + *result = hipCabs(x); + } + __global__ void hipCabs_kernel_v6(Dummy* result, hipDoubleComplex x) { *result = hipCabs(x); } + void hipCabsf_v1(float* result, hipFloatComplex* x) { *result = hipCabsf(x); } + void hipCabsf_v2(float* result, hipDoubleComplex x) { *result = hipCabsf(x); } + void hipCabsf_v3(float* result, float x) { *result = hipCabsf(x); } + void hipCabsf_v4(float* result, Dummy x) { *result = hipCabsf(x); } + void hipCabsf_v5(hipFloatComplex* result, hipFloatComplex x) { *result = hipCabsf(x); } + void hipCabsf_v6(Dummy* result, hipFloatComplex x) { *result = hipCabsf(x); } + void hipCabs_v1(double* result, hipDoubleComplex* x) { *result = hipCabs(x); } + void hipCabs_v2(double* result, hipFloatComplex x) { *result = hipCabs(x); } + void hipCabs_v3(double* result, double x) { *result = hipCabs(x); } + void hipCabs_v4(double* result, Dummy x) { *result = hipCabs(x); } + void hipCabs_v5(hipDoubleComplex* result, hipDoubleComplex x) { *result = hipCabs(x); } + void hipCabs_v6(Dummy* result, hipDoubleComplex x) { *result = hipCabs(x); } +)"}; + +static constexpr auto kComplexSqabs{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void hipCsqabsf_kernel_v1(float* result, hipFloatComplex* x) { *result = hipCsqabsf(x); } + __global__ void hipCsqabsf_kernel_v2(float* result, hipDoubleComplex x) { *result = hipCsqabsf(x); } + __global__ void hipCsqabsf_kernel_v3(float* result, float x) { *result = hipCsqabsf(x); } + __global__ void hipCsqabsf_kernel_v4(float* result, Dummy x) { *result = hipCsqabsf(x); } + __global__ void hipCsqabsf_kernel_v5(hipFloatComplex* result, hipFloatComplex x) { + *result = hipCsqabsf(x); + } + __global__ void hipCsqabsf_kernel_v6(Dummy* result, hipFloatComplex x) { *result = hipCsqabs(x); } + __global__ void hipCsqabs_kernel_v1(double* result, hipDoubleComplex* x) { *result = hipCsqabs(x); } + __global__ void hipCsqabs_kernel_v2(double* result, hipFloatComplex x) { *result = hipCsqabs(x); } + __global__ void hipCsqabs_kernel_v3(double* result, double x) { *result = hipCsqabs(x); } + __global__ void hipCsqabs_kernel_v4(double* result, Dummy x) { *result = hipCsqabs(x); } + __global__ void hipCsqabs_kernel_v5(hipDoubleComplex* result, hipDoubleComplex x) { + *result = hipCsqabs(x); + } + __global__ void hipCsqabs_kernel_v6(Dummy* result, hipDoubleComplex x) { *result = hipCsqabs(x); } + void hipCsqabsf_v1(float* result, hipFloatComplex* x) { *result = hipCsqabsf(x); } + void hipCsqabsf_v2(float* result, hipDoubleComplex x) { *result = hipCsqabsf(x); } + void hipCsqabsf_v3(float* result, float x) { *result = hipCsqabsf(x); } + void hipCsqabsf_v4(float* result, Dummy x) { *result = hipCsqabsf(x); } + void hipCsqabsf_v5(hipFloatComplex* result, hipFloatComplex x) { *result = hipCsqabsf(x); } + void hipCsqabsf_v6(Dummy* result, hipFloatComplex x) { *result = hipCsqabsf(x); } + void hipCsqabs_v1(double* result, hipDoubleComplex* x) { *result = hipCsqabs(x); } + void hipCsqabs_v2(double* result, hipFloatComplex x) { *result = hipCsqabs(x); } + void hipCsqabs_v3(double* result, double x) { *result = hipCsqabs(x); } + void hipCsqabs_v4(double* result, Dummy x) { *result = hipCsqabs(x); } + void hipCsqabs_v5(hipDoubleComplex* result, hipDoubleComplex x) { *result = hipCsqabs(x); } + void hipCsqabs_v6(Dummy* result, hipDoubleComplex x) { *result = hipCsqabs(x); } +)"}; diff --git a/projects/hip-tests/catch/unit/complex/complex_negative_kernels_2Arg_double.cc b/projects/hip-tests/catch/unit/complex/complex_negative_kernels_2Arg_double.cc new file mode 100644 index 0000000000..c3e29c812c --- /dev/null +++ b/projects/hip-tests/catch/unit/complex/complex_negative_kernels_2Arg_double.cc @@ -0,0 +1,105 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include + +class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} +}; + +#define NEGATIVE_SHELL_TWO_ARG_DOUBLE(func_name) \ + __global__ void func_name##_kernel_v1(hipDoubleComplex* result, hipDoubleComplex* x, \ + hipDoubleComplex y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v2(hipDoubleComplex* result, hipDoubleComplex x, \ + hipDoubleComplex* y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v3(hipDoubleComplex* result, double x, hipDoubleComplex y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v4(hipDoubleComplex* result, hipDoubleComplex x, double y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v5(hipDoubleComplex* result, hipFloatComplex x, \ + hipDoubleComplex y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v6(hipDoubleComplex* result, hipDoubleComplex x, \ + hipFloatComplex y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v7(hipDoubleComplex* result, Dummy x, hipDoubleComplex y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v8(hipDoubleComplex* result, hipDoubleComplex x, Dummy y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v9(double* result, hipDoubleComplex x, hipDoubleComplex y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v10(hipFloatComplex* result, hipDoubleComplex x, \ + hipDoubleComplex y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v11(Dummy* result, hipDoubleComplex x, hipDoubleComplex y) { \ + *result = func_name(x, y); \ + } \ + void func_name##_v1(hipDoubleComplex* result, hipDoubleComplex* x, hipDoubleComplex y) { \ + *result = func_name(x, y); \ + } \ + void func_name##_v2(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex* y) { \ + *result = func_name(x, y); \ + } \ + void func_name##_v3(hipDoubleComplex* result, double x, hipDoubleComplex y) { \ + *result = func_name(x, y); \ + } \ + void func_name##_v4(hipDoubleComplex* result, hipDoubleComplex x, double y) { \ + *result = func_name(x, y); \ + } \ + void func_name##_v5(hipDoubleComplex* result, hipFloatComplex x, hipDoubleComplex y) { \ + *result = func_name(x, y); \ + } \ + void func_name##_v6(hipDoubleComplex* result, hipDoubleComplex x, hipFloatComplex y) { \ + *result = func_name(x, y); \ + } \ + void func_name##_v7(hipDoubleComplex* result, Dummy x, hipDoubleComplex y) { \ + *result = func_name(x, y); \ + } \ + void func_name##_v8(hipDoubleComplex* result, hipDoubleComplex x, Dummy y) { \ + *result = func_name(x, y); \ + } \ + void func_name##_v9(double* result, hipDoubleComplex x, hipDoubleComplex y) { \ + *result = func_name(x, y); \ + } \ + void func_name##_v10(hipFloatComplex* result, hipDoubleComplex x, hipDoubleComplex y) { \ + *result = func_name(x, y); \ + } \ + void func_name##_v11(Dummy* result, hipDoubleComplex x, hipDoubleComplex y) { \ + *result = func_name(x, y); \ + } + +NEGATIVE_SHELL_TWO_ARG_DOUBLE(hipCadd) +NEGATIVE_SHELL_TWO_ARG_DOUBLE(hipCsub) +NEGATIVE_SHELL_TWO_ARG_DOUBLE(hipCmul) +NEGATIVE_SHELL_TWO_ARG_DOUBLE(hipCdiv) \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/complex/complex_negative_kernels_2Arg_float.cc b/projects/hip-tests/catch/unit/complex/complex_negative_kernels_2Arg_float.cc new file mode 100644 index 0000000000..5979155aef --- /dev/null +++ b/projects/hip-tests/catch/unit/complex/complex_negative_kernels_2Arg_float.cc @@ -0,0 +1,105 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include + +class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} +}; + +#define NEGATIVE_SHELL_TWO_ARG_FLOAT(func_name) \ + __global__ void func_name##_kernel_v1(hipFloatComplex* result, hipFloatComplex* x, \ + hipFloatComplex y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v2(hipFloatComplex* result, hipFloatComplex x, \ + hipFloatComplex* y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v3(hipFloatComplex* result, float x, hipFloatComplex y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v4(hipFloatComplex* result, hipFloatComplex x, float y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v5(hipFloatComplex* result, hipDoubleComplex x, \ + hipFloatComplex y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v6(hipFloatComplex* result, hipFloatComplex x, \ + hipDoubleComplex y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v7(hipFloatComplex* result, Dummy x, hipFloatComplex y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v8(hipFloatComplex* result, hipFloatComplex x, Dummy y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v9(float* result, hipFloatComplex x, hipFloatComplex y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v10(hipDoubleComplex* result, hipFloatComplex x, \ + hipFloatComplex y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v11(Dummy* result, hipFloatComplex x, hipFloatComplex y) { \ + *result = func_name(x, y); \ + } \ + void func_name##_v1(hipFloatComplex* result, hipFloatComplex* x, hipFloatComplex y) { \ + *result = func_name(x, y); \ + } \ + void func_name##_v2(hipFloatComplex* result, hipFloatComplex x, hipFloatComplex* y) { \ + *result = func_name(x, y); \ + } \ + void func_name##_v3(hipFloatComplex* result, float x, hipFloatComplex y) { \ + *result = func_name(x, y); \ + } \ + void func_name##_v4(hipFloatComplex* result, hipFloatComplex x, float y) { \ + *result = func_name(x, y); \ + } \ + void func_name##_v5(hipFloatComplex* result, hipDoubleComplex x, hipFloatComplex y) { \ + *result = func_name(x, y); \ + } \ + void func_name##_v6(hipFloatComplex* result, hipFloatComplex x, hipDoubleComplex y) { \ + *result = func_name(x, y); \ + } \ + void func_name##_v7(hipFloatComplex* result, Dummy x, hipFloatComplex y) { \ + *result = func_name(x, y); \ + } \ + void func_name##_v8(hipFloatComplex* result, hipFloatComplex x, Dummy y) { \ + *result = func_name(x, y); \ + } \ + void func_name##_v9(float* result, hipFloatComplex x, hipFloatComplex y) { \ + *result = func_name(x, y); \ + } \ + void func_name##_v10(hipDoubleComplex* result, hipFloatComplex x, hipFloatComplex y) { \ + *result = func_name(x, y); \ + } \ + void func_name##_v11(Dummy* result, hipFloatComplex x, hipFloatComplex y) { \ + *result = func_name(x, y); \ + } + +NEGATIVE_SHELL_TWO_ARG_FLOAT(hipCaddf) +NEGATIVE_SHELL_TWO_ARG_FLOAT(hipCsubf) +NEGATIVE_SHELL_TWO_ARG_FLOAT(hipCmulf) +NEGATIVE_SHELL_TWO_ARG_FLOAT(hipCdivf) \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/complex/complex_negative_kernels_2Arg_rtc.hh b/projects/hip-tests/catch/unit/complex/complex_negative_kernels_2Arg_rtc.hh new file mode 100644 index 0000000000..f3576a1ec8 --- /dev/null +++ b/projects/hip-tests/catch/unit/complex/complex_negative_kernels_2Arg_rtc.hh @@ -0,0 +1,620 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +static constexpr auto kComplexAdd{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void hipCaddf_kernel_v1(hipFloatComplex* result, hipFloatComplex* x, + hipFloatComplex y) { + *result = hipCaddf(x, y); + } + __global__ void hipCaddf_kernel_v2(hipFloatComplex* result, hipFloatComplex x, + hipFloatComplex* y) { + *result = hipCaddf(x, y); + } + __global__ void hipCaddf_kernel_v3(hipFloatComplex* result, float x, hipFloatComplex y) { + *result = hipCaddf(x, y); + } + __global__ void hipCaddf_kernel_v4(hipFloatComplex* result, hipFloatComplex x, float y) { + *result = hipCaddf(x, y); + } + __global__ void hipCaddf_kernel_v5(hipFloatComplex* result, hipDoubleComplex x, + hipFloatComplex y) { + *result = hipCaddf(x, y); + } + __global__ void hipCaddf_kernel_v6(hipFloatComplex* result, hipFloatComplex x, + hipDoubleComplex y) { + *result = hipCaddf(x, y); + } + __global__ void hipCaddf_kernel_v7(hipFloatComplex* result, Dummy x, hipFloatComplex y) { + *result = hipCaddf(x, y); + } + __global__ void hipCaddf_kernel_v8(hipFloatComplex* result, hipFloatComplex x, Dummy y) { + *result = hipCaddf(x, y); + } + __global__ void hipCaddf_kernel_v9(float* result, hipFloatComplex x, hipFloatComplex y) { + *result = hipCaddf(x, y); + } + __global__ void hipCaddf_kernel_v10(hipDoubleComplex* result, hipFloatComplex x, + hipFloatComplex y) { + *result = hipCaddf(x, y); + } + __global__ void hipCaddf_kernel_v11(Dummy* result, hipFloatComplex x, hipFloatComplex y) { + *result = hipCaddf(x, y); + } + __global__ void hipCadd_kernel_v1(hipDoubleComplex* result, hipDoubleComplex* x, + hipDoubleComplex y) { + *result = hipCadd(x, y); + } + __global__ void hipCadd_kernel_v2(hipDoubleComplex* result, hipDoubleComplex x, + hipDoubleComplex* y) { + *result = hipCadd(x, y); + } + __global__ void hipCadd_kernel_v3(hipDoubleComplex* result, double x, hipDoubleComplex y) { + *result = hipCadd(x, y); + } + __global__ void hipCadd_kernel_v4(hipDoubleComplex* result, hipDoubleComplex x, double y) { + *result = hipCadd(x, y); + } + __global__ void hipCadd_kernel_v5(hipDoubleComplex* result, hipFloatComplex x, + hipDoubleComplex y) { + *result = hipCadd(x, y); + } + __global__ void hipCadd_kernel_v6(hipDoubleComplex* result, hipDoubleComplex x, + hipFloatComplex y) { + *result = hipCadd(x, y); + } + __global__ void hipCadd_kernel_v7(hipDoubleComplex* result, Dummy x, hipDoubleComplex y) { + *result = hipCadd(x, y); + } + __global__ void hipCadd_kernel_v8(hipDoubleComplex* result, hipDoubleComplex x, Dummy y) { + *result = hipCadd(x, y); + } + __global__ void hipCadd_kernel_v9(double* result, hipDoubleComplex x, hipDoubleComplex y) { + *result = hipCadd(x, y); + } + __global__ void hipCadd_kernel_v10(hipFloatComplex* result, hipDoubleComplex x, + hipDoubleComplex y) { + *result = hipCadd(x, y); + } + __global__ void hipCadd_kernel_v11(Dummy* result, hipDoubleComplex x, hipDoubleComplex y) { + *result = hipCadd(x, y); + } + void hipCaddf_v1(hipFloatComplex* result, hipFloatComplex* x, hipFloatComplex y) { + *result = hipCaddf(x, y); + } + void hipCaddf_v2(hipFloatComplex* result, hipFloatComplex x, hipFloatComplex* y) { + *result = hipCaddf(x, y); + } + void hipCaddf_v3(hipFloatComplex* result, float x, hipFloatComplex y) { + *result = hipCaddf(x, y); + } + void hipCaddf_v4(hipFloatComplex* result, hipFloatComplex x, float y) { + *result = hipCaddf(x, y); + } + void hipCaddf_v5(hipFloatComplex* result, hipDoubleComplex x, hipFloatComplex y) { + *result = hipCaddf(x, y); + } + void hipCaddf_v6(hipFloatComplex* result, hipFloatComplex x, hipDoubleComplex y) { + *result = hipCaddf(x, y); + } + void hipCaddf_v7(hipFloatComplex* result, Dummy x, hipFloatComplex y) { + *result = hipCaddf(x, y); + } + void hipCaddf_v8(hipFloatComplex* result, hipFloatComplex x, Dummy y) { + *result = hipCaddf(x, y); + } + void hipCaddf_v9(float* result, hipFloatComplex x, hipFloatComplex y) { + *result = hipCaddf(x, y); + } + void hipCaddf_v10(hipDoubleComplex* result, hipFloatComplex x, hipFloatComplex y) { + *result = hipCaddf(x, y); + } + void hipCaddf_v11(Dummy* result, hipFloatComplex x, hipFloatComplex y) { + *result = hipCaddf(x, y); + } + void hipCadd_v1(hipDoubleComplex* result, hipDoubleComplex* x, hipDoubleComplex y) { + *result = hipCadd(x, y); + } + void hipCadd_v2(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex* y) { + *result = hipCadd(x, y); + } + void hipCadd_v3(hipDoubleComplex* result, double x, hipDoubleComplex y) { + *result = hipCadd(x, y); + } + void hipCadd_v4(hipDoubleComplex* result, hipDoubleComplex x, double y) { + *result = hipCadd(x, y); + } + void hipCadd_v5(hipDoubleComplex* result, hipFloatComplex x, hipDoubleComplex y) { + *result = hipCadd(x, y); + } + void hipCadd_v6(hipDoubleComplex* result, hipDoubleComplex x, hipFloatComplex y) { + *result = hipCadd(x, y); + } + void hipCadd_v7(hipDoubleComplex* result, Dummy x, hipDoubleComplex y) { + *result = hipCadd(x, y); + } + void hipCadd_v8(hipDoubleComplex* result, hipDoubleComplex x, Dummy y) { + *result = hipCadd(x, y); + } + void hipCadd_v9(double* result, hipDoubleComplex x, hipDoubleComplex y) { + *result = hipCadd(x, y); + } + void hipCadd_v10(hipFloatComplex* result, hipDoubleComplex x, hipDoubleComplex y) { + *result = hipCadd(x, y); + } + void hipCadd_v11(Dummy* result, hipDoubleComplex x, hipDoubleComplex y) { + *result = hipCadd(x, y); + } +)"}; + +static constexpr auto kComplexSub{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void hipCsubf_kernel_v1(hipFloatComplex* result, hipFloatComplex* x, + hipFloatComplex y) { + *result = hipCsubf(x, y); + } + __global__ void hipCsubf_kernel_v2(hipFloatComplex* result, hipFloatComplex x, + hipFloatComplex* y) { + *result = hipCsubf(x, y); + } + __global__ void hipCsubf_kernel_v3(hipFloatComplex* result, float x, hipFloatComplex y) { + *result = hipCsubf(x, y); + } + __global__ void hipCsubf_kernel_v4(hipFloatComplex* result, hipFloatComplex x, float y) { + *result = hipCsubf(x, y); + } + __global__ void hipCsubf_kernel_v5(hipFloatComplex* result, hipDoubleComplex x, + hipFloatComplex y) { + *result = hipCsubf(x, y); + } + __global__ void hipCsubf_kernel_v6(hipFloatComplex* result, hipFloatComplex x, + hipDoubleComplex y) { + *result = hipCsubf(x, y); + } + __global__ void hipCsubf_kernel_v7(hipFloatComplex* result, Dummy x, hipFloatComplex y) { + *result = hipCsubf(x, y); + } + __global__ void hipCsubf_kernel_v8(hipFloatComplex* result, hipFloatComplex x, Dummy y) { + *result = hipCsubf(x, y); + } + __global__ void hipCsubf_kernel_v9(float* result, hipFloatComplex x, hipFloatComplex y) { + *result = hipCsubf(x, y); + } + __global__ void hipCsubf_kernel_v10(hipDoubleComplex* result, hipFloatComplex x, + hipFloatComplex y) { + *result = hipCsubf(x, y); + } + __global__ void hipCsubf_kernel_v11(Dummy* result, hipFloatComplex x, hipFloatComplex y) { + *result = hipCsubf(x, y); + } + __global__ void hipCsub_kernel_v1(hipDoubleComplex* result, hipDoubleComplex* x, + hipDoubleComplex y) { + *result = hipCsub(x, y); + } + __global__ void hipCsub_kernel_v2(hipDoubleComplex* result, hipDoubleComplex x, + hipDoubleComplex* y) { + *result = hipCsub(x, y); + } + __global__ void hipCsub_kernel_v3(hipDoubleComplex* result, double x, hipDoubleComplex y) { + *result = hipCsub(x, y); + } + __global__ void hipCsub_kernel_v4(hipDoubleComplex* result, hipDoubleComplex x, double y) { + *result = hipCsub(x, y); + } + __global__ void hipCsub_kernel_v5(hipDoubleComplex* result, hipFloatComplex x, + hipDoubleComplex y) { + *result = hipCsub(x, y); + } + __global__ void hipCsub_kernel_v6(hipDoubleComplex* result, hipDoubleComplex x, + hipFloatComplex y) { + *result = hipCsub(x, y); + } + __global__ void hipCsub_kernel_v7(hipDoubleComplex* result, Dummy x, hipDoubleComplex y) { + *result = hipCsub(x, y); + } + __global__ void hipCsub_kernel_v8(hipDoubleComplex* result, hipDoubleComplex x, Dummy y) { + *result = hipCsub(x, y); + } + __global__ void hipCsub_kernel_v9(double* result, hipDoubleComplex x, hipDoubleComplex y) { + *result = hipCsub(x, y); + } + __global__ void hipCsub_kernel_v10(hipFloatComplex* result, hipDoubleComplex x, + hipDoubleComplex y) { + *result = hipCsub(x, y); + } + __global__ void hipCsub_kernel_v11(Dummy* result, hipDoubleComplex x, hipDoubleComplex y) { + *result = hipCsub(x, y); + } + void hipCsubf_v1(hipFloatComplex* result, hipFloatComplex* x, hipFloatComplex y) { + *result = hipCsubf(x, y); + } + void hipCsubf_v2(hipFloatComplex* result, hipFloatComplex x, hipFloatComplex* y) { + *result = hipCsubf(x, y); + } + void hipCsubf_v3(hipFloatComplex* result, float x, hipFloatComplex y) { + *result = hipCsubf(x, y); + } + void hipCsubf_v4(hipFloatComplex* result, hipFloatComplex x, float y) { + *result = hipCsubf(x, y); + } + void hipCsubf_v5(hipFloatComplex* result, hipDoubleComplex x, hipFloatComplex y) { + *result = hipCsubf(x, y); + } + void hipCsubf_v6(hipFloatComplex* result, hipFloatComplex x, hipDoubleComplex y) { + *result = hipCsubf(x, y); + } + void hipCsubf_v7(hipFloatComplex* result, Dummy x, hipFloatComplex y) { + *result = hipCsubf(x, y); + } + void hipCsubf_v8(hipFloatComplex* result, hipFloatComplex x, Dummy y) { + *result = hipCsubf(x, y); + } + void hipCaddf_v9(float* result, hipFloatComplex x, hipFloatComplex y) { + *result = hipCaddf(x, y); + } + void hipCsubf_v10(hipDoubleComplex* result, hipFloatComplex x, hipFloatComplex y) { + *result = hipCsubf(x, y); + } + void hipCsubf_v11(Dummy* result, hipFloatComplex x, hipFloatComplex y) { + *result = hipCsubf(x, y); + } + void hipCsub_v1(hipDoubleComplex* result, hipDoubleComplex* x, hipDoubleComplex y) { + *result = hipCsub(x, y); + } + void hipCsub_v2(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex* y) { + *result = hipCsub(x, y); + } + void hipCsub_v3(hipDoubleComplex* result, double x, hipDoubleComplex y) { + *result = hipCsub(x, y); + } + void hipCsub_v4(hipDoubleComplex* result, hipDoubleComplex x, double y) { + *result = hipCsub(x, y); + } + void hipCsub_v5(hipDoubleComplex* result, hipFloatComplex x, hipDoubleComplex y) { + *result = hipCsub(x, y); + } + void hipCsub_v6(hipDoubleComplex* result, hipDoubleComplex x, hipFloatComplex y) { + *result = hipCsub(x, y); + } + void hipCsub_v7(hipDoubleComplex* result, Dummy x, hipDoubleComplex y) { + *result = hipCsub(x, y); + } + void hipCsub_v8(hipDoubleComplex* result, hipDoubleComplex x, Dummy y) { + *result = hipCsub(x, y); + } + void hipCsub_v9(double* result, hipDoubleComplex x, hipDoubleComplex y) { + *result = hipCsub(x, y); + } + void hipCsub_v10(hipFloatComplex* result, hipDoubleComplex x, hipDoubleComplex y) { + *result = hipCsub(x, y); + } + void hipCsub_v11(Dummy* result, hipDoubleComplex x, hipDoubleComplex y) { + *result = hipCsub(x, y); + } +)"}; + +static constexpr auto kComplexMul{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void hipCmulf_kernel_v1(hipFloatComplex* result, hipFloatComplex* x, + hipFloatComplex y) { + *result = hipCmulf(x, y); + } + __global__ void hipCmulf_kernel_v2(hipFloatComplex* result, hipFloatComplex x, + hipFloatComplex* y) { + *result = hipCmulf(x, y); + } + __global__ void hipCmulf_kernel_v3(hipFloatComplex* result, float x, hipFloatComplex y) { + *result = hipCmulf(x, y); + } + __global__ void hipCmulf_kernel_v4(hipFloatComplex* result, hipFloatComplex x, float y) { + *result = hipCmulf(x, y); + } + __global__ void hipCmulf_kernel_v5(hipFloatComplex* result, hipDoubleComplex x, + hipFloatComplex y) { + *result = hipCmulf(x, y); + } + __global__ void hipCmulf_kernel_v6(hipFloatComplex* result, hipFloatComplex x, + hipDoubleComplex y) { + *result = hipCmulf(x, y); + } + __global__ void hipCmulf_kernel_v7(hipFloatComplex* result, Dummy x, hipFloatComplex y) { + *result = hipCmulf(x, y); + } + __global__ void hipCmulf_kernel_v8(hipFloatComplex* result, hipFloatComplex x, Dummy y) { + *result = hipCmulf(x, y); + } + __global__ void hipCmulf_kernel_v9(float* result, hipFloatComplex x, hipFloatComplex y) { + *result = hipCmulf(x, y); + } + __global__ void hipCmulf_kernel_v10(hipDoubleComplex* result, hipFloatComplex x, + hipFloatComplex y) { + *result = hipCmulf(x, y); + } + __global__ void hipCmulf_kernel_v11(Dummy* result, hipFloatComplex x, hipFloatComplex y) { + *result = hipCmulf(x, y); + } + __global__ void hipCmul_kernel_v1(hipDoubleComplex* result, hipDoubleComplex* x, + hipDoubleComplex y) { + *result = hipCmul(x, y); + } + __global__ void hipCmul_kernel_v2(hipDoubleComplex* result, hipDoubleComplex x, + hipDoubleComplex* y) { + *result = hipCmul(x, y); + } + __global__ void hipCmul_kernel_v3(hipDoubleComplex* result, double x, hipDoubleComplex y) { + *result = hipCmul(x, y); + } + __global__ void hipCmul_kernel_v4(hipDoubleComplex* result, hipDoubleComplex x, double y) { + *result = hipCmul(x, y); + } + __global__ void hipCmul_kernel_v5(hipDoubleComplex* result, hipFloatComplex x, + hipDoubleComplex y) { + *result = hipCmul(x, y); + } + __global__ void hipCmul_kernel_v6(hipDoubleComplex* result, hipDoubleComplex x, + hipFloatComplex y) { + *result = hipCmul(x, y); + } + __global__ void hipCmul_kernel_v7(hipDoubleComplex* result, Dummy x, hipDoubleComplex y) { + *result = hipCmul(x, y); + } + __global__ void hipCmul_kernel_v8(hipDoubleComplex* result, hipDoubleComplex x, Dummy y) { + *result = hipCmul(x, y); + } + __global__ void hipCmul_kernel_v9(double* result, hipDoubleComplex x, hipDoubleComplex y) { + *result = hipCmul(x, y); + } + __global__ void hipCmul_kernel_v10(hipFloatComplex* result, hipDoubleComplex x, + hipDoubleComplex y) { + *result = hipCmul(x, y); + } + __global__ void hipCmul_kernel_v11(Dummy* result, hipDoubleComplex x, hipDoubleComplex y) { + *result = hipCmul(x, y); + } + void hipCmulf_v1(hipFloatComplex* result, hipFloatComplex* x, hipFloatComplex y) { + *result = hipCmulf(x, y); + } + void hipCmulf_v2(hipFloatComplex* result, hipFloatComplex x, hipFloatComplex* y) { + *result = hipCmulf(x, y); + } + void hipCmulf_v3(hipFloatComplex* result, float x, hipFloatComplex y) { + *result = hipCmulf(x, y); + } + void hipCmulf_v4(hipFloatComplex* result, hipFloatComplex x, float y) { + *result = hipCmulf(x, y); + } + void hipCmulf_v5(hipFloatComplex* result, hipDoubleComplex x, hipFloatComplex y) { + *result = hipCmulf(x, y); + } + void hipCmulf_v6(hipFloatComplex* result, hipFloatComplex x, hipDoubleComplex y) { + *result = hipCmulf(x, y); + } + void hipCmulf_v7(hipFloatComplex* result, Dummy x, hipFloatComplex y) { + *result = hipCmulf(x, y); + } + void hipCmulf_v8(hipFloatComplex* result, hipFloatComplex x, Dummy y) { + *result = hipCmulf(x, y); + } + void hipCmulf_v9(float* result, hipFloatComplex x, hipFloatComplex y) { + *result = hipCmulf(x, y); + } + void hipCmulf_v10(hipDoubleComplex* result, hipFloatComplex x, hipFloatComplex y) { + *result = hipCmulf(x, y); + } + void hipCmulf_v11(Dummy* result, hipFloatComplex x, hipFloatComplex y) { + *result = hipCmulf(x, y); + } + void hipCmul_v1(hipDoubleComplex* result, hipDoubleComplex* x, hipDoubleComplex y) { + *result = hipCmul(x, y); + } + void hipCmul_v2(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex* y) { + *result = hipCmul(x, y); + } + void hipCmul_v3(hipDoubleComplex* result, double x, hipDoubleComplex y) { + *result = hipCmul(x, y); + } + void hipCmul_v4(hipDoubleComplex* result, hipDoubleComplex x, double y) { + *result = hipCmul(x, y); + } + void hipCmul_v5(hipDoubleComplex* result, hipFloatComplex x, hipDoubleComplex y) { + *result = hipCmul(x, y); + } + void hipCmul_v6(hipDoubleComplex* result, hipDoubleComplex x, hipFloatComplex y) { + *result = hipCmul(x, y); + } + void hipCmul_v7(hipDoubleComplex* result, Dummy x, hipDoubleComplex y) { + *result = hipCmul(x, y); + } + void hipCmul_v8(hipDoubleComplex* result, hipDoubleComplex x, Dummy y) { + *result = hipCmul(x, y); + } + void hipCmul_v9(double* result, hipDoubleComplex x, hipDoubleComplex y) { + *result = hipCmul(x, y); + } + void hipCmul_v10(hipFloatComplex* result, hipDoubleComplex x, hipDoubleComplex y) { + *result = hipCmul(x, y); + } + void hipCmul_v11(Dummy* result, hipDoubleComplex x, hipDoubleComplex y) { + *result = hipCmul(x, y); + } +)"}; + +static constexpr auto kComplexDiv{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void hipCdivf_kernel_v1(hipFloatComplex* result, hipFloatComplex* x, + hipFloatComplex y) { + *result = hipCdivf(x, y); + } + __global__ void hipCdivf_kernel_v2(hipFloatComplex* result, hipFloatComplex x, + hipFloatComplex* y) { + *result = hipCdivf(x, y); + } + __global__ void hipCdivf_kernel_v3(hipFloatComplex* result, float x, hipFloatComplex y) { + *result = hipCdivf(x, y); + } + __global__ void hipCdivf_kernel_v4(hipFloatComplex* result, hipFloatComplex x, float y) { + *result = hipCdivf(x, y); + } + __global__ void hipCdivf_kernel_v5(hipFloatComplex* result, hipDoubleComplex x, + hipFloatComplex y) { + *result = hipCdivf(x, y); + } + __global__ void hipCdivf_kernel_v6(hipFloatComplex* result, hipFloatComplex x, + hipDoubleComplex y) { + *result = hipCdivf(x, y); + } + __global__ void hipCdivf_kernel_v7(hipFloatComplex* result, Dummy x, hipFloatComplex y) { + *result = hipCdivf(x, y); + } + __global__ void hipCdivf_kernel_v8(hipFloatComplex* result, hipFloatComplex x, Dummy y) { + *result = hipCdivf(x, y); + } + __global__ void hipCdivf_kernel_v9(float* result, hipFloatComplex x, hipFloatComplex y) { + *result = hipCdivf(x, y); + } + __global__ void hipCdivf_kernel_v10(hipDoubleComplex* result, hipFloatComplex x, + hipFloatComplex y) { + *result = hipCdivf(x, y); + } + __global__ void hipCdivf_kernel_v11(Dummy* result, hipFloatComplex x, hipFloatComplex y) { + *result = hipCdivf(x, y); + } + __global__ void hipCdiv_kernel_v1(hipDoubleComplex* result, hipDoubleComplex* x, + hipDoubleComplex y) { + *result = hipCdiv(x, y); + } + __global__ void hipCdiv_kernel_v2(hipDoubleComplex* result, hipDoubleComplex x, + hipDoubleComplex* y) { + *result = hipCdiv(x, y); + } + __global__ void hipCdiv_kernel_v3(hipDoubleComplex* result, double x, hipDoubleComplex y) { + *result = hipCdiv(x, y); + } + __global__ void hipCdiv_kernel_v4(hipDoubleComplex* result, hipDoubleComplex x, double y) { + *result = hipCdiv(x, y); + } + __global__ void hipCdiv_kernel_v5(hipDoubleComplex* result, hipFloatComplex x, + hipDoubleComplex y) { + *result = hipCdiv(x, y); + } + __global__ void hipCdiv_kernel_v6(hipDoubleComplex* result, hipDoubleComplex x, + hipFloatComplex y) { + *result = hipCdiv(x, y); + } + __global__ void hipCdiv_kernel_v7(hipDoubleComplex* result, Dummy x, hipDoubleComplex y) { + *result = hipCdiv(x, y); + } + __global__ void hipCdiv_kernel_v8(hipDoubleComplex* result, hipDoubleComplex x, Dummy y) { + *result = hipCdiv(x, y); + } + __global__ void hipCdiv_kernel_v9(double* result, hipDoubleComplex x, hipDoubleComplex y) { + *result = hipCdiv(x, y); + } + __global__ void hipCdiv_kernel_v10(hipFloatComplex* result, hipDoubleComplex x, + hipDoubleComplex y) { + *result = hipCdiv(x, y); + } + __global__ void hipCdiv_kernel_v11(Dummy* result, hipDoubleComplex x, hipDoubleComplex y) { + *result = hipCdiv(x, y); + } + void hipCdivf_v1(hipFloatComplex* result, hipFloatComplex* x, hipFloatComplex y) { + *result = hipCdivf(x, y); + } + void hipCdivf_v2(hipFloatComplex* result, hipFloatComplex x, hipFloatComplex* y) { + *result = hipCdivf(x, y); + } + void hipCdivf_v3(hipFloatComplex* result, float x, hipFloatComplex y) { + *result = hipCdivf(x, y); + } + void hipCdivf_v4(hipFloatComplex* result, hipFloatComplex x, float y) { + *result = hipCdivf(x, y); + } + void hipCdivf_v5(hipFloatComplex* result, hipDoubleComplex x, hipFloatComplex y) { + *result = hipCdivf(x, y); + } + void hipCdivf_v6(hipFloatComplex* result, hipFloatComplex x, hipDoubleComplex y) { + *result = hipCdivf(x, y); + } + void hipCdivf_v7(hipFloatComplex* result, Dummy x, hipFloatComplex y) { + *result = hipCdivf(x, y); + } + void hipCdivf_v8(hipFloatComplex* result, hipFloatComplex x, Dummy y) { + *result = hipCdivf(x, y); + } + void hipCdivf_v9(float* result, hipFloatComplex x, hipFloatComplex y) { + *result = hipCdivf(x, y); + } + void hipCdivf_v10(hipDoubleComplex* result, hipFloatComplex x, hipFloatComplex y) { + *result = hipCdivf(x, y); + } + void hipCdivf_v11(Dummy* result, hipFloatComplex x, hipFloatComplex y) { + *result = hipCdivf(x, y); + } + void hipCdiv_v1(hipDoubleComplex* result, hipDoubleComplex* x, hipDoubleComplex y) { + *result = hipCdiv(x, y); + } + void hipCdiv_v2(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex* y) { + *result = hipCdiv(x, y); + } + void hipCdiv_v3(hipDoubleComplex* result, double x, hipDoubleComplex y) { + *result = hipCdiv(x, y); + } + void hipCdiv_v4(hipDoubleComplex* result, hipDoubleComplex x, double y) { + *result = hipCdiv(x, y); + } + void hipCdiv_v5(hipDoubleComplex* result, hipFloatComplex x, hipDoubleComplex y) { + *result = hipCdiv(x, y); + } + void hipCdiv_v6(hipDoubleComplex* result, hipDoubleComplex x, hipFloatComplex y) { + *result = hipCdiv(x, y); + } + void hipCdiv_v7(hipDoubleComplex* result, Dummy x, hipDoubleComplex y) { + *result = hipCdiv(x, y); + } + void hipCdiv_v8(hipDoubleComplex* result, hipDoubleComplex x, Dummy y) { + *result = hipCdiv(x, y); + } + void hipCdiv_v9(double* result, hipDoubleComplex x, hipDoubleComplex y) { + *result = hipCdiv(x, y); + } + void hipCdiv_v10(hipFloatComplex* result, hipDoubleComplex x, hipDoubleComplex y) { + *result = hipCdiv(x, y); + } + void hipCdiv_v11(Dummy* result, hipDoubleComplex x, hipDoubleComplex y) { + *result = hipCdiv(x, y); + } +)"}; diff --git a/projects/hip-tests/catch/unit/complex/complex_negative_kernels_3Arg.cc b/projects/hip-tests/catch/unit/complex/complex_negative_kernels_3Arg.cc new file mode 100644 index 0000000000..ce26e50af0 --- /dev/null +++ b/projects/hip-tests/catch/unit/complex/complex_negative_kernels_3Arg.cc @@ -0,0 +1,247 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include + +class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} +}; + +__global__ void hipCfmaf_kernel_v1(hipComplex* result, hipFloatComplex* x, hipFloatComplex y, + hipFloatComplex z) { + *result = hipCfmaf(x, y, z); +} +__global__ void hipCfmaf_kernel_v2(hipComplex* result, hipFloatComplex x, hipFloatComplex* y, + hipFloatComplex z) { + *result = hipCfmaf(x, y, z); +} +__global__ void hipCfmaf_kernel_v3(hipComplex* result, hipFloatComplex x, hipFloatComplex y, + hipFloatComplex* z) { + *result = hipCfmaf(x, y, z); +} +__global__ void hipCfmaf_kernel_v4(hipComplex* result, float x, hipFloatComplex y, + hipFloatComplex z) { + *result = hipCfmaf(x, y, z); +} +__global__ void hipCfmaf_kernel_v5(hipComplex* result, hipFloatComplex x, float y, + hipFloatComplex z) { + *result = hipCfmaf(x, y, z); +} +__global__ void hipCfmaf_kernel_v6(hipComplex* result, hipFloatComplex x, hipFloatComplex y, + float z) { + *result = hipCfmaf(x, y, z); +} +__global__ void hipCfmaf_kernel_v7(hipComplex* result, hipDoubleComplex x, hipFloatComplex y, + hipFloatComplex z) { + *result = hipCfmaf(x, y, z); +} +__global__ void hipCfmaf_kernel_v8(hipComplex* result, hipFloatComplex x, hipDoubleComplex y, + hipFloatComplex z) { + *result = hipCfmaf(x, y, z); +} +__global__ void hipCfmaf_kernel_v9(hipComplex* result, hipFloatComplex x, hipFloatComplex y, + hipDoubleComplex z) { + *result = hipCfmaf(x, y, z); +} +__global__ void hipCfmaf_kernel_v10(hipComplex* result, Dummy x, hipFloatComplex y, + hipFloatComplex z) { + *result = hipCfmaf(x, y, z); +} +__global__ void hipCfmaf_kernel_v11(hipComplex* result, hipFloatComplex x, Dummy y, + hipFloatComplex z) { + *result = hipCfmaf(x, y, z); +} +__global__ void hipCfmaf_kernel_v12(hipComplex* result, hipFloatComplex x, hipFloatComplex y, + Dummy z) { + *result = hipCfmaf(x, y, z); +} +__global__ void hipCfmaf_kernel_v13(float* result, hipFloatComplex x, hipFloatComplex y, + hipFloatComplex z) { + *result = hipCfmaf(x, y, z); +} +__global__ void hipCfmaf_kernel_v14(hipDoubleComplex* result, hipFloatComplex x, hipFloatComplex y, + hipFloatComplex z) { + *result = hipCfmaf(x, y, z); +} +__global__ void hipCfmaf_kernel_v15(Dummy* result, hipFloatComplex x, hipFloatComplex y, + hipFloatComplex z) { + *result = hipCfmaf(x, y, z); +} +void hipCfmaf_v1(hipComplex* result, hipFloatComplex* x, hipFloatComplex y, hipFloatComplex z) { + *result = hipCfmaf(x, y, z); +} +void hipCfmaf_v2(hipComplex* result, hipFloatComplex x, hipFloatComplex* y, hipFloatComplex z) { + *result = hipCfmaf(x, y, z); +} +void hipCfmaf_v3(hipComplex* result, hipFloatComplex x, hipFloatComplex y, hipFloatComplex* z) { + *result = hipCfmaf(x, y, z); +} +void hipCfmaf_v4(hipComplex* result, float x, hipFloatComplex y, hipFloatComplex z) { + *result = hipCfmaf(x, y, z); +} +void hipCfmaf_v5(hipComplex* result, hipFloatComplex x, float y, hipFloatComplex z) { + *result = hipCfmaf(x, y, z); +} +void hipCfmaf_v6(hipComplex* result, hipFloatComplex x, hipFloatComplex y, float z) { + *result = hipCfmaf(x, y, z); +} +void hipCfmaf_v7(hipComplex* result, hipDoubleComplex x, hipFloatComplex y, hipFloatComplex z) { + *result = hipCfmaf(x, y, z); +} +void hipCfmaf_v8(hipComplex* result, hipFloatComplex x, hipDoubleComplex y, hipFloatComplex z) { + *result = hipCfmaf(x, y, z); +} +void hipCfmaf_v9(hipComplex* result, hipFloatComplex x, hipFloatComplex y, hipDoubleComplex z) { + *result = hipCfmaf(x, y, z); +} +void hipCfmaf_v10(hipComplex* result, Dummy x, hipFloatComplex y, hipFloatComplex z) { + *result = hipCfmaf(x, y, z); +} +void hipCfmaf_v11(hipComplex* result, hipFloatComplex x, Dummy y, hipFloatComplex z) { + *result = hipCfmaf(x, y, z); +} +void hipCfmaf_v12(hipComplex* result, hipFloatComplex x, hipFloatComplex y, Dummy z) { + *result = hipCfmaf(x, y, z); +} +void hipCfmaf_v13(float* result, hipFloatComplex x, hipFloatComplex y, hipFloatComplex z) { + *result = hipCfmaf(x, y, z); +} +void hipCfmaf_v14(hipDoubleComplex* result, hipFloatComplex x, hipFloatComplex y, + hipFloatComplex z) { + *result = hipCfmaf(x, y, z); +} +void hipCfmaf_v15(Dummy* result, hipFloatComplex x, hipFloatComplex y, hipFloatComplex z) { + *result = hipCfmaf(x, y, z); +} + +__global__ void hipCfma_kernel_v1(hipDoubleComplex* result, hipDoubleComplex* x, hipDoubleComplex y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); +} +__global__ void hipCfma_kernel_v2(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex* y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); +} +__global__ void hipCfma_kernel_v3(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex y, + hipDoubleComplex* z) { + *result = hipCfma(x, y, z); +} +__global__ void hipCfma_kernel_v4(hipDoubleComplex* result, double x, hipDoubleComplex y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); +} +__global__ void hipCfma_kernel_v5(hipDoubleComplex* result, hipDoubleComplex x, double y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); +} +__global__ void hipCfma_kernel_v6(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex y, + double z) { + *result = hipCfma(x, y, z); +} +__global__ void hipCfma_kernel_v7(hipDoubleComplex* result, hipFloatComplex x, hipDoubleComplex y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); +} +__global__ void hipCfma_kernel_v8(hipDoubleComplex* result, hipDoubleComplex x, hipFloatComplex y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); +} +__global__ void hipCfma_kernel_v9(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex y, + hipFloatComplex z) { + *result = hipCfma(x, y, z); +} +__global__ void hipCfma_kernel_v10(hipDoubleComplex* result, Dummy x, hipDoubleComplex y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); +} +__global__ void hipCfma_kernel_v11(hipDoubleComplex* result, hipDoubleComplex x, Dummy y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); +} +__global__ void hipCfma_kernel_v12(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex y, + Dummy z) { + *result = hipCfma(x, y, z); +} +__global__ void hipCfma_kernel_v13(double* result, hipDoubleComplex x, hipDoubleComplex y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); +} +__global__ void hipCfma_kernel_v14(hipFloatComplex* result, hipDoubleComplex x, hipDoubleComplex y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); +} +__global__ void hipCfma_kernel_v15(Dummy* result, hipDoubleComplex x, hipDoubleComplex y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); +} +void hipCfma_v1(hipDoubleComplex* result, hipDoubleComplex* x, hipDoubleComplex y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); +} +void hipCfma_v2(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex* y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); +} +void hipCfma_v3(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex y, + hipDoubleComplex* z) { + *result = hipCfma(x, y, z); +} +void hipCfma_v4(hipDoubleComplex* result, double x, hipDoubleComplex y, hipDoubleComplex z) { + *result = hipCfma(x, y, z); +} +void hipCfma_v5(hipDoubleComplex* result, hipDoubleComplex x, double y, hipDoubleComplex z) { + *result = hipCfma(x, y, z); +} +void hipCfma_v6(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex y, double z) { + *result = hipCfma(x, y, z); +} +void hipCfma_v7(hipDoubleComplex* result, hipFloatComplex x, hipDoubleComplex y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); +} +void hipCfma_v8(hipDoubleComplex* result, hipDoubleComplex x, hipFloatComplex y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); +} +void hipCfma_v9(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex y, + hipFloatComplex z) { + *result = hipCfma(x, y, z); +} +void hipCfma_v10(hipDoubleComplex* result, Dummy x, hipDoubleComplex y, hipDoubleComplex z) { + *result = hipCfma(x, y, z); +} +void hipCfma_v11(hipDoubleComplex* result, hipDoubleComplex x, Dummy y, hipDoubleComplex z) { + *result = hipCfma(x, y, z); +} +void hipCfma_v12(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex y, Dummy z) { + *result = hipCfma(x, y, z); +} +void hipCfma_v13(double* result, hipDoubleComplex x, hipDoubleComplex y, hipDoubleComplex z) { + *result = hipCfma(x, y, z); +} +void hipCfma_v14(hipFloatComplex* result, hipDoubleComplex x, hipDoubleComplex y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); +} +void hipCfma_v15(Dummy* result, hipDoubleComplex x, hipDoubleComplex y, hipDoubleComplex z) { + *result = hipCfma(x, y, z); +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/complex/complex_negative_kernels_3Arg_rtc.hh b/projects/hip-tests/catch/unit/complex/complex_negative_kernels_3Arg_rtc.hh new file mode 100644 index 0000000000..caf12a31c8 --- /dev/null +++ b/projects/hip-tests/catch/unit/complex/complex_negative_kernels_3Arg_rtc.hh @@ -0,0 +1,246 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +static constexpr auto kComplexFma{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void hipCfmaf_kernel_v1(hipComplex* result, hipFloatComplex* x, hipFloatComplex y, + hipFloatComplex z) { + *result = hipCfmaf(x, y, z); + } + __global__ void hipCfmaf_kernel_v2(hipComplex* result, hipFloatComplex x, hipFloatComplex* y, + hipFloatComplex z) { + *result = hipCfmaf(x, y, z); + } + __global__ void hipCfmaf_kernel_v3(hipComplex* result, hipFloatComplex x, hipFloatComplex y, + hipFloatComplex* z) { + *result = hipCfmaf(x, y, z); + } + __global__ void hipCfmaf_kernel_v4(hipComplex* result, float x, hipFloatComplex y, + hipFloatComplex z) { + *result = hipCfmaf(x, y, z); + } + __global__ void hipCfmaf_kernel_v5(hipComplex* result, hipFloatComplex x, float y, + hipFloatComplex z) { + *result = hipCfmaf(x, y, z); + } + __global__ void hipCfmaf_kernel_v6(hipComplex* result, hipFloatComplex x, hipFloatComplex y, + float z) { + *result = hipCfmaf(x, y, z); + } + __global__ void hipCfmaf_kernel_v7(hipComplex* result, hipDoubleComplex x, hipFloatComplex y, + hipFloatComplex z) { + *result = hipCfmaf(x, y, z); + } + __global__ void hipCfmaf_kernel_v8(hipComplex* result, hipFloatComplex x, hipDoubleComplex y, + hipFloatComplex z) { + *result = hipCfmaf(x, y, z); + } + __global__ void hipCfmaf_kernel_v9(hipComplex* result, hipFloatComplex x, hipFloatComplex y, + hipDoubleComplex z) { + *result = hipCfmaf(x, y, z); + } + __global__ void hipCfmaf_kernel_v10(hipComplex* result, Dummy x, hipFloatComplex y, + hipFloatComplex z) { + *result = hipCfmaf(x, y, z); + } + __global__ void hipCfmaf_kernel_v11(hipComplex* result, hipFloatComplex x, Dummy y, + hipFloatComplex z) { + *result = hipCfmaf(x, y, z); + } + __global__ void hipCfmaf_kernel_v12(hipComplex* result, hipFloatComplex x, hipFloatComplex y, + Dummy z) { + *result = hipCfmaf(x, y, z); + } + __global__ void hipCfmaf_kernel_v13(float* result, hipFloatComplex x, hipFloatComplex y, + hipFloatComplex z) { + *result = hipCfmaf(x, y, z); + } + __global__ void hipCfmaf_kernel_v14(hipDoubleComplex* result, hipFloatComplex x, hipFloatComplex y, + hipFloatComplex z) { + *result = hipCfmaf(x, y, z); + } + __global__ void hipCfmaf_kernel_v15(Dummy* result, hipFloatComplex x, hipFloatComplex y, + hipFloatComplex z) { + *result = hipCfmaf(x, y, z); + } + __global__ void hipCfma_kernel_v1(hipDoubleComplex* result, hipDoubleComplex* x, hipDoubleComplex y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); + } + __global__ void hipCfma_kernel_v2(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex* y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); + } + __global__ void hipCfma_kernel_v3(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex y, + hipDoubleComplex* z) { + *result = hipCfma(x, y, z); + } + __global__ void hipCfma_kernel_v4(hipDoubleComplex* result, double x, hipDoubleComplex y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); + } + __global__ void hipCfma_kernel_v5(hipDoubleComplex* result, hipDoubleComplex x, double y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); + } + __global__ void hipCfma_kernel_v6(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex y, + double z) { + *result = hipCfma(x, y, z); + } + __global__ void hipCfma_kernel_v7(hipDoubleComplex* result, hipFloatComplex x, hipDoubleComplex y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); + } + __global__ void hipCfma_kernel_v8(hipDoubleComplex* result, hipDoubleComplex x, hipFloatComplex y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); + } + __global__ void hipCfma_kernel_v9(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex y, + hipFloatComplex z) { + *result = hipCfma(x, y, z); + } + __global__ void hipCfma_kernel_v10(hipDoubleComplex* result, Dummy x, hipDoubleComplex y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); + } + __global__ void hipCfma_kernel_v11(hipDoubleComplex* result, hipDoubleComplex x, Dummy y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); + } + __global__ void hipCfma_kernel_v12(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex y, + Dummy z) { + *result = hipCfma(x, y, z); + } + __global__ void hipCfma_kernel_v13(double* result, hipDoubleComplex x, hipDoubleComplex y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); + } + __global__ void hipCfma_kernel_v14(hipFloatComplex* result, hipDoubleComplex x, hipDoubleComplex y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); + } + __global__ void hipCfma_kernel_v15(Dummy* result, hipDoubleComplex x, hipDoubleComplex y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); + } + void hipCfmaf_v1(hipComplex* result, hipFloatComplex* x, hipFloatComplex y, hipFloatComplex z) { + *result = hipCfmaf(x, y, z); + } + void hipCfmaf_v2(hipComplex* result, hipFloatComplex x, hipFloatComplex* y, hipFloatComplex z) { + *result = hipCfmaf(x, y, z); + } + void hipCfmaf_v3(hipComplex* result, hipFloatComplex x, hipFloatComplex y, hipFloatComplex* z) { + *result = hipCfmaf(x, y, z); + } + void hipCfmaf_v4(hipComplex* result, float x, hipFloatComplex y, hipFloatComplex z) { + *result = hipCfmaf(x, y, z); + } + void hipCfmaf_v5(hipComplex* result, hipFloatComplex x, float y, hipFloatComplex z) { + *result = hipCfmaf(x, y, z); + } + void hipCfmaf_v6(hipComplex* result, hipFloatComplex x, hipFloatComplex y, float z) { + *result = hipCfmaf(x, y, z); + } + void hipCfmaf_v7(hipComplex* result, hipDoubleComplex x, hipFloatComplex y, hipFloatComplex z) { + *result = hipCfmaf(x, y, z); + } + void hipCfmaf_v8(hipComplex* result, hipFloatComplex x, hipDoubleComplex y, hipFloatComplex z) { + *result = hipCfmaf(x, y, z); + } + void hipCfmaf_v9(hipComplex* result, hipFloatComplex x, hipFloatComplex y, hipDoubleComplex z) { + *result = hipCfmaf(x, y, z); + } + void hipCfmaf_v10(hipComplex* result, Dummy x, hipFloatComplex y, hipFloatComplex z) { + *result = hipCfmaf(x, y, z); + } + void hipCfmaf_v11(hipComplex* result, hipFloatComplex x, Dummy y, hipFloatComplex z) { + *result = hipCfmaf(x, y, z); + } + void hipCfmaf_v12(hipComplex* result, hipFloatComplex x, hipFloatComplex y, Dummy z) { + *result = hipCfmaf(x, y, z); + } + void hipCfmaf_v13(float* result, hipFloatComplex x, hipFloatComplex y, hipFloatComplex z) { + *result = hipCfmaf(x, y, z); + } + void hipCfmaf_v14(hipDoubleComplex* result, hipFloatComplex x, hipFloatComplex y, + hipFloatComplex z) { + *result = hipCfmaf(x, y, z); + } + void hipCfmaf_v15(Dummy* result, hipFloatComplex x, hipFloatComplex y, hipFloatComplex z) { + *result = hipCfmaf(x, y, z); + } + void hipCfma_v1(hipDoubleComplex* result, hipDoubleComplex* x, hipDoubleComplex y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); + } + void hipCfma_v2(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex* y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); + } + void hipCfma_v3(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex y, + hipDoubleComplex* z) { + *result = hipCfma(x, y, z); + } + void hipCfma_v4(hipDoubleComplex* result, double x, hipDoubleComplex y, hipDoubleComplex z) { + *result = hipCfma(x, y, z); + } + void hipCfma_v5(hipDoubleComplex* result, hipDoubleComplex x, double y, hipDoubleComplex z) { + *result = hipCfma(x, y, z); + } + void hipCfma_v6(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex y, double z) { + *result = hipCfma(x, y, z); + } + void hipCfma_v7(hipDoubleComplex* result, hipFloatComplex x, hipDoubleComplex y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); + } + void hipCfma_v8(hipDoubleComplex* result, hipDoubleComplex x, hipFloatComplex y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); + } + void hipCfma_v9(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex y, + hipFloatComplex z) { + *result = hipCfma(x, y, z); + } + void hipCfma_v10(hipDoubleComplex* result, Dummy x, hipDoubleComplex y, hipDoubleComplex z) { + *result = hipCfma(x, y, z); + } + void hipCfma_v11(hipDoubleComplex* result, hipDoubleComplex x, Dummy y, hipDoubleComplex z) { + *result = hipCfma(x, y, z); + } + void hipCfma_v12(hipDoubleComplex* result, hipDoubleComplex x, hipDoubleComplex y, Dummy z) { + *result = hipCfma(x, y, z); + } + void hipCfma_v13(double* result, hipDoubleComplex x, hipDoubleComplex y, hipDoubleComplex z) { + *result = hipCfma(x, y, z); + } + void hipCfma_v14(hipFloatComplex* result, hipDoubleComplex x, hipDoubleComplex y, + hipDoubleComplex z) { + *result = hipCfma(x, y, z); + } + void hipCfma_v15(Dummy* result, hipDoubleComplex x, hipDoubleComplex y, hipDoubleComplex z) { + *result = hipCfma(x, y, z); + } +)"};