SWDEV-1 - complex hiprtc negative tests

Change-Id: I22a200b1e762651689659a7a404b5f27f1ab01a7

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",

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",

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}

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);
 }
 
 /**

catch/unit/complex/complex_basic_common.hh

@@ -83,7 +83,13 @@ template <int expected_errors_num> 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{};

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);
   }