Reapply "[clang][HIP] Make some math not not work with AMDGCN SPIR-V #128360" #129306

AlexVlx · 2025-02-28T21:01:35Z

This reapplies #128360, the only change being that the modified tests also checks for the availability of the SPIRV target.

…l_is_bad_for_spirv

) Do not hardcode `address_space(5)` (`private`) in the ROCDL interface, as that breaks SPIRV generation (the latter uses 0). Add test. In the long run we should stop using ROCDL inline.

llvmbot · 2025-02-28T21:02:10Z

@llvm/pr-subscribers-clang

Author: Alex Voicu (AlexVlx)

Changes

This reapplies #128360, the only change being that the modified tests also checks for the availability of the SPIRV target.

Patch is 183.47 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/129306.diff

3 Files Affected:

(modified) clang/lib/Headers/__clang_hip_libdevice_declares.h (+12-20)
(modified) clang/lib/Headers/__clang_hip_math.h (+12-16)
(modified) clang/test/Headers/__clang_hip_math.hip (+1656)

diff --git a/clang/lib/Headers/__clang_hip_libdevice_declares.h b/clang/lib/Headers/__clang_hip_libdevice_declares.h
index f15198b3d9f93..fa8d918248dd0 100644
--- a/clang/lib/Headers/__clang_hip_libdevice_declares.h
+++ b/clang/lib/Headers/__clang_hip_libdevice_declares.h
@@ -14,6 +14,8 @@
 #include "hip/hip_version.h"
 #endif // __has_include("hip/hip_version.h")
 
+#define __PRIVATE_AS __attribute__((opencl_private))
+
 #ifdef __cplusplus
 extern "C" {
 #endif
@@ -55,8 +57,7 @@ __device__ __attribute__((const)) float __ocml_fmax_f32(float, float);
 __device__ __attribute__((const)) float __ocml_fmin_f32(float, float);
 __device__ __attribute__((const)) __device__ float __ocml_fmod_f32(float,
                                                                    float);
-__device__ float __ocml_frexp_f32(float,
-                                  __attribute__((address_space(5))) int *);
+__device__ float __ocml_frexp_f32(float, __PRIVATE_AS int *);
 __device__ __attribute__((const)) float __ocml_hypot_f32(float, float);
 __device__ __attribute__((const)) int __ocml_ilogb_f32(float);
 __device__ __attribute__((const)) int __ocml_isfinite_f32(float);
@@ -74,8 +75,7 @@ __device__ __attribute__((pure)) float __ocml_native_log2_f32(float);
 __device__ __attribute__((const)) float __ocml_logb_f32(float);
 __device__ __attribute__((pure)) float __ocml_log_f32(float);
 __device__ __attribute__((pure)) float __ocml_native_log_f32(float);
-__device__ float __ocml_modf_f32(float,
-                                 __attribute__((address_space(5))) float *);
+__device__ float __ocml_modf_f32(float, __PRIVATE_AS float *);
 __device__ __attribute__((const)) float __ocml_nearbyint_f32(float);
 __device__ __attribute__((const)) float __ocml_nextafter_f32(float, float);
 __device__ __attribute__((const)) float __ocml_len3_f32(float, float, float);
@@ -87,8 +87,7 @@ __device__ __attribute__((pure)) float __ocml_pow_f32(float, float);
 __device__ __attribute__((pure)) float __ocml_pown_f32(float, int);
 __device__ __attribute__((pure)) float __ocml_rcbrt_f32(float);
 __device__ __attribute__((const)) float __ocml_remainder_f32(float, float);
-__device__ float __ocml_remquo_f32(float, float,
-                                   __attribute__((address_space(5))) int *);
+__device__ float __ocml_remquo_f32(float, float, __PRIVATE_AS int *);
 __device__ __attribute__((const)) float __ocml_rhypot_f32(float, float);
 __device__ __attribute__((const)) float __ocml_rint_f32(float);
 __device__ __attribute__((const)) float __ocml_rlen3_f32(float, float, float);
@@ -99,10 +98,8 @@ __device__ __attribute__((pure)) float __ocml_rsqrt_f32(float);
 __device__ __attribute__((const)) float __ocml_scalb_f32(float, float);
 __device__ __attribute__((const)) float __ocml_scalbn_f32(float, int);
 __device__ __attribute__((const)) int __ocml_signbit_f32(float);
-__device__ float __ocml_sincos_f32(float,
-                                   __attribute__((address_space(5))) float *);
-__device__ float __ocml_sincospi_f32(float,
-                                     __attribute__((address_space(5))) float *);
+__device__ float __ocml_sincos_f32(float, __PRIVATE_AS float *);
+__device__ float __ocml_sincospi_f32(float, __PRIVATE_AS float *);
 __device__ float __ocml_sin_f32(float);
 __device__ float __ocml_native_sin_f32(float);
 __device__ __attribute__((pure)) float __ocml_sinh_f32(float);
@@ -176,8 +173,7 @@ __device__ __attribute__((const)) double __ocml_fma_f64(double, double, double);
 __device__ __attribute__((const)) double __ocml_fmax_f64(double, double);
 __device__ __attribute__((const)) double __ocml_fmin_f64(double, double);
 __device__ __attribute__((const)) double __ocml_fmod_f64(double, double);
-__device__ double __ocml_frexp_f64(double,
-                                   __attribute__((address_space(5))) int *);
+__device__ double __ocml_frexp_f64(double, __PRIVATE_AS int *);
 __device__ __attribute__((const)) double __ocml_hypot_f64(double, double);
 __device__ __attribute__((const)) int __ocml_ilogb_f64(double);
 __device__ __attribute__((const)) int __ocml_isfinite_f64(double);
@@ -192,8 +188,7 @@ __device__ __attribute__((pure)) double __ocml_log1p_f64(double);
 __device__ __attribute__((pure)) double __ocml_log2_f64(double);
 __device__ __attribute__((const)) double __ocml_logb_f64(double);
 __device__ __attribute__((pure)) double __ocml_log_f64(double);
-__device__ double __ocml_modf_f64(double,
-                                  __attribute__((address_space(5))) double *);
+__device__ double __ocml_modf_f64(double, __PRIVATE_AS double *);
 __device__ __attribute__((const)) double __ocml_nearbyint_f64(double);
 __device__ __attribute__((const)) double __ocml_nextafter_f64(double, double);
 __device__ __attribute__((const)) double __ocml_len3_f64(double, double,
@@ -206,8 +201,7 @@ __device__ __attribute__((pure)) double __ocml_pow_f64(double, double);
 __device__ __attribute__((pure)) double __ocml_pown_f64(double, int);
 __device__ __attribute__((pure)) double __ocml_rcbrt_f64(double);
 __device__ __attribute__((const)) double __ocml_remainder_f64(double, double);
-__device__ double __ocml_remquo_f64(double, double,
-                                    __attribute__((address_space(5))) int *);
+__device__ double __ocml_remquo_f64(double, double, __PRIVATE_AS int *);
 __device__ __attribute__((const)) double __ocml_rhypot_f64(double, double);
 __device__ __attribute__((const)) double __ocml_rint_f64(double);
 __device__ __attribute__((const)) double __ocml_rlen3_f64(double, double,
@@ -219,10 +213,8 @@ __device__ __attribute__((pure)) double __ocml_rsqrt_f64(double);
 __device__ __attribute__((const)) double __ocml_scalb_f64(double, double);
 __device__ __attribute__((const)) double __ocml_scalbn_f64(double, int);
 __device__ __attribute__((const)) int __ocml_signbit_f64(double);
-__device__ double __ocml_sincos_f64(double,
-                                    __attribute__((address_space(5))) double *);
-__device__ double
-__ocml_sincospi_f64(double, __attribute__((address_space(5))) double *);
+__device__ double __ocml_sincos_f64(double, __PRIVATE_AS double *);
+__device__ double __ocml_sincospi_f64(double, __PRIVATE_AS double *);
 __device__ double __ocml_sin_f64(double);
 __device__ __attribute__((pure)) double __ocml_sinh_f64(double);
 __device__ double __ocml_sinpi_f64(double);
diff --git a/clang/lib/Headers/__clang_hip_math.h b/clang/lib/Headers/__clang_hip_math.h
index 8468751d9de26..bf8517bc3a507 100644
--- a/clang/lib/Headers/__clang_hip_math.h
+++ b/clang/lib/Headers/__clang_hip_math.h
@@ -33,6 +33,9 @@
 #define __DEVICE__ static __device__ inline __attribute__((always_inline))
 #endif
 
+#pragma push_macro("__PRIVATE_AS")
+
+#define __PRIVATE_AS __attribute__((opencl_private))
 // Device library provides fast low precision and slow full-recision
 // implementations for some functions. Which one gets selected depends on
 // __CLANG_GPU_APPROX_TRANSCENDENTALS__ which gets defined by clang if
@@ -512,8 +515,7 @@ float modff(float __x, float *__iptr) {
 #ifdef __OPENMP_AMDGCN__
 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
 #endif
-  float __r =
-      __ocml_modf_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
+  float __r = __ocml_modf_f32(__x, (__PRIVATE_AS float *)&__tmp);
   *__iptr = __tmp;
   return __r;
 }
@@ -595,8 +597,7 @@ float remquof(float __x, float __y, int *__quo) {
 #ifdef __OPENMP_AMDGCN__
 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
 #endif
-  float __r = __ocml_remquo_f32(
-      __x, __y, (__attribute__((address_space(5))) int *)&__tmp);
+  float __r = __ocml_remquo_f32(__x, __y, (__PRIVATE_AS int *)&__tmp);
   *__quo = __tmp;
 
   return __r;
@@ -657,8 +658,7 @@ void sincosf(float __x, float *__sinptr, float *__cosptr) {
 #ifdef __CLANG_CUDA_APPROX_TRANSCENDENTALS__
   __sincosf(__x, __sinptr, __cosptr);
 #else
-  *__sinptr =
-      __ocml_sincos_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
+  *__sinptr = __ocml_sincos_f32(__x, (__PRIVATE_AS float *)&__tmp);
   *__cosptr = __tmp;
 #endif
 }
@@ -669,8 +669,7 @@ void sincospif(float __x, float *__sinptr, float *__cosptr) {
 #ifdef __OPENMP_AMDGCN__
 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
 #endif
-  *__sinptr = __ocml_sincospi_f32(
-      __x, (__attribute__((address_space(5))) float *)&__tmp);
+  *__sinptr = __ocml_sincospi_f32(__x, (__PRIVATE_AS float *)&__tmp);
   *__cosptr = __tmp;
 }
 
@@ -913,8 +912,7 @@ double modf(double __x, double *__iptr) {
 #ifdef __OPENMP_AMDGCN__
 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
 #endif
-  double __r =
-      __ocml_modf_f64(__x, (__attribute__((address_space(5))) double *)&__tmp);
+  double __r = __ocml_modf_f64(__x, (__PRIVATE_AS double *)&__tmp);
   *__iptr = __tmp;
 
   return __r;
@@ -1004,8 +1002,7 @@ double remquo(double __x, double __y, int *__quo) {
 #ifdef __OPENMP_AMDGCN__
 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
 #endif
-  double __r = __ocml_remquo_f64(
-      __x, __y, (__attribute__((address_space(5))) int *)&__tmp);
+  double __r = __ocml_remquo_f64(__x, __y, (__PRIVATE_AS int *)&__tmp);
   *__quo = __tmp;
 
   return __r;
@@ -1065,8 +1062,7 @@ void sincos(double __x, double *__sinptr, double *__cosptr) {
 #ifdef __OPENMP_AMDGCN__
 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
 #endif
-  *__sinptr = __ocml_sincos_f64(
-      __x, (__attribute__((address_space(5))) double *)&__tmp);
+  *__sinptr = __ocml_sincos_f64(__x, (__PRIVATE_AS double *)&__tmp);
   *__cosptr = __tmp;
 }
 
@@ -1076,8 +1072,7 @@ void sincospi(double __x, double *__sinptr, double *__cosptr) {
 #ifdef __OPENMP_AMDGCN__
 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
 #endif
-  *__sinptr = __ocml_sincospi_f64(
-      __x, (__attribute__((address_space(5))) double *)&__tmp);
+  *__sinptr = __ocml_sincospi_f64(__x, (__PRIVATE_AS double *)&__tmp);
   *__cosptr = __tmp;
 }
 
@@ -1322,6 +1317,7 @@ __host__ inline static int max(int __arg1, int __arg2) {
 #endif
 
 #pragma pop_macro("__DEVICE__")
+#pragma pop_macro("__PRIVATE_AS")
 #pragma pop_macro("__RETURN_TYPE")
 #pragma pop_macro("__FAST_OR_SLOW")
 
diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip
index e4254d1e64bec..d448ab134ca4d 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -1,5 +1,6 @@
 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
 // REQUIRES: amdgpu-registered-target
+// REQUIRES: spirv-registered-target
 
 // Test without OCML_BASIC_ROUNDED_OPERATIONS
 // RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
@@ -26,6 +27,14 @@
 // RUN:   -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -O1 -fgpu-approx-transcendentals -o - \
 // RUN:   -D__HIPCC_RTC__ | FileCheck -check-prefixes=CHECK,APPROX %s
 
+// Check that we use the AMDGCNSPIRV address space map
+// RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
+// RUN:   -internal-isystem %S/../../lib/Headers/cuda_wrappers \
+// RUN:   -internal-isystem %S/Inputs/include \
+// RUN:   -triple spirv64-amd-amdhsa -aux-triple x86_64-unknown-unknown \
+// RUN:   -emit-llvm %s -fcuda-is-device -O1 -o - \
+// RUN:   -D__HIPCC_RTC__ | FileCheck -check-prefixes=AMDGCNSPIRV %s
+
 #define BOOL_TYPE int
 typedef unsigned long long uint64_t;
 
@@ -57,6 +66,30 @@ typedef unsigned long long uint64_t;
 // CHECK-NEXT:    [[RETVAL_2_I:%.*]] = phi i64 [ 0, [[CLEANUP_I]] ], [ [[__R_0_I]], [[WHILE_COND_I]] ]
 // CHECK-NEXT:    ret i64 [[RETVAL_2_I]]
 //
+// AMDGCNSPIRV-LABEL: @test___make_mantissa_base8(
+// AMDGCNSPIRV-NEXT:  entry:
+// AMDGCNSPIRV-NEXT:    br label [[WHILE_COND_I:%.*]]
+// AMDGCNSPIRV:       while.cond.i:
+// AMDGCNSPIRV-NEXT:    [[__TAGP_ADDR_0_I:%.*]] = phi ptr addrspace(4) [ [[P:%.*]], [[ENTRY:%.*]] ], [ [[__TAGP_ADDR_1_I:%.*]], [[WHILE_BODY_I:%.*]] ]
+// AMDGCNSPIRV-NEXT:    [[__R_0_I:%.*]] = phi i64 [ 0, [[ENTRY]] ], [ [[__R_1_I:%.*]], [[WHILE_BODY_I]] ]
+// AMDGCNSPIRV-NEXT:    [[TMP0:%.*]] = load i8, ptr addrspace(4) [[__TAGP_ADDR_0_I]], align 1, !tbaa [[TBAA5:![0-9]+]]
+// AMDGCNSPIRV-NEXT:    [[CMP_NOT_I:%.*]] = icmp eq i8 [[TMP0]], 0
+// AMDGCNSPIRV-NEXT:    br i1 [[CMP_NOT_I]], label [[_ZL21__MAKE_MANTISSA_BASE8PKC_EXIT:%.*]], label [[WHILE_BODY_I]]
+// AMDGCNSPIRV:       while.body.i:
+// AMDGCNSPIRV-NEXT:    [[TMP1:%.*]] = and i8 [[TMP0]], -8
+// AMDGCNSPIRV-NEXT:    [[OR_COND_I:%.*]] = icmp eq i8 [[TMP1]], 48
+// AMDGCNSPIRV-NEXT:    [[MUL_I:%.*]] = shl i64 [[__R_0_I]], 3
+// AMDGCNSPIRV-NEXT:    [[CONV5_I:%.*]] = zext nneg i8 [[TMP0]] to i64
+// AMDGCNSPIRV-NEXT:    [[ADD_I:%.*]] = add i64 [[MUL_I]], -48
+// AMDGCNSPIRV-NEXT:    [[SUB_I:%.*]] = add i64 [[ADD_I]], [[CONV5_I]]
+// AMDGCNSPIRV-NEXT:    [[__TAGP_ADDR_1_I_IDX:%.*]] = zext i1 [[OR_COND_I]] to i64
+// AMDGCNSPIRV-NEXT:    [[__TAGP_ADDR_1_I]] = getelementptr inbounds nuw i8, ptr addrspace(4) [[__TAGP_ADDR_0_I]], i64 [[__TAGP_ADDR_1_I_IDX]]
+// AMDGCNSPIRV-NEXT:    [[__R_1_I]] = select i1 [[OR_COND_I]], i64 [[SUB_I]], i64 [[__R_0_I]]
+// AMDGCNSPIRV-NEXT:    br i1 [[OR_COND_I]], label [[WHILE_COND_I]], label [[_ZL21__MAKE_MANTISSA_BASE8PKC_EXIT]], !llvm.loop [[LOOP8:![0-9]+]]
+// AMDGCNSPIRV:       _ZL21__make_mantissa_base8PKc.exit:
+// AMDGCNSPIRV-NEXT:    [[RETVAL_2_I:%.*]] = phi i64 [ 0, [[WHILE_BODY_I]] ], [ [[__R_0_I]], [[WHILE_COND_I]] ]
+// AMDGCNSPIRV-NEXT:    ret i64 [[RETVAL_2_I]]
+//
 extern "C" __device__ uint64_t test___make_mantissa_base8(const char *p) {
   return __make_mantissa_base8(p);
 }
@@ -89,6 +122,30 @@ extern "C" __device__ uint64_t test___make_mantissa_base8(const char *p) {
 // CHECK-NEXT:    [[RETVAL_2_I:%.*]] = phi i64 [ 0, [[CLEANUP_I]] ], [ [[__R_0_I]], [[WHILE_COND_I]] ]
 // CHECK-NEXT:    ret i64 [[RETVAL_2_I]]
 //
+// AMDGCNSPIRV-LABEL: @test___make_mantissa_base10(
+// AMDGCNSPIRV-NEXT:  entry:
+// AMDGCNSPIRV-NEXT:    br label [[WHILE_COND_I:%.*]]
+// AMDGCNSPIRV:       while.cond.i:
+// AMDGCNSPIRV-NEXT:    [[__TAGP_ADDR_0_I:%.*]] = phi ptr addrspace(4) [ [[P:%.*]], [[ENTRY:%.*]] ], [ [[__TAGP_ADDR_1_I:%.*]], [[WHILE_BODY_I:%.*]] ]
+// AMDGCNSPIRV-NEXT:    [[__R_0_I:%.*]] = phi i64 [ 0, [[ENTRY]] ], [ [[__R_1_I:%.*]], [[WHILE_BODY_I]] ]
+// AMDGCNSPIRV-NEXT:    [[TMP0:%.*]] = load i8, ptr addrspace(4) [[__TAGP_ADDR_0_I]], align 1, !tbaa [[TBAA5]]
+// AMDGCNSPIRV-NEXT:    [[CMP_NOT_I:%.*]] = icmp eq i8 [[TMP0]], 0
+// AMDGCNSPIRV-NEXT:    br i1 [[CMP_NOT_I]], label [[_ZL22__MAKE_MANTISSA_BASE10PKC_EXIT:%.*]], label [[WHILE_BODY_I]]
+// AMDGCNSPIRV:       while.body.i:
+// AMDGCNSPIRV-NEXT:    [[TMP1:%.*]] = add i8 [[TMP0]], -48
+// AMDGCNSPIRV-NEXT:    [[OR_COND_I:%.*]] = icmp ult i8 [[TMP1]], 10
+// AMDGCNSPIRV-NEXT:    [[MUL_I:%.*]] = mul i64 [[__R_0_I]], 10
+// AMDGCNSPIRV-NEXT:    [[CONV5_I:%.*]] = zext nneg i8 [[TMP0]] to i64
+// AMDGCNSPIRV-NEXT:    [[ADD_I:%.*]] = add i64 [[MUL_I]], -48
+// AMDGCNSPIRV-NEXT:    [[SUB_I:%.*]] = add i64 [[ADD_I]], [[CONV5_I]]
+// AMDGCNSPIRV-NEXT:    [[__TAGP_ADDR_1_I_IDX:%.*]] = zext i1 [[OR_COND_I]] to i64
+// AMDGCNSPIRV-NEXT:    [[__TAGP_ADDR_1_I]] = getelementptr inbounds nuw i8, ptr addrspace(4) [[__TAGP_ADDR_0_I]], i64 [[__TAGP_ADDR_1_I_IDX]]
+// AMDGCNSPIRV-NEXT:    [[__R_1_I]] = select i1 [[OR_COND_I]], i64 [[SUB_I]], i64 [[__R_0_I]]
+// AMDGCNSPIRV-NEXT:    br i1 [[OR_COND_I]], label [[WHILE_COND_I]], label [[_ZL22__MAKE_MANTISSA_BASE10PKC_EXIT]], !llvm.loop [[LOOP11:![0-9]+]]
+// AMDGCNSPIRV:       _ZL22__make_mantissa_base10PKc.exit:
+// AMDGCNSPIRV-NEXT:    [[RETVAL_2_I:%.*]] = phi i64 [ 0, [[WHILE_BODY_I]] ], [ [[__R_0_I]], [[WHILE_COND_I]] ]
+// AMDGCNSPIRV-NEXT:    ret i64 [[RETVAL_2_I]]
+//
 extern "C" __device__ uint64_t test___make_mantissa_base10(const char *p) {
   return __make_mantissa_base10(p);
 }
@@ -131,6 +188,44 @@ extern "C" __device__ uint64_t test___make_mantissa_base10(const char *p) {
 // CHECK-NEXT:    [[RETVAL_2_I:%.*]] = phi i64 [ 0, [[CLEANUP_I]] ], [ [[__R_0_I]], [[WHILE_COND_I]] ]
 // CHECK-NEXT:    ret i64 [[RETVAL_2_I]]
 //
+// AMDGCNSPIRV-LABEL: @test___make_mantissa_base16(
+// AMDGCNSPIRV-NEXT:  entry:
+// AMDGCNSPIRV-NEXT:    br label [[WHILE_COND_I:%.*]]
+// AMDGCNSPIRV:       while.cond.i:
+// AMDGCNSPIRV-NEXT:    [[__TAGP_ADDR_0_I:%.*]] = phi ptr addrspace(4) [ [[P:%.*]], [[ENTRY:%.*]] ], [ [[__TAGP_ADDR_1_I:%.*]], [[CLEANUP_I:%.*]] ]
+// AMDGCNSPIRV-NEXT:    [[__R_0_I:%.*]] = phi i64 [ 0, [[ENTRY]] ], [ [[__R_2_I:%.*]], [[CLEANUP_I]] ]
+// AMDGCNSPIRV-NEXT:    [[TMP0:%.*]] = load i8, ptr addrspace(4) [[__TAGP_ADDR_0_I]], align 1, !tbaa [[TBAA5]]
+// AMDGCNSPIRV-NEXT:    [[CMP_NOT_I:%.*]] = icmp eq i8 [[TMP0]], 0
+// AMDGCNSPIRV-NEXT:    br i1 [[CMP_NOT_I]], label [[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// AMDGCNSPIRV:       while.body.i:
+// AMDGCNSPIRV-NEXT:    [[TMP1:%.*]] = add i8 [[TMP0]], -48
+// AMDGCNSPIRV-NEXT:    [[OR_COND_I:%.*]] = icmp ult i8 [[TMP1]], 10
+// AMDGCNSPIRV-NEXT:    br i1 [[OR_COND_I]], label [[IF_END31_I:%.*]], label [[IF_ELSE_I:%.*]]
+// AMDGCNSPIRV:       if.else.i:
+// AMDGCNSPIRV-NEXT:    [[TMP2:%.*]] = add i8 [[TMP0]], -97
+// AMDGCNSPIRV-NEXT:    [[OR_COND33_I:%.*]] = icmp ult i8 [[TMP2]], 6
+// AMDGCNSPIRV-NEXT:    br i1 [[OR_COND33_I]], label [[IF_END31_I]], label [[IF_ELSE17_I:%.*]]
+// AMDGCNSPIRV:       if.else17.i:
+// AMDGCNSPIRV-NEXT:    [[TMP3:%.*]] = add i8 [[TMP0]], -65
+// AMDGCNSPIRV-NEXT:    [[OR_COND34_I:%.*]] = icmp ult i8 [[TMP3]], 6
+// AMDGCNSPIRV-NEXT:    br i1 [[OR_COND34_I]], label [[IF_END31_I]], label [[CLEANUP_I]]
+// AMDGCNSPIRV:       if.end31.i:
+// AMDGCNSPIRV-NEXT:    [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I]] ], [ -87, [[IF_ELSE_I]] ], [ -55, [[IF_ELSE17_I]] ]
+// AMDGCNSPIRV-NEXT:    [[MUL24_I:%.*]] = shl i64 [[__R_0_I]], 4
+// AMDGCNSPIRV-NEXT:    [[CONV25_I:%.*]] = zext nneg i8 [[TMP0]] to i64
+// AMDGCNSPIRV-NEXT:    [[ADD26_I:%.*]] = add i64 [[MUL24_I]], [[DOTSINK]]
+// AMDGCNSPIRV-NEXT:    [[ADD28_I:%.*]] = add i64 [[ADD26_I]], [[CONV25_I]]
+// AMDGCNSPIRV-NEXT:    [[INCDEC_PTR_I:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(4) [[__TAGP_ADDR_0_I]], i64 1
+// AMDGCNSPIRV-NEXT:    br label [[CLEANUP_I]]
+// AMDGCNSPIRV:       cleanup.i:
+// AMDGCNSPIRV-NEXT:    [[__TAGP_ADDR_1_I]] = phi ptr addrspace(4) [ [[INCDEC_PTR_I]], [[IF_END31_I]] ], [ [[__TAGP_ADDR_0_I]], [[IF_ELSE17_I]] ]
+// AMDGCNSPIRV-NEXT:    [[__R_2_I]] = phi i64 [ [[ADD28_I]], [[IF_END31_I]] ], [ [[__R_0_I]], [[IF_ELSE17_I]] ]
+// AMDGCNSPIRV-NEXT:    [[COND_I:%.*]] = phi i1 [ true, [[IF_END31_I]] ], [ false, [[IF_ELSE17_I]] ]
+// AMDGCNSPIRV-NEXT:    br i1 [[COND_I]], label [[WHILE_COND_I]], label [[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT]], !llvm.loop [[LOOP12:![0-9]+]]
+// AMDGCNSPIRV:       _ZL22__make_mantissa_base16PKc.exit:
+// AMDGCNSPIRV-NEXT:    [[RETVAL_2_I:%.*]] = phi i64 [ 0, [[CLEANUP_I]] ], [ [[__R_0_I]], [[WHILE_COND_I]] ]
+// AMDGCNSPIRV-NEXT:    ret i64 [[RETVAL_2_I]]
+//
 extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) {
   return __make_mantissa_base16(p);
 }
@@ -226,6 +321,89 @@ extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) {
 // CHECK-NEXT:    [[RETVAL_0_I:%.*]] = phi i64 [ 0, [[CLEANUP_I_I]] ], [ [[__R_0_I_I]], [[WHILE_COND_I_I]] ], [ 0, [[CLEANUP_I36_I]] ], [ [[__R_0_I32_I]], [[WHILE_COND_I30_I]] ], [ 0, [[CLEANUP_I20_I]] ], [ [[__R_0_I16_I]], [[WHILE_COND_I14_I]] ]
 // CHECK-NEXT:    ret i64 [[RETVAL_0_I]]
 //
+// AMDGCNSPIRV-LABEL: @test___make_mantissa(
+// AMDGCNSPIRV-NEXT:  entry:
+// AMDGCNSPIRV-NEXT:    [[TMP0:%.*]] = load i8, ptr addrspace(4) [[P:%.*]], align 1, !tbaa [[TBAA5]]
+// AMDGCNSPIRV-NEXT:    [[CMP_I:%.*]] = icmp eq i8 [[TMP0]], 48
+// AMDGCNSPIRV-NEXT:    br i1 [[CMP_I]], label [[IF_THEN_I:%.*]], label [[WHILE_COND_I14_I:%.*]]
+// AMDGCNSPIRV:       if.then.i:
+// AMDGCNSPIRV-NEXT:    [[INCDEC_PTR_I:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(4) [[P]], i64 1
+// AMDGCNSPIRV-NEXT:    [[TMP1:%.*]] = load i8, ptr addrspace(4) [[INCDEC_PTR_I]], align 1, !tbaa [[TBAA5]]
+// AMDGCNSPIRV-NEXT:    switch i8 [[TMP1]], label [[WHILE_COND_I_I:%.*]] [
+// AMDGCNSPIRV-NEXT:      i8 120, label [[WHILE_COND_I28_I_PREHEADER:%.*]]
+// AMDGCNSPIRV-NEXT:      i8 88, label [[WHILE_COND_I28_I_PREHEADER]]
+// AMDGCNSPIRV-NEXT:    ]
+// AMDGCNSPIRV:       while.cond.i28.i.preheader:
+// AMDGCNSPIRV-NEXT:    br label [[WHILE_COND_I28_I:%.*]]
+// AMDGCNSPIRV:       while.con...
[truncated]

llvmbot · 2025-02-28T21:02:11Z

@llvm/pr-subscribers-backend-x86

Author: Alex Voicu (AlexVlx)

Changes

This reapplies #128360, the only change being that the modified tests also checks for the availability of the SPIRV target.

Patch is 183.47 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/129306.diff

3 Files Affected:

(modified) clang/lib/Headers/__clang_hip_libdevice_declares.h (+12-20)
(modified) clang/lib/Headers/__clang_hip_math.h (+12-16)
(modified) clang/test/Headers/__clang_hip_math.hip (+1656)

diff --git a/clang/lib/Headers/__clang_hip_libdevice_declares.h b/clang/lib/Headers/__clang_hip_libdevice_declares.h
index f15198b3d9f93..fa8d918248dd0 100644
--- a/clang/lib/Headers/__clang_hip_libdevice_declares.h
+++ b/clang/lib/Headers/__clang_hip_libdevice_declares.h
@@ -14,6 +14,8 @@
 #include "hip/hip_version.h"
 #endif // __has_include("hip/hip_version.h")
 
+#define __PRIVATE_AS __attribute__((opencl_private))
+
 #ifdef __cplusplus
 extern "C" {
 #endif
@@ -55,8 +57,7 @@ __device__ __attribute__((const)) float __ocml_fmax_f32(float, float);
 __device__ __attribute__((const)) float __ocml_fmin_f32(float, float);
 __device__ __attribute__((const)) __device__ float __ocml_fmod_f32(float,
                                                                    float);
-__device__ float __ocml_frexp_f32(float,
-                                  __attribute__((address_space(5))) int *);
+__device__ float __ocml_frexp_f32(float, __PRIVATE_AS int *);
 __device__ __attribute__((const)) float __ocml_hypot_f32(float, float);
 __device__ __attribute__((const)) int __ocml_ilogb_f32(float);
 __device__ __attribute__((const)) int __ocml_isfinite_f32(float);
@@ -74,8 +75,7 @@ __device__ __attribute__((pure)) float __ocml_native_log2_f32(float);
 __device__ __attribute__((const)) float __ocml_logb_f32(float);
 __device__ __attribute__((pure)) float __ocml_log_f32(float);
 __device__ __attribute__((pure)) float __ocml_native_log_f32(float);
-__device__ float __ocml_modf_f32(float,
-                                 __attribute__((address_space(5))) float *);
+__device__ float __ocml_modf_f32(float, __PRIVATE_AS float *);
 __device__ __attribute__((const)) float __ocml_nearbyint_f32(float);
 __device__ __attribute__((const)) float __ocml_nextafter_f32(float, float);
 __device__ __attribute__((const)) float __ocml_len3_f32(float, float, float);
@@ -87,8 +87,7 @@ __device__ __attribute__((pure)) float __ocml_pow_f32(float, float);
 __device__ __attribute__((pure)) float __ocml_pown_f32(float, int);
 __device__ __attribute__((pure)) float __ocml_rcbrt_f32(float);
 __device__ __attribute__((const)) float __ocml_remainder_f32(float, float);
-__device__ float __ocml_remquo_f32(float, float,
-                                   __attribute__((address_space(5))) int *);
+__device__ float __ocml_remquo_f32(float, float, __PRIVATE_AS int *);
 __device__ __attribute__((const)) float __ocml_rhypot_f32(float, float);
 __device__ __attribute__((const)) float __ocml_rint_f32(float);
 __device__ __attribute__((const)) float __ocml_rlen3_f32(float, float, float);
@@ -99,10 +98,8 @@ __device__ __attribute__((pure)) float __ocml_rsqrt_f32(float);
 __device__ __attribute__((const)) float __ocml_scalb_f32(float, float);
 __device__ __attribute__((const)) float __ocml_scalbn_f32(float, int);
 __device__ __attribute__((const)) int __ocml_signbit_f32(float);
-__device__ float __ocml_sincos_f32(float,
-                                   __attribute__((address_space(5))) float *);
-__device__ float __ocml_sincospi_f32(float,
-                                     __attribute__((address_space(5))) float *);
+__device__ float __ocml_sincos_f32(float, __PRIVATE_AS float *);
+__device__ float __ocml_sincospi_f32(float, __PRIVATE_AS float *);
 __device__ float __ocml_sin_f32(float);
 __device__ float __ocml_native_sin_f32(float);
 __device__ __attribute__((pure)) float __ocml_sinh_f32(float);
@@ -176,8 +173,7 @@ __device__ __attribute__((const)) double __ocml_fma_f64(double, double, double);
 __device__ __attribute__((const)) double __ocml_fmax_f64(double, double);
 __device__ __attribute__((const)) double __ocml_fmin_f64(double, double);
 __device__ __attribute__((const)) double __ocml_fmod_f64(double, double);
-__device__ double __ocml_frexp_f64(double,
-                                   __attribute__((address_space(5))) int *);
+__device__ double __ocml_frexp_f64(double, __PRIVATE_AS int *);
 __device__ __attribute__((const)) double __ocml_hypot_f64(double, double);
 __device__ __attribute__((const)) int __ocml_ilogb_f64(double);
 __device__ __attribute__((const)) int __ocml_isfinite_f64(double);
@@ -192,8 +188,7 @@ __device__ __attribute__((pure)) double __ocml_log1p_f64(double);
 __device__ __attribute__((pure)) double __ocml_log2_f64(double);
 __device__ __attribute__((const)) double __ocml_logb_f64(double);
 __device__ __attribute__((pure)) double __ocml_log_f64(double);
-__device__ double __ocml_modf_f64(double,
-                                  __attribute__((address_space(5))) double *);
+__device__ double __ocml_modf_f64(double, __PRIVATE_AS double *);
 __device__ __attribute__((const)) double __ocml_nearbyint_f64(double);
 __device__ __attribute__((const)) double __ocml_nextafter_f64(double, double);
 __device__ __attribute__((const)) double __ocml_len3_f64(double, double,
@@ -206,8 +201,7 @@ __device__ __attribute__((pure)) double __ocml_pow_f64(double, double);
 __device__ __attribute__((pure)) double __ocml_pown_f64(double, int);
 __device__ __attribute__((pure)) double __ocml_rcbrt_f64(double);
 __device__ __attribute__((const)) double __ocml_remainder_f64(double, double);
-__device__ double __ocml_remquo_f64(double, double,
-                                    __attribute__((address_space(5))) int *);
+__device__ double __ocml_remquo_f64(double, double, __PRIVATE_AS int *);
 __device__ __attribute__((const)) double __ocml_rhypot_f64(double, double);
 __device__ __attribute__((const)) double __ocml_rint_f64(double);
 __device__ __attribute__((const)) double __ocml_rlen3_f64(double, double,
@@ -219,10 +213,8 @@ __device__ __attribute__((pure)) double __ocml_rsqrt_f64(double);
 __device__ __attribute__((const)) double __ocml_scalb_f64(double, double);
 __device__ __attribute__((const)) double __ocml_scalbn_f64(double, int);
 __device__ __attribute__((const)) int __ocml_signbit_f64(double);
-__device__ double __ocml_sincos_f64(double,
-                                    __attribute__((address_space(5))) double *);
-__device__ double
-__ocml_sincospi_f64(double, __attribute__((address_space(5))) double *);
+__device__ double __ocml_sincos_f64(double, __PRIVATE_AS double *);
+__device__ double __ocml_sincospi_f64(double, __PRIVATE_AS double *);
 __device__ double __ocml_sin_f64(double);
 __device__ __attribute__((pure)) double __ocml_sinh_f64(double);
 __device__ double __ocml_sinpi_f64(double);
diff --git a/clang/lib/Headers/__clang_hip_math.h b/clang/lib/Headers/__clang_hip_math.h
index 8468751d9de26..bf8517bc3a507 100644
--- a/clang/lib/Headers/__clang_hip_math.h
+++ b/clang/lib/Headers/__clang_hip_math.h
@@ -33,6 +33,9 @@
 #define __DEVICE__ static __device__ inline __attribute__((always_inline))
 #endif
 
+#pragma push_macro("__PRIVATE_AS")
+
+#define __PRIVATE_AS __attribute__((opencl_private))
 // Device library provides fast low precision and slow full-recision
 // implementations for some functions. Which one gets selected depends on
 // __CLANG_GPU_APPROX_TRANSCENDENTALS__ which gets defined by clang if
@@ -512,8 +515,7 @@ float modff(float __x, float *__iptr) {
 #ifdef __OPENMP_AMDGCN__
 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
 #endif
-  float __r =
-      __ocml_modf_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
+  float __r = __ocml_modf_f32(__x, (__PRIVATE_AS float *)&__tmp);
   *__iptr = __tmp;
   return __r;
 }
@@ -595,8 +597,7 @@ float remquof(float __x, float __y, int *__quo) {
 #ifdef __OPENMP_AMDGCN__
 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
 #endif
-  float __r = __ocml_remquo_f32(
-      __x, __y, (__attribute__((address_space(5))) int *)&__tmp);
+  float __r = __ocml_remquo_f32(__x, __y, (__PRIVATE_AS int *)&__tmp);
   *__quo = __tmp;
 
   return __r;
@@ -657,8 +658,7 @@ void sincosf(float __x, float *__sinptr, float *__cosptr) {
 #ifdef __CLANG_CUDA_APPROX_TRANSCENDENTALS__
   __sincosf(__x, __sinptr, __cosptr);
 #else
-  *__sinptr =
-      __ocml_sincos_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
+  *__sinptr = __ocml_sincos_f32(__x, (__PRIVATE_AS float *)&__tmp);
   *__cosptr = __tmp;
 #endif
 }
@@ -669,8 +669,7 @@ void sincospif(float __x, float *__sinptr, float *__cosptr) {
 #ifdef __OPENMP_AMDGCN__
 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
 #endif
-  *__sinptr = __ocml_sincospi_f32(
-      __x, (__attribute__((address_space(5))) float *)&__tmp);
+  *__sinptr = __ocml_sincospi_f32(__x, (__PRIVATE_AS float *)&__tmp);
   *__cosptr = __tmp;
 }
 
@@ -913,8 +912,7 @@ double modf(double __x, double *__iptr) {
 #ifdef __OPENMP_AMDGCN__
 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
 #endif
-  double __r =
-      __ocml_modf_f64(__x, (__attribute__((address_space(5))) double *)&__tmp);
+  double __r = __ocml_modf_f64(__x, (__PRIVATE_AS double *)&__tmp);
   *__iptr = __tmp;
 
   return __r;
@@ -1004,8 +1002,7 @@ double remquo(double __x, double __y, int *__quo) {
 #ifdef __OPENMP_AMDGCN__
 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
 #endif
-  double __r = __ocml_remquo_f64(
-      __x, __y, (__attribute__((address_space(5))) int *)&__tmp);
+  double __r = __ocml_remquo_f64(__x, __y, (__PRIVATE_AS int *)&__tmp);
   *__quo = __tmp;
 
   return __r;
@@ -1065,8 +1062,7 @@ void sincos(double __x, double *__sinptr, double *__cosptr) {
 #ifdef __OPENMP_AMDGCN__
 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
 #endif
-  *__sinptr = __ocml_sincos_f64(
-      __x, (__attribute__((address_space(5))) double *)&__tmp);
+  *__sinptr = __ocml_sincos_f64(__x, (__PRIVATE_AS double *)&__tmp);
   *__cosptr = __tmp;
 }
 
@@ -1076,8 +1072,7 @@ void sincospi(double __x, double *__sinptr, double *__cosptr) {
 #ifdef __OPENMP_AMDGCN__
 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
 #endif
-  *__sinptr = __ocml_sincospi_f64(
-      __x, (__attribute__((address_space(5))) double *)&__tmp);
+  *__sinptr = __ocml_sincospi_f64(__x, (__PRIVATE_AS double *)&__tmp);
   *__cosptr = __tmp;
 }
 
@@ -1322,6 +1317,7 @@ __host__ inline static int max(int __arg1, int __arg2) {
 #endif
 
 #pragma pop_macro("__DEVICE__")
+#pragma pop_macro("__PRIVATE_AS")
 #pragma pop_macro("__RETURN_TYPE")
 #pragma pop_macro("__FAST_OR_SLOW")
 
diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip
index e4254d1e64bec..d448ab134ca4d 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -1,5 +1,6 @@
 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
 // REQUIRES: amdgpu-registered-target
+// REQUIRES: spirv-registered-target
 
 // Test without OCML_BASIC_ROUNDED_OPERATIONS
 // RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
@@ -26,6 +27,14 @@
 // RUN:   -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -O1 -fgpu-approx-transcendentals -o - \
 // RUN:   -D__HIPCC_RTC__ | FileCheck -check-prefixes=CHECK,APPROX %s
 
+// Check that we use the AMDGCNSPIRV address space map
+// RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
+// RUN:   -internal-isystem %S/../../lib/Headers/cuda_wrappers \
+// RUN:   -internal-isystem %S/Inputs/include \
+// RUN:   -triple spirv64-amd-amdhsa -aux-triple x86_64-unknown-unknown \
+// RUN:   -emit-llvm %s -fcuda-is-device -O1 -o - \
+// RUN:   -D__HIPCC_RTC__ | FileCheck -check-prefixes=AMDGCNSPIRV %s
+
 #define BOOL_TYPE int
 typedef unsigned long long uint64_t;
 
@@ -57,6 +66,30 @@ typedef unsigned long long uint64_t;
 // CHECK-NEXT:    [[RETVAL_2_I:%.*]] = phi i64 [ 0, [[CLEANUP_I]] ], [ [[__R_0_I]], [[WHILE_COND_I]] ]
 // CHECK-NEXT:    ret i64 [[RETVAL_2_I]]
 //
+// AMDGCNSPIRV-LABEL: @test___make_mantissa_base8(
+// AMDGCNSPIRV-NEXT:  entry:
+// AMDGCNSPIRV-NEXT:    br label [[WHILE_COND_I:%.*]]
+// AMDGCNSPIRV:       while.cond.i:
+// AMDGCNSPIRV-NEXT:    [[__TAGP_ADDR_0_I:%.*]] = phi ptr addrspace(4) [ [[P:%.*]], [[ENTRY:%.*]] ], [ [[__TAGP_ADDR_1_I:%.*]], [[WHILE_BODY_I:%.*]] ]
+// AMDGCNSPIRV-NEXT:    [[__R_0_I:%.*]] = phi i64 [ 0, [[ENTRY]] ], [ [[__R_1_I:%.*]], [[WHILE_BODY_I]] ]
+// AMDGCNSPIRV-NEXT:    [[TMP0:%.*]] = load i8, ptr addrspace(4) [[__TAGP_ADDR_0_I]], align 1, !tbaa [[TBAA5:![0-9]+]]
+// AMDGCNSPIRV-NEXT:    [[CMP_NOT_I:%.*]] = icmp eq i8 [[TMP0]], 0
+// AMDGCNSPIRV-NEXT:    br i1 [[CMP_NOT_I]], label [[_ZL21__MAKE_MANTISSA_BASE8PKC_EXIT:%.*]], label [[WHILE_BODY_I]]
+// AMDGCNSPIRV:       while.body.i:
+// AMDGCNSPIRV-NEXT:    [[TMP1:%.*]] = and i8 [[TMP0]], -8
+// AMDGCNSPIRV-NEXT:    [[OR_COND_I:%.*]] = icmp eq i8 [[TMP1]], 48
+// AMDGCNSPIRV-NEXT:    [[MUL_I:%.*]] = shl i64 [[__R_0_I]], 3
+// AMDGCNSPIRV-NEXT:    [[CONV5_I:%.*]] = zext nneg i8 [[TMP0]] to i64
+// AMDGCNSPIRV-NEXT:    [[ADD_I:%.*]] = add i64 [[MUL_I]], -48
+// AMDGCNSPIRV-NEXT:    [[SUB_I:%.*]] = add i64 [[ADD_I]], [[CONV5_I]]
+// AMDGCNSPIRV-NEXT:    [[__TAGP_ADDR_1_I_IDX:%.*]] = zext i1 [[OR_COND_I]] to i64
+// AMDGCNSPIRV-NEXT:    [[__TAGP_ADDR_1_I]] = getelementptr inbounds nuw i8, ptr addrspace(4) [[__TAGP_ADDR_0_I]], i64 [[__TAGP_ADDR_1_I_IDX]]
+// AMDGCNSPIRV-NEXT:    [[__R_1_I]] = select i1 [[OR_COND_I]], i64 [[SUB_I]], i64 [[__R_0_I]]
+// AMDGCNSPIRV-NEXT:    br i1 [[OR_COND_I]], label [[WHILE_COND_I]], label [[_ZL21__MAKE_MANTISSA_BASE8PKC_EXIT]], !llvm.loop [[LOOP8:![0-9]+]]
+// AMDGCNSPIRV:       _ZL21__make_mantissa_base8PKc.exit:
+// AMDGCNSPIRV-NEXT:    [[RETVAL_2_I:%.*]] = phi i64 [ 0, [[WHILE_BODY_I]] ], [ [[__R_0_I]], [[WHILE_COND_I]] ]
+// AMDGCNSPIRV-NEXT:    ret i64 [[RETVAL_2_I]]
+//
 extern "C" __device__ uint64_t test___make_mantissa_base8(const char *p) {
   return __make_mantissa_base8(p);
 }
@@ -89,6 +122,30 @@ extern "C" __device__ uint64_t test___make_mantissa_base8(const char *p) {
 // CHECK-NEXT:    [[RETVAL_2_I:%.*]] = phi i64 [ 0, [[CLEANUP_I]] ], [ [[__R_0_I]], [[WHILE_COND_I]] ]
 // CHECK-NEXT:    ret i64 [[RETVAL_2_I]]
 //
+// AMDGCNSPIRV-LABEL: @test___make_mantissa_base10(
+// AMDGCNSPIRV-NEXT:  entry:
+// AMDGCNSPIRV-NEXT:    br label [[WHILE_COND_I:%.*]]
+// AMDGCNSPIRV:       while.cond.i:
+// AMDGCNSPIRV-NEXT:    [[__TAGP_ADDR_0_I:%.*]] = phi ptr addrspace(4) [ [[P:%.*]], [[ENTRY:%.*]] ], [ [[__TAGP_ADDR_1_I:%.*]], [[WHILE_BODY_I:%.*]] ]
+// AMDGCNSPIRV-NEXT:    [[__R_0_I:%.*]] = phi i64 [ 0, [[ENTRY]] ], [ [[__R_1_I:%.*]], [[WHILE_BODY_I]] ]
+// AMDGCNSPIRV-NEXT:    [[TMP0:%.*]] = load i8, ptr addrspace(4) [[__TAGP_ADDR_0_I]], align 1, !tbaa [[TBAA5]]
+// AMDGCNSPIRV-NEXT:    [[CMP_NOT_I:%.*]] = icmp eq i8 [[TMP0]], 0
+// AMDGCNSPIRV-NEXT:    br i1 [[CMP_NOT_I]], label [[_ZL22__MAKE_MANTISSA_BASE10PKC_EXIT:%.*]], label [[WHILE_BODY_I]]
+// AMDGCNSPIRV:       while.body.i:
+// AMDGCNSPIRV-NEXT:    [[TMP1:%.*]] = add i8 [[TMP0]], -48
+// AMDGCNSPIRV-NEXT:    [[OR_COND_I:%.*]] = icmp ult i8 [[TMP1]], 10
+// AMDGCNSPIRV-NEXT:    [[MUL_I:%.*]] = mul i64 [[__R_0_I]], 10
+// AMDGCNSPIRV-NEXT:    [[CONV5_I:%.*]] = zext nneg i8 [[TMP0]] to i64
+// AMDGCNSPIRV-NEXT:    [[ADD_I:%.*]] = add i64 [[MUL_I]], -48
+// AMDGCNSPIRV-NEXT:    [[SUB_I:%.*]] = add i64 [[ADD_I]], [[CONV5_I]]
+// AMDGCNSPIRV-NEXT:    [[__TAGP_ADDR_1_I_IDX:%.*]] = zext i1 [[OR_COND_I]] to i64
+// AMDGCNSPIRV-NEXT:    [[__TAGP_ADDR_1_I]] = getelementptr inbounds nuw i8, ptr addrspace(4) [[__TAGP_ADDR_0_I]], i64 [[__TAGP_ADDR_1_I_IDX]]
+// AMDGCNSPIRV-NEXT:    [[__R_1_I]] = select i1 [[OR_COND_I]], i64 [[SUB_I]], i64 [[__R_0_I]]
+// AMDGCNSPIRV-NEXT:    br i1 [[OR_COND_I]], label [[WHILE_COND_I]], label [[_ZL22__MAKE_MANTISSA_BASE10PKC_EXIT]], !llvm.loop [[LOOP11:![0-9]+]]
+// AMDGCNSPIRV:       _ZL22__make_mantissa_base10PKc.exit:
+// AMDGCNSPIRV-NEXT:    [[RETVAL_2_I:%.*]] = phi i64 [ 0, [[WHILE_BODY_I]] ], [ [[__R_0_I]], [[WHILE_COND_I]] ]
+// AMDGCNSPIRV-NEXT:    ret i64 [[RETVAL_2_I]]
+//
 extern "C" __device__ uint64_t test___make_mantissa_base10(const char *p) {
   return __make_mantissa_base10(p);
 }
@@ -131,6 +188,44 @@ extern "C" __device__ uint64_t test___make_mantissa_base10(const char *p) {
 // CHECK-NEXT:    [[RETVAL_2_I:%.*]] = phi i64 [ 0, [[CLEANUP_I]] ], [ [[__R_0_I]], [[WHILE_COND_I]] ]
 // CHECK-NEXT:    ret i64 [[RETVAL_2_I]]
 //
+// AMDGCNSPIRV-LABEL: @test___make_mantissa_base16(
+// AMDGCNSPIRV-NEXT:  entry:
+// AMDGCNSPIRV-NEXT:    br label [[WHILE_COND_I:%.*]]
+// AMDGCNSPIRV:       while.cond.i:
+// AMDGCNSPIRV-NEXT:    [[__TAGP_ADDR_0_I:%.*]] = phi ptr addrspace(4) [ [[P:%.*]], [[ENTRY:%.*]] ], [ [[__TAGP_ADDR_1_I:%.*]], [[CLEANUP_I:%.*]] ]
+// AMDGCNSPIRV-NEXT:    [[__R_0_I:%.*]] = phi i64 [ 0, [[ENTRY]] ], [ [[__R_2_I:%.*]], [[CLEANUP_I]] ]
+// AMDGCNSPIRV-NEXT:    [[TMP0:%.*]] = load i8, ptr addrspace(4) [[__TAGP_ADDR_0_I]], align 1, !tbaa [[TBAA5]]
+// AMDGCNSPIRV-NEXT:    [[CMP_NOT_I:%.*]] = icmp eq i8 [[TMP0]], 0
+// AMDGCNSPIRV-NEXT:    br i1 [[CMP_NOT_I]], label [[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// AMDGCNSPIRV:       while.body.i:
+// AMDGCNSPIRV-NEXT:    [[TMP1:%.*]] = add i8 [[TMP0]], -48
+// AMDGCNSPIRV-NEXT:    [[OR_COND_I:%.*]] = icmp ult i8 [[TMP1]], 10
+// AMDGCNSPIRV-NEXT:    br i1 [[OR_COND_I]], label [[IF_END31_I:%.*]], label [[IF_ELSE_I:%.*]]
+// AMDGCNSPIRV:       if.else.i:
+// AMDGCNSPIRV-NEXT:    [[TMP2:%.*]] = add i8 [[TMP0]], -97
+// AMDGCNSPIRV-NEXT:    [[OR_COND33_I:%.*]] = icmp ult i8 [[TMP2]], 6
+// AMDGCNSPIRV-NEXT:    br i1 [[OR_COND33_I]], label [[IF_END31_I]], label [[IF_ELSE17_I:%.*]]
+// AMDGCNSPIRV:       if.else17.i:
+// AMDGCNSPIRV-NEXT:    [[TMP3:%.*]] = add i8 [[TMP0]], -65
+// AMDGCNSPIRV-NEXT:    [[OR_COND34_I:%.*]] = icmp ult i8 [[TMP3]], 6
+// AMDGCNSPIRV-NEXT:    br i1 [[OR_COND34_I]], label [[IF_END31_I]], label [[CLEANUP_I]]
+// AMDGCNSPIRV:       if.end31.i:
+// AMDGCNSPIRV-NEXT:    [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I]] ], [ -87, [[IF_ELSE_I]] ], [ -55, [[IF_ELSE17_I]] ]
+// AMDGCNSPIRV-NEXT:    [[MUL24_I:%.*]] = shl i64 [[__R_0_I]], 4
+// AMDGCNSPIRV-NEXT:    [[CONV25_I:%.*]] = zext nneg i8 [[TMP0]] to i64
+// AMDGCNSPIRV-NEXT:    [[ADD26_I:%.*]] = add i64 [[MUL24_I]], [[DOTSINK]]
+// AMDGCNSPIRV-NEXT:    [[ADD28_I:%.*]] = add i64 [[ADD26_I]], [[CONV25_I]]
+// AMDGCNSPIRV-NEXT:    [[INCDEC_PTR_I:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(4) [[__TAGP_ADDR_0_I]], i64 1
+// AMDGCNSPIRV-NEXT:    br label [[CLEANUP_I]]
+// AMDGCNSPIRV:       cleanup.i:
+// AMDGCNSPIRV-NEXT:    [[__TAGP_ADDR_1_I]] = phi ptr addrspace(4) [ [[INCDEC_PTR_I]], [[IF_END31_I]] ], [ [[__TAGP_ADDR_0_I]], [[IF_ELSE17_I]] ]
+// AMDGCNSPIRV-NEXT:    [[__R_2_I]] = phi i64 [ [[ADD28_I]], [[IF_END31_I]] ], [ [[__R_0_I]], [[IF_ELSE17_I]] ]
+// AMDGCNSPIRV-NEXT:    [[COND_I:%.*]] = phi i1 [ true, [[IF_END31_I]] ], [ false, [[IF_ELSE17_I]] ]
+// AMDGCNSPIRV-NEXT:    br i1 [[COND_I]], label [[WHILE_COND_I]], label [[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT]], !llvm.loop [[LOOP12:![0-9]+]]
+// AMDGCNSPIRV:       _ZL22__make_mantissa_base16PKc.exit:
+// AMDGCNSPIRV-NEXT:    [[RETVAL_2_I:%.*]] = phi i64 [ 0, [[CLEANUP_I]] ], [ [[__R_0_I]], [[WHILE_COND_I]] ]
+// AMDGCNSPIRV-NEXT:    ret i64 [[RETVAL_2_I]]
+//
 extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) {
   return __make_mantissa_base16(p);
 }
@@ -226,6 +321,89 @@ extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) {
 // CHECK-NEXT:    [[RETVAL_0_I:%.*]] = phi i64 [ 0, [[CLEANUP_I_I]] ], [ [[__R_0_I_I]], [[WHILE_COND_I_I]] ], [ 0, [[CLEANUP_I36_I]] ], [ [[__R_0_I32_I]], [[WHILE_COND_I30_I]] ], [ 0, [[CLEANUP_I20_I]] ], [ [[__R_0_I16_I]], [[WHILE_COND_I14_I]] ]
 // CHECK-NEXT:    ret i64 [[RETVAL_0_I]]
 //
+// AMDGCNSPIRV-LABEL: @test___make_mantissa(
+// AMDGCNSPIRV-NEXT:  entry:
+// AMDGCNSPIRV-NEXT:    [[TMP0:%.*]] = load i8, ptr addrspace(4) [[P:%.*]], align 1, !tbaa [[TBAA5]]
+// AMDGCNSPIRV-NEXT:    [[CMP_I:%.*]] = icmp eq i8 [[TMP0]], 48
+// AMDGCNSPIRV-NEXT:    br i1 [[CMP_I]], label [[IF_THEN_I:%.*]], label [[WHILE_COND_I14_I:%.*]]
+// AMDGCNSPIRV:       if.then.i:
+// AMDGCNSPIRV-NEXT:    [[INCDEC_PTR_I:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(4) [[P]], i64 1
+// AMDGCNSPIRV-NEXT:    [[TMP1:%.*]] = load i8, ptr addrspace(4) [[INCDEC_PTR_I]], align 1, !tbaa [[TBAA5]]
+// AMDGCNSPIRV-NEXT:    switch i8 [[TMP1]], label [[WHILE_COND_I_I:%.*]] [
+// AMDGCNSPIRV-NEXT:      i8 120, label [[WHILE_COND_I28_I_PREHEADER:%.*]]
+// AMDGCNSPIRV-NEXT:      i8 88, label [[WHILE_COND_I28_I_PREHEADER]]
+// AMDGCNSPIRV-NEXT:    ]
+// AMDGCNSPIRV:       while.cond.i28.i.preheader:
+// AMDGCNSPIRV-NEXT:    br label [[WHILE_COND_I28_I:%.*]]
+// AMDGCNSPIRV:       while.con...
[truncated]

llvm-ci · 2025-03-01T00:17:05Z

LLVM Buildbot has detected a new failure on builder llvm-x86_64-debian-dylib running on gribozavr4 while building clang at step 7 "test-build-unified-tree-check-llvm".

Full details are available at: https://lab.llvm.org/buildbot/#/builders/60/builds/20773

Here is the relevant piece of the build log for the reference

Step 7 (test-build-unified-tree-check-llvm) failure: test (failure)
******************** TEST 'LLVM :: tools/llvm-exegesis/RISCV/rvv/reduction.test' FAILED ********************
Exit Code: 1

Command Output (stderr):
--
RUN: at line 1: /b/1/llvm-x86_64-debian-dylib/build/bin/llvm-exegesis -mtriple=riscv64 -mcpu=sifive-p670 -benchmark-phase=assemble-measured-code --mode=latency --opcode-name=PseudoVWREDSUMU_VS_M8_E32 --min-instructions=100 |     /b/1/llvm-x86_64-debian-dylib/build/bin/FileCheck /b/1/llvm-x86_64-debian-dylib/llvm-project/llvm/test/tools/llvm-exegesis/RISCV/rvv/reduction.test
+ /b/1/llvm-x86_64-debian-dylib/build/bin/FileCheck /b/1/llvm-x86_64-debian-dylib/llvm-project/llvm/test/tools/llvm-exegesis/RISCV/rvv/reduction.test
+ /b/1/llvm-x86_64-debian-dylib/build/bin/llvm-exegesis -mtriple=riscv64 -mcpu=sifive-p670 -benchmark-phase=assemble-measured-code --mode=latency --opcode-name=PseudoVWREDSUMU_VS_M8_E32 --min-instructions=100
/b/1/llvm-x86_64-debian-dylib/llvm-project/llvm/test/tools/llvm-exegesis/RISCV/rvv/reduction.test:7:14: error: CHECK-NOT: excluded string found in input
# CHECK-NOT: V[[REG:[0-9]+]] V[[REG]] V{{[0-9]+}}M8 V[[REG]]
             ^
<stdin>:5:31: note: found here
 - 'PseudoVWREDSUMU_VS_M8_E32 V7 V7 V0M8 V7 i_0xffffffffffffffff i_0x5 i_0x0'
                              ^~~~~~~~~~~~~
<stdin>:5:32: note: captured var "REG"
 - 'PseudoVWREDSUMU_VS_M8_E32 V7 V7 V0M8 V7 i_0xffffffffffffffff i_0x5 i_0x0'
                               ^

Input file: <stdin>
Check file: /b/1/llvm-x86_64-debian-dylib/llvm-project/llvm/test/tools/llvm-exegesis/RISCV/rvv/reduction.test

-dump-input=help explains the following input dump.

Input was:
<<<<<<
         1: --- 
         2: mode: latency 
         3: key: 
         4:  instructions: 
         5:  - 'PseudoVWREDSUMU_VS_M8_E32 V7 V7 V0M8 V7 i_0xffffffffffffffff i_0x5 i_0x0' 
not:7'0                                   !~~~~~~~~~~~~                                    error: no match expected
not:7'1                                    !                                               captured var "REG"
         6:  config: 'vtype = {AVL: VLMAX, SEW: e32, Policy: tu/mu}' 
         7:  register_initial_values: 
         8:  - 'V7=0x0' 
         9:  - 'V0M8=0x0' 
        10: cpu_name: sifive-p670 
         .
         .
         .
>>>>>>

--

********************

…lvm#128360" (llvm#129306) This reapplies llvm#128360, the only change being that the modified tests also checks for the availability of the SPIRV target.

llvm#2066) …llvm#128360" (llvm#129306) This reapplies llvm#128360, the only change being that the modified tests also checks for the availability of the SPIRV target. Co-authored-by: Alex Voicu <[email protected]>

…lvm#128360" (llvm#129306) This reapplies llvm#128360, the only change being that the modified tests also checks for the availability of the SPIRV target.

AlexVlx added 8 commits February 22, 2025 17:34

"Generalise" math function signatures taking an explicit AS.

184e0f3

Merge branch 'main' of https://github.com/llvm/llvm-project into rocd…

8ade6ad

…l_is_bad_for_spirv

Apply suggestion.

b53179a

Add missing paired attribute.

4b012cd

Merge branch 'main' of https://github.com/llvm/llvm-project into rocd…

0020b73

…l_is_bad_for_spirv

Ensure SPIR-V target is available.

36e550c

Merge branch 'main' of https://github.com/llvm/llvm-project into rocd…

e905031

…l_is_bad_for_spirv

[clang][HIP] Make some math not not work with AMDGCN SPIR-V (llvm#128360

c719e91

) Do not hardcode `address_space(5)` (`private`) in the ROCDL interface, as that breaks SPIRV generation (the latter uses 0). Add test. In the long run we should stop using ROCDL inline.

llvmbot added clang Clang issues not falling into any other category backend:X86 clang:headers Headers provided by Clang, e.g. for intrinsics labels Feb 28, 2025

AlexVlx merged commit 21a0500 into llvm:main Feb 28, 2025
15 checks passed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

Reapply "[clang][HIP] Make some math not not work with AMDGCN SPIR-V #128360" #129306

Reapply "[clang][HIP] Make some math not not work with AMDGCN SPIR-V #128360" #129306

Uh oh!

AlexVlx commented Feb 28, 2025

Uh oh!

llvmbot commented Feb 28, 2025

Uh oh!

llvmbot commented Feb 28, 2025

Uh oh!

Uh oh!

llvm-ci commented Mar 1, 2025

Uh oh!

Uh oh!

Reapply "[clang][HIP] Make some math not not work with AMDGCN SPIR-V #128360" #129306

Reapply "[clang][HIP] Make some math not not work with AMDGCN SPIR-V #128360" #129306

Uh oh!

Conversation

AlexVlx commented Feb 28, 2025

Uh oh!

llvmbot commented Feb 28, 2025

Uh oh!

llvmbot commented Feb 28, 2025

Uh oh!

Uh oh!

llvm-ci commented Mar 1, 2025

Uh oh!

Uh oh!