Rebase HIP 6.x + Update hip-tests

.clang-tidy

@@ -1,5 +1,5 @@
 ---
-Checks:          'readability-identifier-naming'
+Checks:          '-*,readability-*,modernize-*,clang-analyzer-*'
 WarningsAsErrors: ''
 HeaderFilterRegex: './src/.*'
 AnalyzeTemporaryDtors: false

.github/workflows/clang-tidy-format.yml

@@ -28,7 +28,7 @@ jobs:
           files-changed-only: true
           lines-changed-only: true
           extensions: 'cc,hh'
-          tidy-checks: 'readability-identifier-naming'
+          tidy-checks: '-*,readability-*,modernize-*,clang-analyzer-*'
       - name: Fail fast?!
         if: steps.linter.outputs.checks-failed > 0
         run: exit 1

CMakeLists.txt

@@ -28,7 +28,7 @@ add_compile_options(-Wno-format-extra-args -mf16c)
 set(CMAKE_EXPORT_COMPILE_COMMANDS ON)
 set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wno-duplicate-decl-specifier \
 -Wno-tautological-constant-compare  -Wno-c++20-extensions -Wno-unused-result \
--Wno-delete-abstract-non-virtual-dtor -Wno-deprecated-declarations -Wunused-command-line-argument")
+-Wno-delete-abstract-non-virtual-dtor -Wno-deprecated-declarations -Wunused-command-line-argument -gdwarf-4")
 
 # end temporary
 cmake_minimum_required(VERSION 3.20 FATAL_ERROR)
@@ -217,6 +217,8 @@ option(CHIP_USE_EXTERNAL_HIP_TESTS "Use Catch2 tests from the hip-tests submodul
 option(CHIP_ENABLE_NON_COMPLIANT_DEVICELIB_CODE "Enable non-compliant devicelib code such as calling LLVM builtins from inside kernel code. Enables certain unsigned long devicelib func variants" OFF)
 option(CHIP_FAST_MATH "Use native_ OpenCL functions which are fast but their precision is implementation defined" OFF)
 option(CHIP_USE_INTEL_USM "When enabled, cl_intel_unified_shared_memory extension, when available, will be used for HIP allocations in the OpenCL backend" ON)
+option(CATCH2_DISCOVER_TESTS_COMPILE_TIME "Discover the tests at compile time" ON)
+
 # This mitigation might be necessary on some systems with an older runtime. 
 # This mitigation makes memory resident (disable swapping) on the GPU
 # This has a significant impact on the cost of a GPU malloc 
@@ -355,6 +357,7 @@ add_dependencies(CHIP hipcc.bin hipconfig.bin)
 
 file(MAKE_DIRECTORY ${HIPCC_BUILD_PATH})
 file(CREATE_LINK hipcc.bin ${HIPCC_BUILD_PATH}/hipcc SYMBOLIC)
+file(CREATE_LINK hipconfig.bin ${HIPCC_BUILD_PATH}/hipconfig SYMBOLIC)
 
 add_subdirectory(bin)
 
@@ -566,7 +569,9 @@ set(_versionInfo "# Auto-generated by cmake\n")
 set(HIP_VERSION_MAJOR 5)
 set(HIP_VERSION_MINOR 1)
 set(HIP_VERSION_PATCH 0)
-set(HIP_VERSION_PATCH_GITHASH 0) # TODO
+set(HIP_VERSION_GITHASH 0)
+set(HIP_VERSION_PATCH_GITHASH "${HIP_VERSION_PATCH}.${HIP_VERSION_GITHASH}")
+set(HIP_VERSION "${HIP_VERSION_MAJOR}.${HIP_VERSION_MINOR}.${HIP_VERSION_PATCH}-${HIP_VERSION_PATCH_GITHASH}")
 
 add_to_config(_versionInfo HIP_VERSION_MAJOR "${HIP_VERSION_MAJOR}")
 add_to_config(_versionInfo HIP_VERSION_MINOR "${HIP_VERSION_MINOR}")
@@ -700,21 +705,13 @@ SET(HIP_PATH ${CMAKE_BINARY_DIR}) # pick up build sources, not install sources
 
 set(SAVED_WARN_DEPRECATED ${CMAKE_WARN_DEPRECATED})
 set(CMAKE_WARN_DEPRECATED OFF)
-if(CHIP_USE_EXTERNAL_HIP_TESTS)
-  add_subdirectory(hip-tests/catch catch)
-else()
-  add_subdirectory(HIP/tests/catch catch)
-endif()
-set(CMAKE_WARN_DEPRECATED ${SAVED_WARN_DEPRECATED})
-
 if(CHIP_BUILD_TESTS)
-  add_subdirectory(tests/cuda)
-  add_subdirectory(tests/devicelib)
-  add_subdirectory(tests/hiprtc)
-  add_subdirectory(tests/fromLibCeed)
-  add_subdirectory(tests/compiler)
-  add_subdirectory(tests/runtime)
-  add_subdirectory(tests/post-install)
+  if(CHIP_USE_EXTERNAL_HIP_TESTS)
+    add_subdirectory(hip-tests/catch catch)
+  else()
+    add_subdirectory(HIP/tests/catch catch)
+  endif()
+  add_subdirectory(tests)
 endif()
 
 if(CHIP_BUILD_SAMPLES)

bitcode/devicelib.cl

@@ -43,6 +43,84 @@
 #error __opencl_c_generic_address_space needed!
 #endif
 
+NOOPT void* device_malloc(unsigned int size) {return (void*)0;};
+NOOPT void device_free(void* ptr) {};
+
+// Given a 32/64-bit value exec mask and an integer value base (between 0 and WAVEFRONT_SIZE),
+// find the n-th (given by offset) set bit in the exec mask from the base bit, and return the bit position.
+// If not found, return -1.
+// In HIP long long is 64-bit integer. In OpenCL it's 128-bit integer.
+EXPORT int __chip__fns64(unsigned long int mask, unsigned int base, int offset) {
+  unsigned long int temp_mask = mask;
+  int temp_offset = offset;
+
+  if (offset == 0) {
+    temp_mask &= (1 << base);
+    temp_offset = 1;
+  }
+  else if (offset < 0) {
+    temp_mask = __builtin_bitreverse64(mask);
+    base = 63 - base;
+    temp_offset = -offset;
+  }
+
+  temp_mask = temp_mask & ((~0ULL) << base);
+  if (__builtin_popcountll(temp_mask) < temp_offset)
+    return -1;
+  int total = 0;
+  for (int i = 0x20; i > 0; i >>= 1) {
+    unsigned long int temp_mask_lo = temp_mask & ((1ULL << i) - 1);
+    int pcnt = __builtin_popcountll(temp_mask_lo);
+    if (pcnt < temp_offset) {
+      temp_mask = temp_mask >> i;
+      temp_offset -= pcnt;
+      total += i;
+    }
+    else {
+      temp_mask = temp_mask_lo;
+    }
+  }
+  if (offset < 0)
+    return 63 - total;
+  else
+    return total;
+}
+
+EXPORT int __chip__fns32(unsigned long int mask, unsigned int base, int offset) {
+  unsigned long int temp_mask = mask;
+  int temp_offset = offset;
+  if (offset == 0) {
+    temp_mask &= (1 << base);
+    temp_offset = 1;
+  }
+  else if (offset < 0) {
+    temp_mask = __builtin_bitreverse64(mask);
+    base = 63 - base;
+    temp_offset = -offset;
+  }
+  temp_mask = temp_mask & ((~0ULL) << base);
+  if (__builtin_popcountll(temp_mask) < temp_offset)
+    return -1;
+  int total = 0;
+  for (int i = 0x20; i > 0; i >>= 1) {
+    unsigned long int temp_mask_lo = temp_mask & ((1ULL << i) - 1);
+    int pcnt = __builtin_popcountll(temp_mask_lo);
+    if (pcnt < temp_offset) {
+      temp_mask = temp_mask >> i;
+      temp_offset -= pcnt;
+      total += i;
+    }
+    else {
+      temp_mask = temp_mask_lo;
+    }
+  }
+  if (offset < 0)
+    return 63 - total;
+  else
+    return total;
+}
+
+
 EXPORT unsigned /* long */ long int
 __chip_umul64hi(unsigned /* long */ long int x,
                 unsigned /* long */ long int y) {

docs/Features.md

@@ -41,6 +41,10 @@ CUDA features not present in HIP are unsupported unless explicitly stated otherw
 
 * few memory APIs (hipMemPrefetchAsync, hipMemAdvise)
 
+* Cooperative Groups API
+
+* Device-side malloc/free
+
 #### partially supported
 
 * Texture Objects of 1D/2D type are supported; 3D, LOD, Grad,

include/hip/devicelib/integer/int_intrinsics.hh

@@ -42,18 +42,33 @@ __bitinsert_u32(unsigned int src0, unsigned int src1, unsigned int offset,
   return __chip_bitinsert_u32(src0, src1, offset, width);
 }
 
-extern "C" __device__ uint64_t __chip_bitinsert_u64(uint64_t src0,
-                                                    uint64_t src1,
-                                                    uint64_t offset,
-                                                    uint64_t width);
-extern "C++" inline __device__ uint64_t __bitinsert_u64(uint64_t src0,
-                                                        uint64_t src1,
-                                                        uint64_t offset,
-                                                        uint64_t width) {
+extern "C" __device__ unsigned long long int
+__chip_bitinsert_u64(unsigned long long int src0, unsigned long long int src1,
+                     unsigned long long int offset,
+                     unsigned long long int width);
+extern "C++" inline __device__ unsigned long long int
+__bitinsert_u64(unsigned long long int src0, unsigned long long int src1,
+                unsigned long long int offset, unsigned long long int width) {
   return __chip_bitinsert_u64(src0, src1, offset, width);
 }
 #endif // CHIP_ENABLE_NON_COMPLIANT_DEVICELIB_CODE
 
+// int was replaced with int
+// int64_t was replaced with long long int
+extern "C" __device__ int __chip__fns32(unsigned long long int mask,
+                                        unsigned int base, int offset);
+extern "C++" inline __device__ int __fns32(unsigned long long int mask,
+                                           unsigned int base, int offset) {
+  return __chip__fns32(mask, base, offset);
+}
+
+extern "C" __device__ int __chip__fns64(unsigned long long int mask,
+                                        unsigned int base, int offset);
+extern "C++" inline __device__ int __fns64(unsigned long long int mask,
+                                           unsigned int base, int offset) {
+  return __chip__fns64(mask, base, offset);
+}
+
 extern "C" __device__ unsigned int __chip_brev(unsigned int x); // Custom
 extern "C++" inline __device__ unsigned int __brev(unsigned int x) {
   return __chip_brev(x);

include/hip/devicelib/type_casting_intrinsics.hh

@@ -25,50 +25,50 @@
 
 #include <hip/devicelib/macros.hh>
 
-extern "C++" inline __device__ float __double2float_rd(double x);
-extern "C++" inline __device__ float __double2float_rn(double x);
-extern "C++" inline __device__ float __double2float_ru(double x);
-extern "C++" inline __device__ float __double2float_rz(double x);
-extern "C++" inline __device__ int __double2hiint(double x);
-extern "C++" inline __device__ int __double2int_rd(double x);
-extern "C++" inline __device__ int __double2int_rn(double x);
-extern "C++" inline __device__ int __double2int_ru(double x);
-extern "C++" inline __device__ int __double2int_rz(double x);
-extern "C++" inline __device__ long long int __double2ll_rd(double x);
-extern "C++" inline __device__ long long int __double2ll_rn(double x);
-extern "C++" inline __device__ long long int __double2ll_ru(double x);
-extern "C++" inline __device__ long long int __double2ll_rz(double x);
-extern "C++" inline __device__ int __double2loint(double x);
-extern "C++" inline __device__ unsigned int __double2uint_rd(double x);
-extern "C++" inline __device__ unsigned int __double2uint_rn(double x);
-extern "C++" inline __device__ unsigned int __double2uint_ru(double x);
-extern "C++" inline __device__ unsigned int __double2uint_rz(double x);
-extern "C++" inline __device__ unsigned long long int __double2ull_rd(double x);
-extern "C++" inline __device__ unsigned long long int __double2ull_rn(double x);
-extern "C++" inline __device__ unsigned long long int __double2ull_ru(double x);
-extern "C++" inline __device__ unsigned long long int __double2ull_rz(double x);
+extern "C++" __device__ float __double2float_rd(double x);
+extern "C++" __device__ float __double2float_rn(double x);
+extern "C++" __device__ float __double2float_ru(double x);
+extern "C++" __device__ float __double2float_rz(double x);
+extern "C++" __device__ int __double2hiint(double x);
+extern "C++" __device__ int __double2int_rd(double x);
+extern "C++" __device__ int __double2int_rn(double x);
+extern "C++" __device__ int __double2int_ru(double x);
+extern "C++" __device__ int __double2int_rz(double x);
+extern "C++" __device__ long long int __double2ll_rd(double x);
+extern "C++" __device__ long long int __double2ll_rn(double x);
+extern "C++" __device__ long long int __double2ll_ru(double x);
+extern "C++" __device__ long long int __double2ll_rz(double x);
+extern "C++" __device__ int __double2loint(double x);
+extern "C++" __device__ unsigned int __double2uint_rd(double x);
+extern "C++" __device__ unsigned int __double2uint_rn(double x);
+extern "C++" __device__ unsigned int __double2uint_ru(double x);
+extern "C++" __device__ unsigned int __double2uint_rz(double x);
+extern "C++" __device__ unsigned long long int __double2ull_rd(double x);
+extern "C++" __device__ unsigned long long int __double2ull_rn(double x);
+extern "C++" __device__ unsigned long long int __double2ull_ru(double x);
+extern "C++" __device__ unsigned long long int __double2ull_rz(double x);
 
 extern "C" __device__ long long int __chip_double_as_longlong(double x);
 extern "C++" inline __device__ long long int __double_as_longlong(double x) {
   return __chip_double_as_longlong(x);
 }
 
-extern "C++" inline __device__ int __float2int_rd(float x);
-extern "C++" inline __device__ int __float2int_rn(float x);
-extern "C++" inline __device__ int __float2int_ru(float);
-extern "C++" inline __device__ int __float2int_rz(float x);
-extern "C++" inline __device__ long long int __float2ll_rd(float x);
-extern "C++" inline __device__ long long int __float2ll_rn(float x);
-extern "C++" inline __device__ long long int __float2ll_ru(float x);
-extern "C++" inline __device__ long long int __float2ll_rz(float x);
-extern "C++" inline __device__ unsigned int __float2uint_rd(float x);
-extern "C++" inline __device__ unsigned int __float2uint_rn(float x);
-extern "C++" inline __device__ unsigned int __float2uint_ru(float x);
-extern "C++" inline __device__ unsigned int __float2uint_rz(float x);
-extern "C++" inline __device__ unsigned long long int __float2ull_rd(float x);
-extern "C++" inline __device__ unsigned long long int __float2ull_rn(float x);
-extern "C++" inline __device__ unsigned long long int __float2ull_ru(float x);
-extern "C++" inline __device__ unsigned long long int __float2ull_rz(float x);
+extern "C++" __device__ int __float2int_rd(float x);
+extern "C++" __device__ int __float2int_rn(float x);
+extern "C++" __device__ int __float2int_ru(float);
+extern "C++" __device__ int __float2int_rz(float x);
+extern "C++" __device__ long long int __float2ll_rd(float x);
+extern "C++" __device__ long long int __float2ll_rn(float x);
+extern "C++" __device__ long long int __float2ll_ru(float x);
+extern "C++" __device__ long long int __float2ll_rz(float x);
+extern "C++" __device__ unsigned int __float2uint_rd(float x);
+extern "C++" __device__ unsigned int __float2uint_rn(float x);
+extern "C++" __device__ unsigned int __float2uint_ru(float x);
+extern "C++" __device__ unsigned int __float2uint_rz(float x);
+extern "C++" __device__ unsigned long long int __float2ull_rd(float x);
+extern "C++" __device__ unsigned long long int __float2ull_rn(float x);
+extern "C++" __device__ unsigned long long int __float2ull_ru(float x);
+extern "C++" __device__ unsigned long long int __float2ull_rz(float x);
 
 extern "C" __device__ int __chip_float_as_int(float x);
 extern "C++" inline __device__ int __float_as_int(float x) {
@@ -80,50 +80,50 @@ extern "C++" inline __device__ unsigned int __float_as_uint(float x) {
   return __chip_float_as_uint(x);
 }
 
-extern "C++" inline __device__ double __hiloint2double(int hi, int lo);
-extern "C++" inline __device__ double __int2double_rn(int x);
-extern "C++" inline __device__ float __int2float_rd(int x);
-extern "C++" inline __device__ float __int2float_rn(int x);
-extern "C++" inline __device__ float __int2float_ru(int x);
-extern "C++" inline __device__ float __int2float_rz(int x);
+extern "C++" __device__ double __hiloint2double(int hi, int lo);
+extern "C++" __device__ double __int2double_rn(int x);
+extern "C++" __device__ float __int2float_rd(int x);
+extern "C++" __device__ float __int2float_rn(int x);
+extern "C++" __device__ float __int2float_ru(int x);
+extern "C++" __device__ float __int2float_rz(int x);
 
 extern "C" __device__ float __chip_int_as_float(int x);
 extern "C++" inline __device__ float __int_as_float(int x) {
   return __chip_int_as_float(x);
 }
 
-extern "C++" inline __device__ double __ll2double_rd(long long int x);
-extern "C++" inline __device__ double __ll2double_rn(long long int x);
-extern "C++" inline __device__ double __ll2double_ru(long long int x);
-extern "C++" inline __device__ double __ll2double_rz(long long int x);
-extern "C++" inline __device__ float __ll2float_rd(long long int x);
-extern "C++" inline __device__ float __ll2float_rn(long long int x);
-extern "C++" inline __device__ float __ll2float_ru(long long int x);
-extern "C++" inline __device__ float __ll2float_rz(long long int x);
+extern "C++" __device__ double __ll2double_rd(long long int x);
+extern "C++" __device__ double __ll2double_rn(long long int x);
+extern "C++" __device__ double __ll2double_ru(long long int x);
+extern "C++" __device__ double __ll2double_rz(long long int x);
+extern "C++" __device__ float __ll2float_rd(long long int x);
+extern "C++" __device__ float __ll2float_rn(long long int x);
+extern "C++" __device__ float __ll2float_ru(long long int x);
+extern "C++" __device__ float __ll2float_rz(long long int x);
 
 extern "C" __device__ double __chip_longlong_as_double(long long int x);
 extern "C++" inline __device__ double __longlong_as_double(long long int x) {
   return __chip_longlong_as_double(x);
 }
 
-extern "C++" inline __device__ double __uint2double_rn(unsigned int x);
-extern "C++" inline __device__ float __uint2float_rd(unsigned int x);
-extern "C++" inline __device__ float __uint2float_rn(unsigned int x);
-extern "C++" inline __device__ float __uint2float_ru(unsigned int x);
-extern "C++" inline __device__ float __uint2float_rz(unsigned int x);
+extern "C++" __device__ double __uint2double_rn(unsigned int x);
+extern "C++" __device__ float __uint2float_rd(unsigned int x);
+extern "C++" __device__ float __uint2float_rn(unsigned int x);
+extern "C++" __device__ float __uint2float_ru(unsigned int x);
+extern "C++" __device__ float __uint2float_rz(unsigned int x);
 
 extern "C" __device__ float __chip_uint_as_float(uint x);
 extern "C++" inline __device__ float __uint_as_float(unsigned int x) {
   return __chip_uint_as_float(x);
 }
 
-extern "C++" inline __device__ double __ull2double_rd(unsigned long long int x);
-extern "C++" inline __device__ double __ull2double_rn(unsigned long long int x);
-extern "C++" inline __device__ double __ull2double_ru(unsigned long long int x);
-extern "C++" inline __device__ double __ull2double_rz(unsigned long long int x);
-extern "C++" inline __device__ float __ull2float_rd(unsigned long long int x);
-extern "C++" inline __device__ float __ull2float_rn(unsigned long long int x);
-extern "C++" inline __device__ float __ull2float_ru(unsigned long long int x);
-extern "C++" inline __device__ float __ull2float_rz(unsigned long long int x);
+extern "C++" __device__ double __ull2double_rd(unsigned long long int x);
+extern "C++" __device__ double __ull2double_rn(unsigned long long int x);
+extern "C++" __device__ double __ull2double_ru(unsigned long long int x);
+extern "C++" __device__ double __ull2double_rz(unsigned long long int x);
+extern "C++" __device__ float __ull2float_rd(unsigned long long int x);
+extern "C++" __device__ float __ull2float_rn(unsigned long long int x);
+extern "C++" __device__ float __ull2float_ru(unsigned long long int x);
+extern "C++" __device__ float __ull2float_rz(unsigned long long int x);
 
 #endif // include guard