EXSWHTEC-316 - Implement tests for Complex type functions (#356)
Change-Id: I67d5e1fd02c7e40319f135bc9ae5bbbde85b5ee7
[ROCm/hip-tests commit: 2f783afe8c]
Этот коммит содержится в:
Nives Vukovic
коммит произвёл
Rakesh Roy
Rakesh Roy
родитель
567cd68705
Коммит
5ef065a74c
@@ -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 ===",
|
||||
|
||||
@@ -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"
|
||||
]
|
||||
|
||||
@@ -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"
|
||||
]
|
||||
}
|
||||
|
||||
@@ -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"
|
||||
]
|
||||
}
|
||||
|
||||
@@ -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.
|
||||
* @}
|
||||
*/
|
||||
|
||||
@@ -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)
|
||||
|
||||
|
||||
@@ -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)
|
||||
|
||||
@@ -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<TestType>(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<TestType>(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<TestType>(input1_r, input1_i);
|
||||
TestType input_val2 = MakeComplexType<TestType>(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<TestType>(input1_r, input1_i);
|
||||
TestType input_val2 = MakeComplexType<TestType>(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<TestType>(input1_r, input1_i);
|
||||
TestType input_val2 = MakeComplexType<TestType>(input2_r, input2_i);
|
||||
TestType input_val3 = MakeComplexType<TestType>(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<TestType>(input1_r, input1_i);
|
||||
TestType input_val2 = MakeComplexType<TestType>(input2_r, input2_i);
|
||||
TestType input_val3 = MakeComplexType<TestType>(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<TestType> result_d(LinearAllocs::hipMalloc, sizeof(TestType));
|
||||
LinearAllocGuard<TestType> result_h(LinearAllocs::hipHostMalloc, sizeof(TestType));
|
||||
|
||||
MakeComplexTypeKernel<TestType><<<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<TestType>(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<hipComplex> result_d(LinearAllocs::hipMalloc, sizeof(hipComplex));
|
||||
LinearAllocGuard<hipComplex> 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<TestType>(input_r, input_i);
|
||||
|
||||
LinearAllocGuard<CastType_t<TestType>> result_d{LinearAllocs::hipMalloc,
|
||||
sizeof(CastType_t<TestType>)};
|
||||
LinearAllocGuard<CastType_t<TestType>> result_h{LinearAllocs::hipHostMalloc,
|
||||
sizeof(CastType_t<TestType>)};
|
||||
|
||||
CastComplexTypeKernel<<<1, 1>>>(result_d.ptr(), input);
|
||||
HIP_CHECK(hipMemcpy(result_h.ptr(), result_d.ptr(), sizeof(CastType_t<TestType>),
|
||||
hipMemcpyDeviceToHost));
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
REQUIRE(result_h.ptr()[0].x == static_cast<decltype(CastType_t<TestType>().x)>(input_r));
|
||||
REQUIRE(result_h.ptr()[0].y == static_cast<decltype(CastType_t<TestType>().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<TestType>(input_r, input_i);
|
||||
|
||||
CastType_t<TestType> result = CastComplexType<CastType_t<TestType>>(input);
|
||||
|
||||
REQUIRE(result.x == static_cast<decltype(CastType_t<TestType>().x)>(input_r));
|
||||
REQUIRE(result.y == static_cast<decltype(CastType_t<TestType>().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);
|
||||
}
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <resource_guards.hh>
|
||||
#include <hip/hip_complex.h>
|
||||
|
||||
template <typename T>
|
||||
__host__ __device__ T MakeComplexType(decltype(T().x) input_val1, decltype(T().x) input_val2) {
|
||||
if constexpr (std::is_same_v<T, hipFloatComplex>) {
|
||||
return make_hipFloatComplex(input_val1, input_val2);
|
||||
} else {
|
||||
return make_hipDoubleComplex(input_val1, input_val2);
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
__global__ void MakeComplexTypeKernel(T* const output_val, decltype(T().x) const input_val1,
|
||||
decltype(T().x) const input_val2) {
|
||||
*output_val = MakeComplexType<T>(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 <typename T> struct CastType {};
|
||||
|
||||
template <> struct CastType<hipFloatComplex> {
|
||||
using type = hipDoubleComplex;
|
||||
};
|
||||
|
||||
template <> struct CastType<hipDoubleComplex> {
|
||||
using type = hipFloatComplex;
|
||||
};
|
||||
|
||||
template <typename T> using CastType_t = typename CastType<T>::type;
|
||||
|
||||
template <typename T1, typename T2> __device__ __host__ T1 CastComplexType(T2 const input_val) {
|
||||
if constexpr (std::is_same_v<hipDoubleComplex, T2>) {
|
||||
return hipComplexDoubleToFloat(input_val);
|
||||
} else if constexpr (std::is_same_v<hipFloatComplex, T2>) {
|
||||
return hipComplexFloatToDouble(input_val);
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T1, typename T2>
|
||||
__global__ void CastComplexTypeKernel(T1* const output_val, T2 const input_val) {
|
||||
*output_val = CastComplexType<T1, T2>(input_val);
|
||||
}
|
||||
|
||||
template <typename T> 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 <int expected_errors_num> 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);
|
||||
}
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_complex.h>
|
||||
|
||||
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);
|
||||
}
|
||||
@@ -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);
|
||||
}
|
||||
)"};
|
||||
@@ -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 <typename T>
|
||||
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 <typename T>
|
||||
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 <typename T>
|
||||
__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<hipFloatComplex, T>) {
|
||||
*output_val = hipCaddf(input_val1, input_val2);
|
||||
} else if constexpr (std::is_same_v<hipDoubleComplex, T>) {
|
||||
*output_val = hipCadd(input_val1, input_val2);
|
||||
}
|
||||
} else if (function == ComplexFunction::kSub) {
|
||||
if constexpr (std::is_same_v<hipFloatComplex, T>) {
|
||||
*output_val = hipCsubf(input_val1, input_val2);
|
||||
} else if constexpr (std::is_same_v<hipDoubleComplex, T>) {
|
||||
*output_val = hipCsub(input_val1, input_val2);
|
||||
}
|
||||
} else if (function == ComplexFunction::kMul) {
|
||||
if constexpr (std::is_same_v<hipFloatComplex, T>) {
|
||||
*output_val = hipCmulf(input_val1, input_val2);
|
||||
} else if constexpr (std::is_same_v<hipDoubleComplex, T>) {
|
||||
*output_val = hipCmul(input_val1, input_val2);
|
||||
}
|
||||
} else if (function == ComplexFunction::kDiv) {
|
||||
if constexpr (std::is_same_v<hipFloatComplex, T>) {
|
||||
*output_val = hipCdivf(input_val1, input_val2);
|
||||
} else if constexpr (std::is_same_v<hipDoubleComplex, T>) {
|
||||
*output_val = hipCdiv(input_val1, input_val2);
|
||||
}
|
||||
} else if (function == ComplexFunction::kConj) {
|
||||
if constexpr (std::is_same_v<hipFloatComplex, T>) {
|
||||
*output_val = hipConjf(input_val1);
|
||||
} else if constexpr (std::is_same_v<hipDoubleComplex, T>) {
|
||||
*output_val = hipConj(input_val1);
|
||||
}
|
||||
} else if (function == ComplexFunction::kFma) {
|
||||
if constexpr (std::is_same_v<hipFloatComplex, T>) {
|
||||
*output_val = hipCfmaf(input_val1, input_val2, input_val3);
|
||||
} else if constexpr (std::is_same_v<hipDoubleComplex, T>) {
|
||||
*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 <typename T>
|
||||
__device__ __host__ void PerformScalarResultFunction(ComplexFunction function,
|
||||
decltype(T().x)* output_val, T input_val) {
|
||||
if (function == ComplexFunction::kReal) {
|
||||
if constexpr (std::is_same_v<hipFloatComplex, T>) {
|
||||
*output_val = hipCrealf(input_val);
|
||||
} else if constexpr (std::is_same_v<hipDoubleComplex, T>) {
|
||||
*output_val = hipCreal(input_val);
|
||||
}
|
||||
} else if (function == ComplexFunction::kImag) {
|
||||
if constexpr (std::is_same_v<hipFloatComplex, T>) {
|
||||
*output_val = hipCimagf(input_val);
|
||||
} else if constexpr (std::is_same_v<hipDoubleComplex, T>) {
|
||||
*output_val = hipCimag(input_val);
|
||||
}
|
||||
} else if (function == ComplexFunction::kAbs) {
|
||||
if constexpr (std::is_same_v<hipFloatComplex, T>) {
|
||||
*output_val = hipCabsf(input_val);
|
||||
} else if constexpr (std::is_same_v<hipDoubleComplex, T>) {
|
||||
*output_val = hipCabs(input_val);
|
||||
}
|
||||
} else if (function == ComplexFunction::kSqabs) {
|
||||
if constexpr (std::is_same_v<hipFloatComplex, T>) {
|
||||
*output_val = hipCsqabsf(input_val);
|
||||
} else if constexpr (std::is_same_v<hipDoubleComplex, T>) {
|
||||
*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 <typename T>
|
||||
__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 <typename T>
|
||||
__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 <typename T> void ComplexFunctionUnaryDeviceTest(ComplexFunction function, T input_val) {
|
||||
if (function == ComplexFunction::kConj) {
|
||||
LinearAllocGuard<T> result_d{LinearAllocs::hipMalloc, sizeof(T)};
|
||||
LinearAllocGuard<T> 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<decltype(T().x)> result_d{LinearAllocs::hipMalloc, sizeof(decltype(T().x))};
|
||||
LinearAllocGuard<decltype(T().x)> 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 <typename T> 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 <typename T>
|
||||
void ComplexFunctionBinaryDeviceTest(ComplexFunction function, T input_val1, T input_val2) {
|
||||
LinearAllocGuard<T> result_d{LinearAllocs::hipMalloc, sizeof(T)};
|
||||
LinearAllocGuard<T> 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 <typename T>
|
||||
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 <typename T>
|
||||
void ComplexFunctionTernaryDeviceTest(ComplexFunction function, T input_val1, T input_val2,
|
||||
T input_val3) {
|
||||
LinearAllocGuard<T> result_d{LinearAllocs::hipMalloc, sizeof(T)};
|
||||
LinearAllocGuard<T> 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 <typename T>
|
||||
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);
|
||||
}
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_complex.h>
|
||||
|
||||
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)
|
||||
@@ -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);
|
||||
}
|
||||
)"};
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_complex.h>
|
||||
|
||||
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)
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_complex.h>
|
||||
|
||||
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)
|
||||
@@ -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); }
|
||||
)"};
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_complex.h>
|
||||
|
||||
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)
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_complex.h>
|
||||
|
||||
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)
|
||||
@@ -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);
|
||||
}
|
||||
)"};
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_complex.h>
|
||||
|
||||
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);
|
||||
}
|
||||
@@ -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);
|
||||
}
|
||||
)"};
|
||||
Ссылка в новой задаче
Block a user