SWDEV-385161 - Deprecate usage of env vars in HIP samples and tests (#363)

* SWDEV-385161 - Deprecate usage of env vars in HIP samples and tests env vars such as HIP_PATH, ROCM_PATH, HIP_COMPILER, HIP_RUNTIME, HSA_PATH etc are removed.

Change-Id: I280e20b291cfb14675846aae72d97fa5b3ad567d

catch/CMakeLists.txt

@@ -9,51 +9,45 @@ project(hiptests)
 # flag to generate standalone exe per src file.
 message(STATUS "STANDALONE_TESTS : ${STANDALONE_TESTS}")
 
-# Check if platform and compiler are set
-if(HIP_PLATFORM STREQUAL "amd")
-    if(HIP_COMPILER STREQUAL "nvcc")
-        message(FATAL_ERROR "Unexpected HIP_COMPILER:${HIP_COMPILER} is set for HIP_PLATFOR:amd")
-    endif()
-elseif(HIP_PLATFORM STREQUAL "nvidia")
-    if(DEFINED HIP_COMPILER AND NOT HIP_COMPILER STREQUAL "nvcc")
-        message(FATAL_ERROR "Unexpected HIP_COMPILER: ${HIP_COMPILER} is set for HIP_PLATFORM:nvidia")
-    endif()
-else()
+# Check if platform is set
+if(NOT HIP_PLATFORM STREQUAL "amd" AND NOT HIP_PLATFORM STREQUAL "nvidia")
     message(FATAL_ERROR "Unexpected HIP_PLATFORM: " ${HIP_PLATFORM})
 endif()
 
-if (WIN32)
+if(WIN32)
     set(EXT ".bat")
 endif()
 
-# Read -DROCM_Path and env{ROCM_PATH}
-if(NOT DEFINED ROCM_PATH)
-    if(DEFINED ENV{ROCM_PATH})
-        set(ROCM_PATH $ENV{ROCM_PATH} CACHE STRING "ROCM Path")
+if(HIP_PLATFORM STREQUAL "amd")
+   if(UNIX AND DEFINED ROCM_PATH)
+      # Read -DROCM_PATH and set CXX_FLAGS for amd platform only
+      set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} --rocm-path=${ROCM_PATH}")
+   endif()
+
+   if(DEFINED HIP_PATH)
+      # Read -DHIP_PATH and set CXX_FLAGS for amd platform only
+      set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} --hip-path=${HIP_PATH}")
+   endif()
+endif()
+
+# Read -DHIP_PATH
+# If not set read env{HIP_PATH} only on Windows
+if(WIN32)
+    if(NOT DEFINED HIP_PATH)
+        if(DEFINED ENV{HIP_PATH})
+            set(HIP_PATH $ENV{HIP_PATH} CACHE STRING "HIP Path")
+        endif()
     endif()
 endif()
 
-# Read -DHIP_Path and env{HIP_PATH}
 if(NOT DEFINED HIP_PATH)
-    if(DEFINED ENV{HIP_PATH})
-        set(HIP_PATH $ENV{HIP_PATH} CACHE STRING "HIP Path")
-    endif()
-endif()
-
-# both are not set
-if(NOT DEFINED HIP_PATH AND NOT DEFINED ROCM_PATH)
-    set(HIP_PATH "/opt/rocm")
-    set(ROCM_PATH "/opt/rocm")
-elseif(DEFINED HIP_PATH AND NOT DEFINED ROCM_PATH)
-    execute_process(COMMAND ${HIP_PATH}/bin/hipconfig${EXT} --rocmpath
-                OUTPUT_VARIABLE ROCM_PATH
-                OUTPUT_STRIP_TRAILING_WHITESPACE)
-elseif(DEFINED ROCM_PATH AND NOT DEFINED HIP_PATH)
-    set(HIP_PATH ${ROCM_PATH})
+   if(DEFINED ROCM_PATH)
+      set(HIP_PATH ${ROCM_PATH})
+   else()
+      set(HIP_PATH "/opt/rocm")
+   endif()
 endif()
 message(STATUS "HIP_PATH: ${HIP_PATH}")
-message(STATUS "ROCM_PATH: ${ROCM_PATH}")
-
 
 set(CMAKE_CXX_COMPILER "${HIP_PATH}/bin/hipcc${EXT}")
 set(CMAKE_C_COMPILER "${HIP_PATH}/bin/hipcc${EXT}")
@@ -62,10 +56,6 @@ execute_process(COMMAND ${HIPCONFIG_EXECUTABLE} --version
                 OUTPUT_VARIABLE HIP_VERSION
                 OUTPUT_STRIP_TRAILING_WHITESPACE)
 
-if(HIP_PLATFORM STREQUAL "amd")
-    # prioritize -DROCM_PATH over env{ROCM_PATH} for amd platform only
-    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} --rocm-path=${ROCM_PATH}")
-endif()
 # enforce c++17
 set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} --std=c++17")
 
@@ -155,8 +145,8 @@ if (WIN32)
   SET(CMAKE_CXX_RESPONSE_FILE_LINK_FLAG "")
 endif()
 
-if(HIP_PLATFORM MATCHES "amd" AND HIP_COMPILER MATCHES "clang")
-    add_compile_options(-Wall -Wextra -pedantic -Werror -Wno-deprecated)
+if(HIP_PLATFORM STREQUAL "amd")
+    add_compile_options(-Wall -Wextra -Werror -Wno-deprecated)
 endif()
 
 cmake_policy(PUSH)
@@ -178,10 +168,8 @@ message(STATUS "CMAKE HIP ARCHITECTURES: ${CMAKE_HIP_ARCHITECTURES}")
 # That results in hipcc building the test for gfx803 (the default target)
 # preference to pass arch -
 # OFFLOAD_ARCH_STR
-# ENV{HCC_AMDGPU_TARGET}
 # rocm_agent_enumerator
 if(NOT DEFINED OFFLOAD_ARCH_STR
-   AND NOT DEFINED ENV{HCC_AMDGPU_TARGET}
    AND EXISTS "${ROCM_PATH}/bin/rocm_agent_enumerator"
    AND HIP_PLATFORM STREQUAL "amd" AND UNIX)
     execute_process(COMMAND ${ROCM_PATH}/bin/rocm_agent_enumerator
@@ -208,10 +196,6 @@ endif()
 
 if(DEFINED OFFLOAD_ARCH_STR)
     set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${OFFLOAD_ARCH_STR} ")
-elseif(DEFINED ENV{HCC_AMDGPU_TARGET})
-    # hipcc pl script appends it to the options
-    set(OFFLOAD_ARCH_STR "--offload-arch=$ENV{HCC_AMDGPU_TARGET}")
-    set(HIP_GPU_ARCH_LIST $ENV{HCC_AMDGPU_TARGET})
 endif()
 message(STATUS "Using offload arch string: ${OFFLOAD_ARCH_STR}")
 
@@ -240,7 +224,7 @@ set(_subdirs ${_autogen} "subdirs(..)\n")
 file(WRITE ${CMAKE_CURRENT_BINARY_DIR}/${CATCH_BUILD_DIR}/CTestTestfile.cmake ${_subdirs})
 
 # Enable device lambda on nvidia platforms
-if(HIP_COMPILER MATCHES "nvcc")
+if(HIP_PLATFORM STREQUAL "nvidia")
     set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} --extended-lambda")
 endif()
 

catch/README.md

@@ -11,7 +11,7 @@ Tests in Catch2 are declared via ```TEST_CASE```.
 [Catch2 Detailed Reference](https://github.com/catchorg/Catch2/blob/v2.13.6/docs/Readme.md#top)
 
 ## Taking care of existing features
-- Don’t build on platform: EXCLUDE_(HIP_PLATFORM/HIP_RUNTIME), can be done via CMAKE. Adding source in if(HIP_PLATFORM == amd/nvidia).
+- Don’t build on platform: EXCLUDE_HIP_PLATFORM, can be done via CMAKE. Adding source in if(HIP_PLATFORM == amd/nvidia).
 - HIPCC_OPTIONS/CLANG Options: Can be done via: set_source_files_properties(src.cc PROPERTIES COMPILE_FLAGS “…”).
 - Additional libraries: Can be done via target_link_libraries()
 - Multiple runs with different args: This can be done by Catch’s Feature: GENERATE(…)

catch/include/hip_test_common.hh

@@ -21,6 +21,7 @@ THE SOFTWARE.
 */
 
 #pragma once
+#pragma clang diagnostic ignored "-Wsign-compare"
 #include "hip_test_context.hh"
 
 #include <catch.hpp>
@@ -355,7 +356,7 @@ class BlockingContext {
   hipStream_t stream;
 
  public:
-  BlockingContext(hipStream_t s) : stream(s), blocked(true) {}
+  BlockingContext(hipStream_t s) : blocked(true), stream(s) {}
 
   BlockingContext(const BlockingContext& in) {
     blocked = in.blocked_val();

catch/include/memcpy3d_tests_common.hh

@@ -21,7 +21,8 @@ THE SOFTWARE.
 */
 
 #pragma once
-
+#pragma clang diagnostic ignored "-Wmissing-field-initializers"
+#pragma clang diagnostic ignored "-Wunused-lambda-capture"
 #include <variant>
 
 #include <hip_test_common.hh>
@@ -31,7 +32,7 @@ THE SOFTWARE.
 
 using PtrVariant = std::variant<hipPitchedPtr, hipArray_t>;
 
-static hipMemcpyKind ReverseMemcpyDirection(const hipMemcpyKind direction) {
+static inline hipMemcpyKind ReverseMemcpyDirection(const hipMemcpyKind direction) {
   switch (direction) {
     case hipMemcpyHostToDevice:
       return hipMemcpyDeviceToHost;
@@ -76,7 +77,7 @@ static bool operator==(const hipExtent& lhs, const hipExtent& rhs) {
   return lhs.width == rhs.width && lhs.height == rhs.height && lhs.depth == rhs.depth;
 }
 
-static bool operator==(const hipMemcpy3DParms& lhs, const hipMemcpy3DParms& rhs) {
+static inline bool operator==(const hipMemcpy3DParms& lhs, const hipMemcpy3DParms& rhs) {
   return lhs.dstArray == rhs.dstArray && lhs.dstPtr == rhs.dstPtr && lhs.dstPos == rhs.dstPos &&
       lhs.srcArray == rhs.srcArray && lhs.srcPtr == rhs.srcPtr && lhs.srcPos == rhs.srcPos &&
       lhs.extent == rhs.extent && lhs.kind == rhs.kind;
@@ -168,7 +169,6 @@ void Memcpy3DDeviceToDeviceShell(F memcpy_func, const hipStream_t kernel_stream
   const auto device_count = HipTest::getDeviceCount();
   const auto src_device = GENERATE_COPY(range(0, device_count));
   const auto dst_device = GENERATE_COPY(range(0, device_count));
-  const size_t src_cols_mult = GENERATE(1, 2);
 
   INFO("Src device: " << src_device << ", Dst device: " << dst_device);
 
@@ -797,4 +797,4 @@ void DrvMemcpy3DArrayDeviceShell(F memcpy_func, const hipStream_t kernel_stream
   };
   PitchedMemoryVerify(host_alloc.ptr(), extent.width, extent.width / sizeof(int), extent.height,
                       extent.depth, f);
-}
+}

catch/include/performance_common.hh

@@ -33,6 +33,9 @@ THE SOFTWARE.
 #include <hip_test_common.hh>
 #include <resource_guards.hh>
 
+#pragma clang diagnostic ignored "-Wunused-but-set-variable"
+#pragma clang diagnostic ignored "-Wunused-function"
+
 #if defined(_WIN32)
 #if defined(_WIN64)
 typedef __int64 ssize_t;

catch/multiproc/hipMemCoherencyTstMProc.cc

@@ -131,7 +131,6 @@ TEST_CASE("Unit_malloc_CoherentTst") {
   p = strstr(prop.gcnArchName, "xnack+");
   if (p) {
     // Test Case execution begins from here
-    int stat = 0;
     int managed = 0;
     HIPCHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
                                     0));
@@ -162,7 +161,6 @@ TEST_CASE("Unit_malloc_CoherentTstWthAdvise") {
   char *p = NULL;
   p = strstr(prop.gcnArchName, "xnack+");
   if (p) {
-    int stat = 0;
     int managed = 0;
     HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
                                     0));
@@ -195,7 +193,6 @@ TEST_CASE("Unit_mmap_CoherentTst") {
   char *p = NULL;
   p = strstr(prop.gcnArchName, "xnack+");
   if (p) {
-    int stat = 0;
     int managed = 0;
     HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
                                     0));
@@ -233,7 +230,6 @@ TEST_CASE("Unit_mmap_CoherentTstWthAdvise") {
   char *p = NULL;
   p = strstr(prop.gcnArchName, "xnack+");
   if (p) {
-    int stat = 0;
     int managed = 0;
     HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
                                     0));

catch/unit/cooperativeGrps/coalesced_groups_shfl_down.cc

@@ -31,7 +31,7 @@ THE SOFTWARE.
 #include <hip/hip_cooperative_groups.h>
 #include <stdio.h>
 #include <vector>
-
+#pragma clang diagnostic ignored "-Wunused-variable"
 using namespace cooperative_groups;
 
 #define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs)

catch/unit/cooperativeGrps/coalesced_groups_shfl_up.cc

@@ -30,7 +30,7 @@ THE SOFTWARE.
 #include <hip/hip_cooperative_groups.h>
 #include <stdio.h>
 #include <vector>
-
+#pragma clang diagnostic ignored "-Wunused-variable"
 using namespace cooperative_groups;
 
 #define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs)
@@ -247,4 +247,4 @@ TEST_CASE("Unit_coalesced_groups_shfl_up") {
   }
 }
 
-/* Kogge-Stone algorithm */
+/* Kogge-Stone algorithm */

catch/unit/cooperativeGrps/grid_group.cc

@@ -143,7 +143,7 @@ TEST_CASE("Unit_Grid_Group_Getters_Positive_Basic") {
   HIP_CHECK(hipDeviceSynchronize());
 
   // Verify grid_group.is_valid() values
-  ArrayAllOf(uint_arr.ptr(), grid.thread_count_, [](uint32_t i) { return 1; });
+  ArrayAllOf(uint_arr.ptr(), grid.thread_count_, [](uint32_t) { return 1; });
 }
 
 /**
@@ -282,4 +282,4 @@ TEST_CASE("Unit_Grid_Group_Sync_Positive_Basic") {
     }
     REQUIRE(uint_arr.ptr()[i * grid.block_count_ + j] == max_in_this_loop - 1);
   }
-}
+}

catch/unit/device/getDeviceCount_exe.cc

@@ -34,10 +34,10 @@ bool UNSETENV(std::string var) {
   return (result == 0) ? true: false;
 }
 
-bool SETENV(std::string var, std::string value, int overwrite) {
+bool SETENV(std::string var, std::string value) {
   int result = -1;
   #ifdef __unix__
-    result = setenv(var.c_str(), value.c_str(), overwrite);
+    result = setenv(var.c_str(), value.c_str(), 1);
   #else
     result = _putenv((var + '=' + value).c_str());
   #endif
@@ -59,7 +59,7 @@ int main(int argc, char** argv) {
   // disable visible_devices env from shell
 #ifdef __HIP_PLATFORM_NVCC__
   UNSETENV("CUDA_VISIBLE_DEVICES");
-  SETENV("CUDA_VISIBLE_DEVICES", argv[1], 1);
+  SETENV("CUDA_VISIBLE_DEVICES", argv[1]);
   auto init_res = hipInit(0);
   if (hipSuccess != init_res) {
     std::cerr << "CUDA INIT API returned : " << hipGetErrorString(init_res) << std::endl;
@@ -68,8 +68,8 @@ int main(int argc, char** argv) {
 #else
   UNSETENV("ROCR_VISIBLE_DEVICES");
   UNSETENV("HIP_VISIBLE_DEVICES");
-  SETENV("ROCR_VISIBLE_DEVICES", argv[1], 1);
-  SETENV("HIP_VISIBLE_DEVICES", argv[1], 1);
+  SETENV("ROCR_VISIBLE_DEVICES", argv[1]);
+  SETENV("HIP_VISIBLE_DEVICES", argv[1]);
 #endif
 
   int count = 0;
@@ -86,4 +86,4 @@ int main(int argc, char** argv) {
   UNSETENV("HIP_VISIBLE_DEVICES");
 #endif
   return count;
-}
+}

catch/unit/device/hipDeviceGetP2PAttribute_exe.cc

@@ -32,10 +32,10 @@ bool UNSETENV(std::string var) {
   return (result == 0) ? true: false;
 }
 
-bool SETENV(std::string var, std::string value, int overwrite) {
+bool SETENV(std::string var, std::string value) {
   int result = -1;
   #ifdef __unix__
-    result = setenv(var.c_str(), value.c_str(), overwrite);
+    result = setenv(var.c_str(), value.c_str(), 1);
   #else
     result = _putenv((var + '=' + value).c_str());
   #endif
@@ -44,10 +44,10 @@ bool SETENV(std::string var, std::string value, int overwrite) {
 
 void inline hideDevices(const char* devices) {
 #if HT_NVIDIA
-  SETENV("CUDA_VISIBLE_DEVICES", devices, 1);
+  SETENV("CUDA_VISIBLE_DEVICES", devices);
 #else
-  SETENV("HIP_VISIBLE_DEVICES", devices, 1);
-  SETENV("ROCR_VISIBLE_DEVICES", devices, 1);
+  SETENV("HIP_VISIBLE_DEVICES", devices);
+  SETENV("ROCR_VISIBLE_DEVICES", devices);
 #endif
 }
 

catch/unit/device/hipDeviceGetUuid.cc

@@ -53,7 +53,7 @@ TEST_CASE("Unit_hipDeviceGetUuid_Positive") {
 
   // Atleast one non zero value
   size_t uuidSize = sizeof(uuid.bytes) / sizeof(uuid.bytes[0]);
-  for (int i = 0; i < uuidSize; i++) {
+  for (size_t i = 0; i < uuidSize; i++) {
     if (uuid.bytes[i] != 0) {
       uuidValid = true;
       break;

catch/unit/deviceLib/BuiltIns_fadd.cc

@@ -198,8 +198,8 @@ TEST_CASE("Unit_BuiltInAtomicAdd_CoherentGlobalMemWithRtc") {
       void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f,
                           HIP_LAUNCH_PARAM_BUFFER_SIZE,
                           &size, HIP_LAUNCH_PARAM_END};
-      hipModuleLaunchKernel(fmaxkernel, 1, 1, 1, 1, 1, 1, 0,
-                            nullptr, nullptr, config_d);
+      HIP_CHECK(hipModuleLaunchKernel(fmaxkernel, 1, 1, 1, 1, 1, 1, 0,
+                            nullptr, nullptr, config_d));
       HIP_CHECK(hipDeviceSynchronize());
       HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost));
       REQUIRE(A_h[0] == INITIAL_VAL);
@@ -278,8 +278,8 @@ TEST_CASE("Unit_BuiltInAtomicAdd_NonCoherentGlobalMemWithRtc") {
       void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f,
                           HIP_LAUNCH_PARAM_BUFFER_SIZE,
                           &size, HIP_LAUNCH_PARAM_END};
-      hipModuleLaunchKernel(fmaxkernel, 1, 1, 1, 1, 1, 1, 0,
-                            nullptr, nullptr, config_d);
+      HIP_CHECK(hipModuleLaunchKernel(fmaxkernel, 1, 1, 1, 1, 1, 1, 0,
+                            nullptr, nullptr, config_d));
       HIP_CHECK(hipDeviceSynchronize());
       HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost));
       REQUIRE(A_h[0] == INITIAL_VAL + INC_VAL);

catch/unit/deviceLib/BuiltIns_fmax.cc

@@ -229,8 +229,8 @@ TEST_CASE("Unit_BuiltinAtomicsRTC_fmaxCoherentGlobalMem") {
       void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f,
         HIP_LAUNCH_PARAM_BUFFER_SIZE,
         &size, HIP_LAUNCH_PARAM_END};
-      hipModuleLaunchKernel(fmaxkernel, 1, 1, 1, 1, 1, 1, 0,
-          nullptr, nullptr, config_d);
+      HIP_CHECK(hipModuleLaunchKernel(fmaxkernel, 1, 1, 1, 1, 1, 1, 0,
+                nullptr, nullptr, config_d));
       HIP_CHECK(hipDeviceSynchronize());
       HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost));
       REQUIRE(*B_h == 0);
@@ -324,8 +324,8 @@ TEST_CASE("Unit_BuiltinAtomicsRTC_fmaxNonCoherentGlobalFlatMem") {
       void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f,
         HIP_LAUNCH_PARAM_BUFFER_SIZE,
         &size, HIP_LAUNCH_PARAM_END};
-      hipModuleLaunchKernel(fmaxkernel, 1, 1, 1, 1, 1, 1, 0,
-          nullptr, nullptr, config_d);
+      HIP_CHECK(hipModuleLaunchKernel(fmaxkernel, 1, 1, 1, 1, 1, 1, 0,
+                nullptr, nullptr, config_d));
       HIP_CHECK(hipDeviceSynchronize());
       HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost));
       REQUIRE(*B_h == INITIAL_VAL);

catch/unit/deviceLib/BuiltIns_fmin.cc

@@ -230,8 +230,8 @@ TEST_CASE("Unit_BuiltinAtomicsRTC__fminCoherentGlobalMem") {
       void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f,
         HIP_LAUNCH_PARAM_BUFFER_SIZE,
         &size, HIP_LAUNCH_PARAM_END};
-      hipModuleLaunchKernel(fmaxkernel, 1, 1, 1, 1, 1, 1, 0,
-          nullptr, nullptr, config_d);
+      HIP_CHECK(hipModuleLaunchKernel(fmaxkernel, 1, 1, 1, 1, 1, 1, 0,
+                nullptr, nullptr, config_d));
       HIP_CHECK(hipDeviceSynchronize());
       HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost));
       REQUIRE(*B_h == 0);
@@ -327,8 +327,8 @@ TEST_CASE("Unit_BuiltinAtomicsRTC_fminNonCoherentGlobalFlatMem") {
       void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f,
         HIP_LAUNCH_PARAM_BUFFER_SIZE,
         &size, HIP_LAUNCH_PARAM_END};
-      hipModuleLaunchKernel(fmaxkernel, 1, 1, 1, 1, 1, 1, 0,
-          nullptr, nullptr, config_d);
+      HIP_CHECK(hipModuleLaunchKernel(fmaxkernel, 1, 1, 1, 1, 1, 1, 0,
+                nullptr, nullptr, config_d));
       HIP_CHECK(hipDeviceSynchronize());
       HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost));
       REQUIRE(*B_h == INITIAL_VAL);

catch/unit/deviceLib/hipStdComplex.cc

@@ -22,6 +22,7 @@ THE SOFTWARE.
 #include <hip_test_checkers.hh>
 #include <complex>
 
+#pragma clang diagnostic ignored "-Wunused-variable"
 // Tolerance for error
 const double tolerance = 1e-6;
 

catch/unit/deviceLib/unsafeAtomicAdd_RTC.cc

@@ -126,8 +126,8 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentRTCnounsafeatomicflag", "",
       void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f,
         HIP_LAUNCH_PARAM_BUFFER_SIZE,
         &size, HIP_LAUNCH_PARAM_END};
-      hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1, 1, 0,
-          nullptr, nullptr, config_d);
+      HIP_CHECK(hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1, 1, 0,
+                nullptr, nullptr, config_d));
       HIP_CHECK(hipDeviceSynchronize());
       REQUIRE(A_h[0] == INITIAL_VAL);
       REQUIRE(*result == 0);
@@ -218,8 +218,8 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentRTCunsafeatomicflag", "",
       void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f,
         HIP_LAUNCH_PARAM_BUFFER_SIZE,
         &size, HIP_LAUNCH_PARAM_END};
-      hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1, 1, 0,
-          nullptr, nullptr, config_d);
+      HIP_CHECK(hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1, 1, 0,
+                nullptr, nullptr, config_d));
       HIP_CHECK(hipDeviceSynchronize());
       REQUIRE(A_h[0] == INITIAL_VAL);
       REQUIRE(*result == 0);
@@ -306,8 +306,8 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentRTCwithoutflag", "",
       void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f,
                           HIP_LAUNCH_PARAM_BUFFER_SIZE,
                           &size, HIP_LAUNCH_PARAM_END};
-      hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1,
-                            1, 0, nullptr, nullptr, config_d);
+      HIP_CHECK(hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1,
+                            1, 0, nullptr, nullptr, config_d));
       HIP_CHECK(hipDeviceSynchronize());
       REQUIRE(A_h[0] == INITIAL_VAL);
       REQUIRE(*result == 0);
@@ -392,8 +392,8 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentRTCnounsafeatomicflag", "",
       void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f,
         HIP_LAUNCH_PARAM_BUFFER_SIZE,
         &size, HIP_LAUNCH_PARAM_END};
-      hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1, 1, 0,
-                            nullptr, nullptr, config_d);
+      HIP_CHECK(hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1, 1, 0,
+                            nullptr, nullptr, config_d));
       HIP_CHECK(hipDeviceSynchronize());
       REQUIRE(A_h[0] == INITIAL_VAL + INCREMENT_VAL);
       REQUIRE(*result == INITIAL_VAL);
@@ -480,8 +480,8 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentRTCunsafeatomicflag", "",
       void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f,
         HIP_LAUNCH_PARAM_BUFFER_SIZE,
         &size, HIP_LAUNCH_PARAM_END};
-      hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1, 1, 0,
-                            nullptr, nullptr, config_d);
+      HIP_CHECK(hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1, 1, 0,
+                            nullptr, nullptr, config_d));
       HIP_CHECK(hipDeviceSynchronize());
       REQUIRE(A_h[0] == INITIAL_VAL + INCREMENT_VAL);
       REQUIRE(*result == INITIAL_VAL);
@@ -568,8 +568,8 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentRTC", "",
       void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f,
                           HIP_LAUNCH_PARAM_BUFFER_SIZE,
                           &size, HIP_LAUNCH_PARAM_END};
-      hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1, 1, 0,
-                            nullptr, nullptr, config_d);
+      HIP_CHECK(hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1, 1, 0,
+                            nullptr, nullptr, config_d));
       HIP_CHECK(hipDeviceSynchronize());
       REQUIRE(A_h[0] == INITIAL_VAL + INCREMENT_VAL);
       REQUIRE(*result == INITIAL_VAL);

catch/unit/graph/hipGraphCycle.cc

@@ -132,7 +132,6 @@ TEST_CASE("Unit_hipGraph_BasicCyclic4") {
 
   HipTest::initArrays<int>(&X_d, &Y_d, nullptr, &X_h, &Y_h, nullptr, N, false);
 
-  constexpr size_t memSetVal = 9;
   hipGraphNode_t kMemCpyH2D_X, memcpyD2D, memcpyD2H_RC, emptyNode1;
 
   HIP_CHECK(hipStreamCreate(&stream));
@@ -192,7 +191,6 @@ TEST_CASE("Unit_hipGraph_BasicCyclic5") {
 
   HipTest::initArrays<int>(&X_d, &Y_d, nullptr, &X_h, &Y_h, nullptr, N, false);
 
-  constexpr size_t memSetVal = 9;
   hipGraphNode_t kMemCpyH2D_X, memcpyD2D, memcpyD2H_RC, emptyNode1, emptyNode2, emptyNode3;
 
   HIP_CHECK(hipStreamCreate(&stream));
@@ -241,4 +239,4 @@ TEST_CASE("Unit_hipGraph_BasicCyclic5") {
   HIP_CHECK(hipGraphExecDestroy(graphExec));
   HIP_CHECK(hipGraphDestroy(graph));
   HIP_CHECK(hipStreamDestroy(stream));
-}
+}

catch/unit/graph/hipStreamBeginCapture.cc

@@ -22,6 +22,7 @@ THE SOFTWARE.
 #include <hip_test_defgroups.hh>
 #include "stream_capture_common.hh" // NOLINT
 
+#pragma clang diagnostic ignored "-Wunused-variable"
 /**
  * @addtogroup hipStreamBeginCapture hipStreamBeginCapture
  * @{

catch/unit/graph/user_object_common.hh

@@ -24,6 +24,8 @@ THE SOFTWARE.
 
 #include <hip_test_common.hh>
 
+#pragma clang diagnostic ignored "-Wunused-function"
+
 struct BoxStruct {
   int count;
   BoxStruct() { INFO("Constructor called for Struct!\n"); }
@@ -58,4 +60,4 @@ void destroyFloatObj(void* ptr) {
   delete ptr2;
 }
 
-}  // anonymous namespace
+}  // anonymous namespace

catch/unit/kernel/hipShflTests.cc

@@ -150,7 +150,6 @@ static void runTest() {
 
   // verify the results
   REQUIRE(errors == compare(TransposeMatrix, cpuTransposeMatrix));
-  double eps = 1.0E-6;
   // free the resources on device side
   HIP_CHECK(hipFree(gpuMatrix));
   HIP_CHECK(hipFree(gpuTransposeMatrix));

catch/unit/memory/hipArrayCommon.hh

@@ -43,12 +43,7 @@ __global__ void readFromTexture(T* output, hipTextureObject_t texObj, size_t wid
   } else {
     const float v = y / (float)height;
     if (textureGather) {
-      // tex2Dgather not supported on __gfx90a__
-      #if !defined(__gfx90a__)
       output[y * width + x] = tex2Dgather<T>(texObj, u, v, ChannelToRead);
-      #else
-      #warning("tex2Dgather not supported on gfx90a");
-      #endif
     } else {
       output[y * width + x] = tex2D<T>(texObj, u, v);
     }

catch/unit/memory/hipArrayCreate.cc

@@ -54,7 +54,7 @@ static void ArrayCreate_DiffSizes(int gpu) {
   std::vector<std::pair<size_t, size_t>> runs {std::make_pair(NUM_W, NUM_H), std::make_pair(BIGNUM_W, BIGNUM_H)};
   for (const auto& size : runs) {
     std::array<HIP_ARRAY, ARRAY_LOOP> array;
-    size_t pavail, avail;
+    size_t pavail;
     HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr));
     HIP_ARRAY_DESCRIPTOR desc;
     desc.NumChannels = 1;

catch/unit/memory/hipFreeAsync.cc

@@ -26,7 +26,7 @@ TEST_CASE("Unit_hipFreeAsync_negative") {
   HIP_CHECK(hipSetDevice(0));
   void* p = nullptr;
   hipStream_t stream{nullptr};
-  hipStreamCreate(&stream);
+  HIP_CHECK(hipStreamCreate(&stream));
 
   SECTION("dev_ptr is nullptr") { REQUIRE(hipFreeAsync(nullptr, stream) != hipSuccess); }
 

catch/unit/memory/hipMalloc3D.cc

@@ -43,7 +43,7 @@ static void MemoryAlloc3DDiffSizes(int gpu) {
     size_t height{sizes}, depth{sizes};
     hipPitchedPtr devPitchedPtr[CHUNK_LOOP];
     hipExtent extent = make_hipExtent(width, height, depth);
-    size_t tot, avail, ptot, pavail;
+    size_t ptot, pavail;
     HIPCHECK(hipMemGetInfo(&pavail, &ptot));
     for (int i = 0; i < CHUNK_LOOP; i++) {
       HIPCHECK(hipMalloc3D(&devPitchedPtr[i], extent));

catch/unit/memory/hipMallocArray.cc

@@ -399,12 +399,22 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_happy", "", uint, int, int4, ushort, sho
     testArrayAsSurface<TestType>(arrayPtr, width, height);
   }
   SECTION("hipArrayTextureGather") {
-    height = 1024;
-    INFO("flag is hipArrayTextureGather");
-    INFO("height: " << height);
+    hipDeviceProp_t prop;
+    int device;
+    HIP_CHECK(hipGetDevice(&device));
+    HIP_CHECK(hipGetDeviceProperties(&prop, device));
+    // tex2Dgather not supported on gfx90a
+    if (std::string(prop.gcnArchName).find("gfx90a") == std::string::npos) {
+      height = 1024;
+      INFO("flag is hipArrayTextureGather");
+      INFO("height: " << height);
 
-    HIP_CHECK(hipMallocArray(&arrayPtr, &desc, width, height, hipArrayTextureGather));
-    testArrayAsTextureWithGather<TestType>(arrayPtr, width, height);
+      HIP_CHECK(hipMallocArray(&arrayPtr, &desc, width, height, hipArrayTextureGather));
+      testArrayAsTextureWithGather<TestType>(arrayPtr, width, height);
+    } else {
+      SUCCEED("tex2Dgather is not supported for gfx90a, Hence"
+               "skipping the testcase for this device " << device);
+    }
   }
 #endif
 

catch/unit/memory/hipMallocConcurrency.cc

@@ -82,7 +82,7 @@ static std::atomic<bool> g_thTestPassed{true};
 /**
  * Validates data consistency on supplied gpu
  */
-static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
+static bool validateMemoryOnGPU(int gpu) {
   int *A_d, *B_d, *C_d;
   int *A_h, *B_h, *C_h;
   bool TestPassed = true;
@@ -148,7 +148,7 @@ static bool regressAllocInLoop(int gpu) {
  * Validates data consistency on supplied gpu
  * In Multithreaded Environment
  */
-static bool validateMemoryOnGpuMThread(int gpu, bool concurOnOneGPU = false) {
+static bool validateMemoryOnGpuMThread(int gpu) {
   int *A_d, *B_d, *C_d;
   int *A_h, *B_h, *C_h;
   bool TestPassed = true;
@@ -213,7 +213,7 @@ static bool regressAllocInLoopMthread(int gpu) {
  * Thread func to regress alloc and check data consistency
  */
 static void threadFunc(int gpu) {
-  g_thTestPassed = regressAllocInLoopMthread(gpu) && validateMemoryOnGpuMThread(gpu, true);
+  g_thTestPassed = regressAllocInLoopMthread(gpu) && validateMemoryOnGpuMThread(gpu);
 
   UNSCOPED_INFO("thread execution status on gpu" << gpu << ":" << g_thTestPassed.load());
 }

catch/unit/memory/hipMemAdviseMmap.cc

@@ -31,7 +31,7 @@ THE SOFTWARE.
 */
 
 TEST_CASE("Unit_hipMemAdvise_MmapMem") {
-  int managed = 0, NUM_ELMS = 212992, PageableMem = 0;
+  int managed = 0, PageableMem = 0;
   INFO("The following are the attribute values related to HMM for"
          " device 0:\n");
   HIP_CHECK(hipDeviceGetAttribute(&managed,
@@ -54,6 +54,7 @@ TEST_CASE("Unit_hipMemAdvise_MmapMem") {
 #ifdef __linux__
     // For now this test is enabled only for linux platforms
     FILE *fptr;
+    int NUM_ELMS = 212992;
     fptr = fopen("ForTest1.txt", "w");
     for (int m = 0; m < NUM_ELMS; ++m) {
       putw(m, fptr);

catch/unit/memory/hipMemAdvise_old.cc

@@ -666,7 +666,6 @@ TEST_CASE("Unit_hipMemAdvise_TstAlignedAllocMem") {
   std::string gfxName(prop.gcnArchName);
 
   if (gfxName.find("xnack+") != std::string::npos) {
-    int stat = 0;
     int managedMem = 0, pageMemAccess = 0;
     HIP_CHECK(hipDeviceGetAttribute(&pageMemAccess,
               hipDeviceAttributePageableMemoryAccess, 0));

catch/unit/memory/hipMemRangeGetAttributes_old.cc

@@ -26,6 +26,7 @@ THE SOFTWARE.
 #include <hip_test_common.hh>
 #define MEM_SIZE 8192
 
+#ifdef __linux__
 static bool CheckError(hipError_t err, int LineNo) {
   if (err == hipSuccess) {
     WARN("Error expected but received hipSuccess at line no.:"
@@ -59,7 +60,6 @@ static int HmmAttrPrint() {
   return managed;
 }
 
-#ifdef __linux__
 /* Test Scenario: Testing basic working of hipMemRangeGetAttributes()
    api with different flags */
 

catch/unit/memory/hipMemcpy3D.cc

@@ -28,6 +28,8 @@ THE SOFTWARE.
 #include <resource_guards.hh>
 #include <utils.hh>
 
+#pragma clang diagnostic ignored "-Wunused-variable"
+
 TEST_CASE("Unit_hipMemcpy3D_Positive_Basic") {
   constexpr bool async = false;
 

catch/unit/memory/hipMemcpy3DAsync.cc

@@ -28,6 +28,8 @@ THE SOFTWARE.
 #include <resource_guards.hh>
 #include <utils.hh>
 
+#pragma clang diagnostic ignored "-Wunused-variable"
+
 TEST_CASE("Unit_hipMemcpy3DAsync_Positive_Basic") {
   constexpr bool async = true;
 

catch/unit/rtc/linker.cc

@@ -11,8 +11,7 @@
 #include <iterator>
 #include <vector>
 
-static constexpr auto NUM_THREADS{128};
-static constexpr auto NUM_BLOCKS{32};
+#pragma clang diagnostic ignored "-Wuninitialized"
 
 static constexpr auto src{
     R"(
@@ -125,4 +124,4 @@ TEST_CASE("Unit_RTC_LinkerAPI") {
   for (size_t i = 0; i < n; ++i) {
     REQUIRE(fabs(a * hX[i] + hY[i] - hOut[i]) <= fabs(hOut[i]) * 1e-6);
   }
-}
+}

catch/unit/stream/hipStreamCreateWithPriority.cc

@@ -686,8 +686,8 @@ void TestForMultipleStreamWithPriority(void) {
   }
   // launch kernels repeatedly on each of the low prioritiy stream
   for (int k = 0; k < LOW_PRIORITY_STREAMCOUNT; ++k) {
-    for (int i = 0; i < size; i += MEMCPYSIZE1) {
-      int j = i / sizeof(T);
+    for (size_t i = 0; i < size; i += MEMCPYSIZE1) {
+      size_t j = i / sizeof(T);
       if (enable_priority_low) {
         hipLaunchKernelGGL((memcpy_kernel<T>), dim3(GRIDSIZE), dim3(BLOCKSIZE),
         0, stream_low[k], dst_d_low[k] + j, src_d_low[k] + j,
@@ -697,8 +697,8 @@ void TestForMultipleStreamWithPriority(void) {
   }
   // launch kernels repeatedly on each of the normal prioritiy stream
   for (int k = 0; k < NORMAL_PRIORITY_STREAMCOUNT; ++k) {
-    for (int i = 0; i < size; i += MEMCPYSIZE1) {
-      int j = i / sizeof(T);
+    for (size_t i = 0; i < size; i += MEMCPYSIZE1) {
+      size_t j = i / sizeof(T);
       if (enable_priority_normal) {
         hipLaunchKernelGGL((memcpy_kernel<T>), dim3(GRIDSIZE), dim3(BLOCKSIZE),
         0, stream_normal[k], dst_d_normal[k] + j, src_d_normal[k] + j,
@@ -708,8 +708,8 @@ void TestForMultipleStreamWithPriority(void) {
   }
   // launch kernels repeatedly on each of the high prioritiy stream
   for (int k = 0; k < HIGH_PRIORITY_STREAMCOUNT; ++k) {
-    for (int i = 0; i < size; i += MEMCPYSIZE1) {
-      int j = i / sizeof(T);
+    for (size_t i = 0; i < size; i += MEMCPYSIZE1) {
+      size_t j = i / sizeof(T);
       if (enable_priority_high) {
         hipLaunchKernelGGL((memcpy_kernel<T>), dim3(GRIDSIZE), dim3(BLOCKSIZE),
         0, stream_high[k], dst_d_high[k] + j, src_d_high[k] + j,

catch/unit/surface/hipSurfaceObj1D.cc

@@ -19,7 +19,7 @@ THE SOFTWARE.
 #include <hip_test_common.hh>
 #include <hip_array_common.hh>
 #include <hip_texture_helper.hh>
-
+#pragma clang diagnostic ignored "-Wunused-variable"
 template <typename T>
 __global__ void
 surf1DKernelR(hipSurfaceObject_t surfaceObject,

catch/unit/surface/hipSurfaceObj2D.cc

@@ -19,7 +19,7 @@ THE SOFTWARE.
 #include <hip_test_common.hh>
 #include <hip_array_common.hh>
 #include <hip_texture_helper.hh>
-
+#pragma clang diagnostic ignored "-Wunused-variable"
 #define LOG_DATA 0
 
 template <typename T>

catch/unit/surface/hipSurfaceObj3D.cc

@@ -19,7 +19,7 @@ THE SOFTWARE.
 #include <hip_test_common.hh>
 #include <hip_array_common.hh>
 #include <hip_texture_helper.hh>
-
+#pragma clang diagnostic ignored "-Wunused-variable"
 template <typename T>
 __global__ void
 surf3DKernelR(hipSurfaceObject_t surfaceObject,

catch/unit/texture/hipTextureMipmapObj2D.cc

@@ -98,7 +98,7 @@ static void runMipMapTest(unsigned int width, unsigned int height, unsigned int
   hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData, textureObject, width,
                      (2 * mipmap_level));
   HIP_CHECK(hipGetLastError()); 
-  hipDeviceSynchronize();
+  HIP_CHECK(hipDeviceSynchronize());
 
   float* hOutputData = reinterpret_cast<float*>(malloc(size));
   REQUIRE(hOutputData != nullptr);

catch/unit/texture/hipTextureObjFetchVector.cc

@@ -21,7 +21,7 @@ THE SOFTWARE.
 #include <hip_array_common.hh>
 #include <vector>
 #include <iostream>
-
+#pragma clang diagnostic ignored "-Wunused-variable"
 template <typename T>
 __global__ void tex1dKernelFetch(T *val, hipTextureObject_t obj, int N) {
 #if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
@@ -79,7 +79,7 @@ bool runTest() {
   // Allocating the required buffer on gpu device
   T *texBuf, *texBufOut;
   T val[N], output[N];
-  hipGetLastError(); // Clear err due to negative tests
+  auto err = hipGetLastError(); // Clear err due to negative tests
   memset(output, 0, sizeof(output));
   std::srand(std::time(nullptr)); // use current time as seed for random generator