From c025def8ac19f3250851373b443ddbf4701cda2b Mon Sep 17 00:00:00 2001 From: Ajay Date: Thu, 5 Sep 2024 17:56:07 -0700 Subject: [PATCH] SWDEV-1 - complex hiprtc negative tests Change-Id: I22a200b1e762651689659a7a404b5f27f1ab01a7 --- catch/hipTestMain/config/config_amd_linux | 5 ---- catch/hipTestMain/config/config_amd_windows | 5 ---- catch/unit/complex/CMakeLists.txt | 1 + catch/unit/complex/complex.cc | 30 +++++++++---------- catch/unit/complex/complex_basic_common.hh | 8 ++++- .../complex_negative_kernels_1Arg_rtc.hh | 6 ++++ 6 files changed, 29 insertions(+), 26 deletions(-) diff --git a/catch/hipTestMain/config/config_amd_linux b/catch/hipTestMain/config/config_amd_linux index 883f0784ac..5da7726c5f 100644 --- a/catch/hipTestMain/config/config_amd_linux +++ b/catch/hipTestMain/config/config_amd_linux @@ -188,11 +188,6 @@ "Unit_hipGraphicsUnregisterResource_Negative_Parameters", "SWDEV-443760: This test fails when device memory is used for kernel args", "=== 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", "Note: Test disabled due to defect - EXSWHTEC-151", "Unit_hipModuleLoad_Negative_Load_From_A_File_That_Is_Not_A_Module", "Note: Following two tests disabled due to defect - EXSWHTEC-153", diff --git a/catch/hipTestMain/config/config_amd_windows b/catch/hipTestMain/config/config_amd_windows index 7f3a42f52b..33d2fb6db8 100644 --- a/catch/hipTestMain/config/config_amd_windows +++ b/catch/hipTestMain/config/config_amd_windows @@ -342,11 +342,6 @@ "Unit_Printf_flags_Sanity_Positive", "Unit_Printf_length_Sanity_Positive", "=== 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", "=== Below tests are failing PSDB ===", "Unit_hipGraphMem_Alloc_Free_NodeGetParams_Functional_3", "Unit_hipGraphAddMemAllocNode_Positive_FreeInGraph", diff --git a/catch/unit/complex/CMakeLists.txt b/catch/unit/complex/CMakeLists.txt index 35223203f9..57d9e7f2dd 100644 --- a/catch/unit/complex/CMakeLists.txt +++ b/catch/unit/complex/CMakeLists.txt @@ -29,6 +29,7 @@ if(HIP_PLATFORM MATCHES "nvidia") elseif(HIP_PLATFORM MATCHES "amd") set(LINKER_LIBS hiprtc) endif() +add_definitions(-DHIP_PATH="${HIP_PATH}") hip_add_exe_to_target(NAME ComplexTest TEST_SRC ${TEST_SRC} diff --git a/catch/unit/complex/complex.cc b/catch/unit/complex/complex.cc index 39d14edfbe..1d8dc0d40a 100644 --- a/catch/unit/complex/complex.cc +++ b/catch/unit/complex/complex.cc @@ -111,11 +111,11 @@ TEMPLATE_TEST_CASE("Unit_Device_Complex_Unary_Host_Sanity_Positive", "", hipFloa * - 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); + ComplexTypeRTCWrapper<42>(kComplexConj); + ComplexTypeRTCWrapper<28>(kComplexReal); + ComplexTypeRTCWrapper<28>(kComplexImag); + ComplexTypeRTCWrapper<28>(kComplexAbs); + ComplexTypeRTCWrapper<29>(kComplexSqabs); } /** @@ -199,10 +199,10 @@ TEMPLATE_TEST_CASE("Unit_Device_Complex_Binary_Host_Sanity_Positive", "", hipFlo * - 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); + ComplexTypeRTCWrapper<66>(kComplexAdd); + ComplexTypeRTCWrapper<66>(kComplexSub); + ComplexTypeRTCWrapper<66>(kComplexMul); + ComplexTypeRTCWrapper<66>(kComplexDiv); } /** @@ -276,7 +276,7 @@ TEMPLATE_TEST_CASE("Unit_Device_Complex_hipCfma_Host_Sanity_Positive", "", hipFl * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_Device_Complex_hipCfma_Negative_Parameters_RTC") { - ComplexTypeRTCWrapper<60>(kComplexFma); + ComplexTypeRTCWrapper<90>(kComplexFma); } /** @@ -398,9 +398,9 @@ TEST_CASE("Unit_Device_make_hipComplex_Host_Positive") { * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_Device_make_Complex_Negative_Parameters_RTC") { - ComplexTypeRTCWrapper<18>(kMakeHipComplex); - ComplexTypeRTCWrapper<18>(kMakeHipFloatComplex); - ComplexTypeRTCWrapper<18>(kMakeHipDoubleComplex); + ComplexTypeRTCWrapper<27>(kMakeHipComplex); + ComplexTypeRTCWrapper<27>(kMakeHipFloatComplex); + ComplexTypeRTCWrapper<27>(kMakeHipDoubleComplex); } /** @@ -474,8 +474,8 @@ TEMPLATE_TEST_CASE("Unit_Device_Complex_Cast_Host_Sanity_Positive", "", hipFloat * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_Device_Complex_Cast_Negative_Parameters_RTC") { - ComplexTypeRTCWrapper<14>(kComplexDoubleToFloat); - ComplexTypeRTCWrapper<14>(kComplexFloatToDouble); + ComplexTypeRTCWrapper<21>(kComplexDoubleToFloat); + ComplexTypeRTCWrapper<21>(kComplexFloatToDouble); } /** diff --git a/catch/unit/complex/complex_basic_common.hh b/catch/unit/complex/complex_basic_common.hh index f0e07d6cc4..d60b9a4f13 100644 --- a/catch/unit/complex/complex_basic_common.hh +++ b/catch/unit/complex/complex_basic_common.hh @@ -83,7 +83,13 @@ template void ComplexTypeRTCWrapper(const char* progra const char* options[] = {args.c_str()}; hiprtcResult result{hiprtcCompileProgram(program, 1, options)}; #else - hiprtcResult result{hiprtcCompileProgram(program, 0, nullptr)}; + // read from CMakeLists.txt file add_definitions + std::string hip_path = HIP_PATH; + // nvcc errors if headers with relative paths are included + std::string opts_pre = "--pre-include="+ hip_path + "/include/hip/nvidia_detail/nvidia_hip_complex.h"; + std::string opts_pos = "--include-path=/usr/local/cuda/include"; + const char* options[] = {"-default-device" , opts_pre.c_str(), opts_pos.c_str()}; + hiprtcResult result{hiprtcCompileProgram(program, 3, options)}; #endif size_t log_size{}; diff --git a/catch/unit/complex/complex_negative_kernels_1Arg_rtc.hh b/catch/unit/complex/complex_negative_kernels_1Arg_rtc.hh index ad631a3701..14b84d6bb2 100644 --- a/catch/unit/complex/complex_negative_kernels_1Arg_rtc.hh +++ b/catch/unit/complex/complex_negative_kernels_1Arg_rtc.hh @@ -20,6 +20,12 @@ THE SOFTWARE. #pragma once static constexpr auto kComplexConj{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void hipConjf_kernel_v1(hipFloatComplex* result, hipFloatComplex* x) { *result = hipConjf(x); }