https://github.com/rampitec updated https://github.com/llvm/llvm-project/pull/215894
>From 155fe33be529d49767fb75d7d82ef523589f23fa Mon Sep 17 00:00:00 2001 From: Stanislav Mekhanoshin <[email protected]> Date: Wed, 12 Aug 2026 13:38:22 -0700 Subject: [PATCH] [AMDGPU] Allow constant folding of bfloat --- clang/test/CodeGen/builtins-nvptx.c | 10 +-- llvm/lib/Analysis/ConstantFolding.cpp | 6 +- .../CodeGen/AMDGPU/llvm.amdgcn.cos.bf16.ll | 47 ++++++++++--- .../CodeGen/AMDGPU/llvm.amdgcn.sin.bf16.ll | 47 ++++++++++--- llvm/test/CodeGen/AMDGPU/llvm.cos.bf16.ll | 47 ++++++++++--- llvm/test/CodeGen/AMDGPU/llvm.sin.bf16.ll | 47 ++++++++++--- .../InstCombine/AMDGPU/amdgcn-intrinsics.ll | 69 ++++++++++++------- .../InstCombine/cos-sin-intrinsic.ll | 36 ++++++++++ 8 files changed, 243 insertions(+), 66 deletions(-) diff --git a/clang/test/CodeGen/builtins-nvptx.c b/clang/test/CodeGen/builtins-nvptx.c index 87be7b46aad8e..0eb9eadb251d6 100644 --- a/clang/test/CodeGen/builtins-nvptx.c +++ b/clang/test/CodeGen/builtins-nvptx.c @@ -1415,13 +1415,13 @@ __device__ void nvvm_cvt_sm100a_sm103a() { #define NANBF16X2 {NANBF16, NANBF16} // CHECK-LABEL: nvvm_abs_neg_bf16_bf16x2_sm80 -__device__ void nvvm_abs_neg_bf16_bf16x2_sm80() { +__device__ void nvvm_abs_neg_bf16_bf16x2_sm80(__bf16 a) { #if __CUDA_ARCH__ >= 800 - // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.fabs.bf16(bfloat 1.000980e-01) - __nvvm_abs_bf16(BF16); - // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.fabs.v2bf16(<2 x bfloat> splat (bfloat 1.000980e-01)) - __nvvm_abs_bf16x2(BF16X2); + // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.fabs.bf16(bfloat + __nvvm_abs_bf16(a); + // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.fabs.v2bf16(<2 x bfloat> + __nvvm_abs_bf16x2({a, a}); // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.neg.bf16(bfloat 1.000980e-01) __nvvm_neg_bf16(BF16); diff --git a/llvm/lib/Analysis/ConstantFolding.cpp b/llvm/lib/Analysis/ConstantFolding.cpp index 3c110b2138e83..59328d0aefed7 100644 --- a/llvm/lib/Analysis/ConstantFolding.cpp +++ b/llvm/lib/Analysis/ConstantFolding.cpp @@ -2209,7 +2209,7 @@ bool llvm::canConstantFoldCallTo(const CallBase *Call, const Function *F) { namespace { Constant *GetConstantFoldFPValue(double V, Type *Ty) { - if (Ty->isHalfTy() || Ty->isFloatTy()) { + if (Ty->isHalfTy() || Ty->isFloatTy() || Ty->isBFloatTy()) { APFloat APF(V); bool unused; APF.convert(Ty->getFltSemantics(), APFloat::rmNearestTiesToEven, &unused); @@ -2217,7 +2217,7 @@ Constant *GetConstantFoldFPValue(double V, Type *Ty) { } if (Ty->isDoubleTy()) return ConstantFP::get(Ty->getContext(), APFloat(V)); - llvm_unreachable("Can only constant fold half/float/double"); + llvm_unreachable("Can only constant fold half/float/double/bfloat"); } #if defined(HAS_IEE754_FLOAT128) && defined(HAS_LOGF128) @@ -2652,7 +2652,7 @@ static Constant *ConstantFoldScalarCall1(StringRef Name, #endif if (!Ty->isHalfTy() && !Ty->isFloatTy() && !Ty->isDoubleTy() && - !Ty->isIntegerTy()) + !Ty->isIntegerTy() && !Ty->isBFloatTy()) return nullptr; // Use internal versions of these intrinsics. diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cos.bf16.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cos.bf16.ll index ec7f254a41bff..b45e8109c9e2e 100644 --- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cos.bf16.ll +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cos.bf16.ll @@ -39,8 +39,8 @@ define amdgpu_kernel void @cos_bf16(ptr addrspace(1) %out, bfloat %src) #1 { ret void } -define amdgpu_kernel void @cos_bf16_constant_4(ptr addrspace(1) %out) #1 { -; FAKE16-LABEL: cos_bf16_constant_4: +define amdgpu_kernel void @cos_bf16_constant_4_strictfp(ptr addrspace(1) %out) #2 { +; FAKE16-LABEL: cos_bf16_constant_4_strictfp: ; FAKE16: ; %bb.0: ; FAKE16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 ; FAKE16-NEXT: s_mov_b64 s[64:65], 0 @@ -53,7 +53,7 @@ define amdgpu_kernel void @cos_bf16_constant_4(ptr addrspace(1) %out) #1 { ; FAKE16-NEXT: global_store_b16 v1, v0, s[0:1] ; FAKE16-NEXT: s_endpgm ; -; REAL16-LABEL: cos_bf16_constant_4: +; REAL16-LABEL: cos_bf16_constant_4_strictfp: ; REAL16: ; %bb.0: ; REAL16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 ; REAL16-NEXT: s_mov_b64 s[64:65], 0 @@ -65,13 +65,13 @@ define amdgpu_kernel void @cos_bf16_constant_4(ptr addrspace(1) %out) #1 { ; REAL16-NEXT: s_wait_kmcnt 0x0 ; REAL16-NEXT: global_store_b16 v1, v0, s[0:1] ; REAL16-NEXT: s_endpgm - %cos = call bfloat @llvm.amdgcn.cos.bf16(bfloat 4.0) #0 + %cos = call bfloat @llvm.amdgcn.cos.bf16(bfloat 4.0) strictfp store bfloat %cos, ptr addrspace(1) %out, align 2 ret void } -define amdgpu_kernel void @cos_bf16_constant_100(ptr addrspace(1) %out) #1 { -; FAKE16-LABEL: cos_bf16_constant_100: +define amdgpu_kernel void @cos_bf16_constant_100_strictfp(ptr addrspace(1) %out) #2 { +; FAKE16-LABEL: cos_bf16_constant_100_strictfp: ; FAKE16: ; %bb.0: ; FAKE16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 ; FAKE16-NEXT: s_mov_b64 s[64:65], 0 @@ -84,7 +84,7 @@ define amdgpu_kernel void @cos_bf16_constant_100(ptr addrspace(1) %out) #1 { ; FAKE16-NEXT: global_store_b16 v1, v0, s[0:1] ; FAKE16-NEXT: s_endpgm ; -; REAL16-LABEL: cos_bf16_constant_100: +; REAL16-LABEL: cos_bf16_constant_100_strictfp: ; REAL16: ; %bb.0: ; REAL16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 ; REAL16-NEXT: s_mov_b64 s[64:65], 0 @@ -96,12 +96,43 @@ define amdgpu_kernel void @cos_bf16_constant_100(ptr addrspace(1) %out) #1 { ; REAL16-NEXT: s_wait_kmcnt 0x0 ; REAL16-NEXT: global_store_b16 v1, v0, s[0:1] ; REAL16-NEXT: s_endpgm - %cos = call bfloat @llvm.amdgcn.cos.bf16(bfloat 100.0) #0 + %cos = call bfloat @llvm.amdgcn.cos.bf16(bfloat 100.0) strictfp + store bfloat %cos, ptr addrspace(1) %out, align 2 + ret void +} + +define amdgpu_kernel void @cos_bf16_constant_0.3(ptr addrspace(1) %out) #1 { +; FAKE16-LABEL: cos_bf16_constant_0.3: +; FAKE16: ; %bb.0: +; FAKE16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 +; FAKE16-NEXT: s_mov_b64 s[64:65], 0 +; FAKE16-NEXT: v_nop +; FAKE16-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE +; FAKE16-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv +; FAKE16-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, 0xffffbea1 +; FAKE16-NEXT: s_wait_kmcnt 0x0 +; FAKE16-NEXT: global_store_b16 v0, v1, s[0:1] +; FAKE16-NEXT: s_endpgm +; +; REAL16-LABEL: cos_bf16_constant_0.3: +; REAL16: ; %bb.0: +; REAL16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 +; REAL16-NEXT: s_mov_b64 s[64:65], 0 +; REAL16-NEXT: v_nop +; REAL16-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE +; REAL16-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv +; REAL16-NEXT: v_mov_b32_e32 v1, 0 +; REAL16-NEXT: v_mov_b16_e32 v0.l, 0xbea1 +; REAL16-NEXT: s_wait_kmcnt 0x0 +; REAL16-NEXT: global_store_b16 v1, v0, s[0:1] +; REAL16-NEXT: s_endpgm + %cos = call bfloat @llvm.amdgcn.cos.bf16(bfloat 0.3) #0 store bfloat %cos, ptr addrspace(1) %out, align 2 ret void } attributes #0 = { nounwind readnone } attributes #1 = { nounwind } +attributes #2 = { nounwind strictfp } ;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line: ; GCN: {{.*}} diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sin.bf16.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sin.bf16.ll index dcce910331713..dedf9593073b1 100644 --- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sin.bf16.ll +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sin.bf16.ll @@ -39,8 +39,8 @@ define amdgpu_kernel void @sin_bf16(ptr addrspace(1) %out, bfloat %src) #1 { ret void } -define amdgpu_kernel void @sin_bf16_constant_4(ptr addrspace(1) %out) #1 { -; FAKE16-LABEL: sin_bf16_constant_4: +define amdgpu_kernel void @sin_bf16_constant_4_strictfp(ptr addrspace(1) %out) #2 { +; FAKE16-LABEL: sin_bf16_constant_4_strictfp: ; FAKE16: ; %bb.0: ; FAKE16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 ; FAKE16-NEXT: s_mov_b64 s[64:65], 0 @@ -53,7 +53,7 @@ define amdgpu_kernel void @sin_bf16_constant_4(ptr addrspace(1) %out) #1 { ; FAKE16-NEXT: global_store_b16 v1, v0, s[0:1] ; FAKE16-NEXT: s_endpgm ; -; REAL16-LABEL: sin_bf16_constant_4: +; REAL16-LABEL: sin_bf16_constant_4_strictfp: ; REAL16: ; %bb.0: ; REAL16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 ; REAL16-NEXT: s_mov_b64 s[64:65], 0 @@ -65,13 +65,13 @@ define amdgpu_kernel void @sin_bf16_constant_4(ptr addrspace(1) %out) #1 { ; REAL16-NEXT: s_wait_kmcnt 0x0 ; REAL16-NEXT: global_store_b16 v1, v0, s[0:1] ; REAL16-NEXT: s_endpgm - %sin = call bfloat @llvm.amdgcn.sin.bf16(bfloat 4.0) #0 + %sin = call bfloat @llvm.amdgcn.sin.bf16(bfloat 4.0) strictfp store bfloat %sin, ptr addrspace(1) %out, align 2 ret void } -define amdgpu_kernel void @sin_bf16_constant_100(ptr addrspace(1) %out) #1 { -; FAKE16-LABEL: sin_bf16_constant_100: +define amdgpu_kernel void @sin_bf16_constant_100_strictfp(ptr addrspace(1) %out) #2 { +; FAKE16-LABEL: sin_bf16_constant_100_strictfp: ; FAKE16: ; %bb.0: ; FAKE16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 ; FAKE16-NEXT: s_mov_b64 s[64:65], 0 @@ -84,7 +84,7 @@ define amdgpu_kernel void @sin_bf16_constant_100(ptr addrspace(1) %out) #1 { ; FAKE16-NEXT: global_store_b16 v1, v0, s[0:1] ; FAKE16-NEXT: s_endpgm ; -; REAL16-LABEL: sin_bf16_constant_100: +; REAL16-LABEL: sin_bf16_constant_100_strictfp: ; REAL16: ; %bb.0: ; REAL16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 ; REAL16-NEXT: s_mov_b64 s[64:65], 0 @@ -96,12 +96,43 @@ define amdgpu_kernel void @sin_bf16_constant_100(ptr addrspace(1) %out) #1 { ; REAL16-NEXT: s_wait_kmcnt 0x0 ; REAL16-NEXT: global_store_b16 v1, v0, s[0:1] ; REAL16-NEXT: s_endpgm - %sin = call bfloat @llvm.amdgcn.sin.bf16(bfloat 100.0) #0 + %sin = call bfloat @llvm.amdgcn.sin.bf16(bfloat 100.0) strictfp + store bfloat %sin, ptr addrspace(1) %out, align 2 + ret void +} + +define amdgpu_kernel void @sin_bf16_constant_0.3(ptr addrspace(1) %out) #1 { +; FAKE16-LABEL: sin_bf16_constant_0.3: +; FAKE16: ; %bb.0: +; FAKE16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 +; FAKE16-NEXT: s_mov_b64 s[64:65], 0 +; FAKE16-NEXT: v_nop +; FAKE16-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE +; FAKE16-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv +; FAKE16-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, 0x3f73 +; FAKE16-NEXT: s_wait_kmcnt 0x0 +; FAKE16-NEXT: global_store_b16 v0, v1, s[0:1] +; FAKE16-NEXT: s_endpgm +; +; REAL16-LABEL: sin_bf16_constant_0.3: +; REAL16: ; %bb.0: +; REAL16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 +; REAL16-NEXT: s_mov_b64 s[64:65], 0 +; REAL16-NEXT: v_nop +; REAL16-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE +; REAL16-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv +; REAL16-NEXT: v_mov_b32_e32 v1, 0 +; REAL16-NEXT: v_mov_b16_e32 v0.l, 0x3f73 +; REAL16-NEXT: s_wait_kmcnt 0x0 +; REAL16-NEXT: global_store_b16 v1, v0, s[0:1] +; REAL16-NEXT: s_endpgm + %sin = call bfloat @llvm.amdgcn.sin.bf16(bfloat 0.3) store bfloat %sin, ptr addrspace(1) %out, align 2 ret void } attributes #0 = { nounwind readnone } attributes #1 = { nounwind } +attributes #2 = { nounwind strictfp } ;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line: ; GCN: {{.*}} diff --git a/llvm/test/CodeGen/AMDGPU/llvm.cos.bf16.ll b/llvm/test/CodeGen/AMDGPU/llvm.cos.bf16.ll index fa5abebf3d593..77764119529ae 100644 --- a/llvm/test/CodeGen/AMDGPU/llvm.cos.bf16.ll +++ b/llvm/test/CodeGen/AMDGPU/llvm.cos.bf16.ll @@ -43,8 +43,8 @@ define amdgpu_kernel void @cos_bf16(ptr addrspace(1) %out, bfloat %src) #1 { ret void } -define amdgpu_kernel void @cos_bf16_constant_4(ptr addrspace(1) %out) #1 { -; FAKE16-LABEL: cos_bf16_constant_4: +define amdgpu_kernel void @cos_bf16_constant_4_strictfp(ptr addrspace(1) %out) #2 { +; FAKE16-LABEL: cos_bf16_constant_4_strictfp: ; FAKE16: ; %bb.0: ; FAKE16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 ; FAKE16-NEXT: s_mov_b64 s[64:65], 0 @@ -57,7 +57,7 @@ define amdgpu_kernel void @cos_bf16_constant_4(ptr addrspace(1) %out) #1 { ; FAKE16-NEXT: global_store_b16 v1, v0, s[0:1] ; FAKE16-NEXT: s_endpgm ; -; REAL16-LABEL: cos_bf16_constant_4: +; REAL16-LABEL: cos_bf16_constant_4_strictfp: ; REAL16: ; %bb.0: ; REAL16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 ; REAL16-NEXT: s_mov_b64 s[64:65], 0 @@ -69,13 +69,13 @@ define amdgpu_kernel void @cos_bf16_constant_4(ptr addrspace(1) %out) #1 { ; REAL16-NEXT: s_wait_kmcnt 0x0 ; REAL16-NEXT: global_store_b16 v1, v0, s[0:1] ; REAL16-NEXT: s_endpgm - %cos = call bfloat @llvm.cos.bf16(bfloat 4.0) #0 + %cos = call bfloat @llvm.cos.bf16(bfloat 4.0) strictfp store bfloat %cos, ptr addrspace(1) %out, align 2 ret void } -define amdgpu_kernel void @cos_bf16_constant_100(ptr addrspace(1) %out) #1 { -; FAKE16-LABEL: cos_bf16_constant_100: +define amdgpu_kernel void @cos_bf16_constant_100_strictfp(ptr addrspace(1) %out) #2 { +; FAKE16-LABEL: cos_bf16_constant_100_strictfp: ; FAKE16: ; %bb.0: ; FAKE16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 ; FAKE16-NEXT: s_mov_b64 s[64:65], 0 @@ -88,7 +88,7 @@ define amdgpu_kernel void @cos_bf16_constant_100(ptr addrspace(1) %out) #1 { ; FAKE16-NEXT: global_store_b16 v1, v0, s[0:1] ; FAKE16-NEXT: s_endpgm ; -; REAL16-LABEL: cos_bf16_constant_100: +; REAL16-LABEL: cos_bf16_constant_100_strictfp: ; REAL16: ; %bb.0: ; REAL16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 ; REAL16-NEXT: s_mov_b64 s[64:65], 0 @@ -100,12 +100,43 @@ define amdgpu_kernel void @cos_bf16_constant_100(ptr addrspace(1) %out) #1 { ; REAL16-NEXT: s_wait_kmcnt 0x0 ; REAL16-NEXT: global_store_b16 v1, v0, s[0:1] ; REAL16-NEXT: s_endpgm - %cos = call bfloat @llvm.cos.bf16(bfloat 100.0) #0 + %cos = call bfloat @llvm.cos.bf16(bfloat 100.0) strictfp + store bfloat %cos, ptr addrspace(1) %out, align 2 + ret void +} + +define amdgpu_kernel void @cos_bf16_constant_0.3(ptr addrspace(1) %out) #1 { +; FAKE16-LABEL: cos_bf16_constant_0.3: +; FAKE16: ; %bb.0: +; FAKE16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 +; FAKE16-NEXT: s_mov_b64 s[64:65], 0 +; FAKE16-NEXT: v_nop +; FAKE16-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE +; FAKE16-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv +; FAKE16-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, 0x3f75 +; FAKE16-NEXT: s_wait_kmcnt 0x0 +; FAKE16-NEXT: global_store_b16 v0, v1, s[0:1] +; FAKE16-NEXT: s_endpgm +; +; REAL16-LABEL: cos_bf16_constant_0.3: +; REAL16: ; %bb.0: +; REAL16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 +; REAL16-NEXT: s_mov_b64 s[64:65], 0 +; REAL16-NEXT: v_nop +; REAL16-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE +; REAL16-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv +; REAL16-NEXT: v_mov_b32_e32 v1, 0 +; REAL16-NEXT: v_mov_b16_e32 v0.l, 0x3f75 +; REAL16-NEXT: s_wait_kmcnt 0x0 +; REAL16-NEXT: global_store_b16 v1, v0, s[0:1] +; REAL16-NEXT: s_endpgm + %cos = call bfloat @llvm.cos.bf16(bfloat 0.3) store bfloat %cos, ptr addrspace(1) %out, align 2 ret void } attributes #0 = { nounwind readnone } attributes #1 = { nounwind } +attributes #2 = { nounwind strictfp } ;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line: ; GCN: {{.*}} diff --git a/llvm/test/CodeGen/AMDGPU/llvm.sin.bf16.ll b/llvm/test/CodeGen/AMDGPU/llvm.sin.bf16.ll index 3625cb7e5dccf..d15f7b6b0bf74 100644 --- a/llvm/test/CodeGen/AMDGPU/llvm.sin.bf16.ll +++ b/llvm/test/CodeGen/AMDGPU/llvm.sin.bf16.ll @@ -43,8 +43,8 @@ define amdgpu_kernel void @sin_bf16(ptr addrspace(1) %out, bfloat %src) #1 { ret void } -define amdgpu_kernel void @sin_bf16_constant_4(ptr addrspace(1) %out) #1 { -; FAKE16-LABEL: sin_bf16_constant_4: +define amdgpu_kernel void @sin_bf16_constant_4_strictfp(ptr addrspace(1) %out) #2 { +; FAKE16-LABEL: sin_bf16_constant_4_strictfp: ; FAKE16: ; %bb.0: ; FAKE16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 ; FAKE16-NEXT: s_mov_b64 s[64:65], 0 @@ -57,7 +57,7 @@ define amdgpu_kernel void @sin_bf16_constant_4(ptr addrspace(1) %out) #1 { ; FAKE16-NEXT: global_store_b16 v1, v0, s[0:1] ; FAKE16-NEXT: s_endpgm ; -; REAL16-LABEL: sin_bf16_constant_4: +; REAL16-LABEL: sin_bf16_constant_4_strictfp: ; REAL16: ; %bb.0: ; REAL16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 ; REAL16-NEXT: s_mov_b64 s[64:65], 0 @@ -69,13 +69,13 @@ define amdgpu_kernel void @sin_bf16_constant_4(ptr addrspace(1) %out) #1 { ; REAL16-NEXT: s_wait_kmcnt 0x0 ; REAL16-NEXT: global_store_b16 v1, v0, s[0:1] ; REAL16-NEXT: s_endpgm - %sin = call bfloat @llvm.sin.bf16(bfloat 4.0) #0 + %sin = call bfloat @llvm.sin.bf16(bfloat 4.0) strictfp store bfloat %sin, ptr addrspace(1) %out, align 2 ret void } -define amdgpu_kernel void @sin_bf16_constant_100(ptr addrspace(1) %out) #1 { -; FAKE16-LABEL: sin_bf16_constant_100: +define amdgpu_kernel void @sin_bf16_constant_100_strictfp(ptr addrspace(1) %out) #2 { +; FAKE16-LABEL: sin_bf16_constant_100_strictfp: ; FAKE16: ; %bb.0: ; FAKE16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 ; FAKE16-NEXT: s_mov_b64 s[64:65], 0 @@ -88,7 +88,7 @@ define amdgpu_kernel void @sin_bf16_constant_100(ptr addrspace(1) %out) #1 { ; FAKE16-NEXT: global_store_b16 v1, v0, s[0:1] ; FAKE16-NEXT: s_endpgm ; -; REAL16-LABEL: sin_bf16_constant_100: +; REAL16-LABEL: sin_bf16_constant_100_strictfp: ; REAL16: ; %bb.0: ; REAL16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 ; REAL16-NEXT: s_mov_b64 s[64:65], 0 @@ -100,12 +100,43 @@ define amdgpu_kernel void @sin_bf16_constant_100(ptr addrspace(1) %out) #1 { ; REAL16-NEXT: s_wait_kmcnt 0x0 ; REAL16-NEXT: global_store_b16 v1, v0, s[0:1] ; REAL16-NEXT: s_endpgm - %sin = call bfloat @llvm.sin.bf16(bfloat 100.0) #0 + %sin = call bfloat @llvm.sin.bf16(bfloat 100.0) strictfp + store bfloat %sin, ptr addrspace(1) %out, align 2 + ret void +} + +define amdgpu_kernel void @sin_bf16_constant_0.3(ptr addrspace(1) %out) #1 { +; FAKE16-LABEL: sin_bf16_constant_0.3: +; FAKE16: ; %bb.0: +; FAKE16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 +; FAKE16-NEXT: s_mov_b64 s[64:65], 0 +; FAKE16-NEXT: v_nop +; FAKE16-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE +; FAKE16-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv +; FAKE16-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, 0x3e98 +; FAKE16-NEXT: s_wait_kmcnt 0x0 +; FAKE16-NEXT: global_store_b16 v0, v1, s[0:1] +; FAKE16-NEXT: s_endpgm +; +; REAL16-LABEL: sin_bf16_constant_0.3: +; REAL16: ; %bb.0: +; REAL16-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0 +; REAL16-NEXT: s_mov_b64 s[64:65], 0 +; REAL16-NEXT: v_nop +; REAL16-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE +; REAL16-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv +; REAL16-NEXT: v_mov_b32_e32 v1, 0 +; REAL16-NEXT: v_mov_b16_e32 v0.l, 0x3e98 +; REAL16-NEXT: s_wait_kmcnt 0x0 +; REAL16-NEXT: global_store_b16 v1, v0, s[0:1] +; REAL16-NEXT: s_endpgm + %sin = call bfloat @llvm.sin.bf16(bfloat 0.3) store bfloat %sin, ptr addrspace(1) %out, align 2 ret void } attributes #0 = { nounwind readnone } attributes #1 = { nounwind } +attributes #2 = { nounwind strictfp } ;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line: ; GCN: {{.*}} diff --git a/llvm/test/Transforms/InstCombine/AMDGPU/amdgcn-intrinsics.ll b/llvm/test/Transforms/InstCombine/AMDGPU/amdgcn-intrinsics.ll index 041a67095791c..14f23c4cc7692 100644 --- a/llvm/test/Transforms/InstCombine/AMDGPU/amdgcn-intrinsics.ll +++ b/llvm/test/Transforms/InstCombine/AMDGPU/amdgcn-intrinsics.ll @@ -43,7 +43,7 @@ define double @test_constant_fold_rcp_f64_1() nounwind { define float @test_constant_fold_rcp_f32_half() nounwind { ; CHECK-LABEL: @test_constant_fold_rcp_f32_half( -; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.rcp.f32(float 5.000000e-01) #[[ATTR18:[0-9]+]] +; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.rcp.f32(float 5.000000e-01) #[[ATTR12:[0-9]+]] ; CHECK-NEXT: ret float [[VAL]] ; %val = call float @llvm.amdgcn.rcp.f32(float 0.5) nounwind readnone @@ -52,7 +52,7 @@ define float @test_constant_fold_rcp_f32_half() nounwind { define double @test_constant_fold_rcp_f64_half() nounwind { ; CHECK-LABEL: @test_constant_fold_rcp_f64_half( -; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.rcp.f64(double 5.000000e-01) #[[ATTR18]] +; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.rcp.f64(double 5.000000e-01) #[[ATTR12]] ; CHECK-NEXT: ret double [[VAL]] ; %val = call double @llvm.amdgcn.rcp.f64(double 0.5) nounwind readnone @@ -61,7 +61,7 @@ define double @test_constant_fold_rcp_f64_half() nounwind { define float @test_constant_fold_rcp_f32_43() nounwind { ; CHECK-LABEL: @test_constant_fold_rcp_f32_43( -; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.rcp.f32(float 4.300000e+01) #[[ATTR18]] +; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.rcp.f32(float 4.300000e+01) #[[ATTR12]] ; CHECK-NEXT: ret float [[VAL]] ; %val = call float @llvm.amdgcn.rcp.f32(float 4.300000e+01) nounwind readnone @@ -70,7 +70,7 @@ define float @test_constant_fold_rcp_f32_43() nounwind { define double @test_constant_fold_rcp_f64_43() nounwind { ; CHECK-LABEL: @test_constant_fold_rcp_f64_43( -; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.rcp.f64(double 4.300000e+01) #[[ATTR18]] +; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.rcp.f64(double 4.300000e+01) #[[ATTR12]] ; CHECK-NEXT: ret double [[VAL]] ; %val = call double @llvm.amdgcn.rcp.f64(double 4.300000e+01) nounwind readnone @@ -79,7 +79,7 @@ define double @test_constant_fold_rcp_f64_43() nounwind { define float @test_constant_fold_rcp_f32_43_strictfp() nounwind strictfp { ; CHECK-LABEL: @test_constant_fold_rcp_f32_43_strictfp( -; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.rcp.f32(float 4.300000e+01) #[[ATTR19:[0-9]+]] +; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.rcp.f32(float 4.300000e+01) #[[ATTR13:[0-9]+]] ; CHECK-NEXT: ret float [[VAL]] ; %val = call float @llvm.amdgcn.rcp.f32(float 4.300000e+01) strictfp nounwind readnone @@ -190,7 +190,7 @@ define half @test_constant_fold_sqrt_f16_0() nounwind { define float @test_constant_fold_sqrt_f32_0() nounwind { ; CHECK-LABEL: @test_constant_fold_sqrt_f32_0( -; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.sqrt.f32(float 0.000000e+00) #[[ATTR18]] +; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.sqrt.f32(float 0.000000e+00) #[[ATTR12]] ; CHECK-NEXT: ret float [[VAL]] ; %val = call float @llvm.amdgcn.sqrt.f32(float 0.0) nounwind readnone @@ -199,7 +199,7 @@ define float @test_constant_fold_sqrt_f32_0() nounwind { define double @test_constant_fold_sqrt_f64_0() nounwind { ; CHECK-LABEL: @test_constant_fold_sqrt_f64_0( -; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.sqrt.f64(double 0.000000e+00) #[[ATTR18]] +; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.sqrt.f64(double 0.000000e+00) #[[ATTR12]] ; CHECK-NEXT: ret double [[VAL]] ; %val = call double @llvm.amdgcn.sqrt.f64(double 0.0) nounwind readnone @@ -216,7 +216,7 @@ define half @test_constant_fold_sqrt_f16_neg0() nounwind { define float @test_constant_fold_sqrt_f32_neg0() nounwind { ; CHECK-LABEL: @test_constant_fold_sqrt_f32_neg0( -; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.sqrt.f32(float -0.000000e+00) #[[ATTR18]] +; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.sqrt.f32(float -0.000000e+00) #[[ATTR12]] ; CHECK-NEXT: ret float [[VAL]] ; %val = call float @llvm.amdgcn.sqrt.f32(float -0.0) nounwind readnone @@ -225,7 +225,7 @@ define float @test_constant_fold_sqrt_f32_neg0() nounwind { define double @test_constant_fold_sqrt_f64_neg0() nounwind { ; CHECK-LABEL: @test_constant_fold_sqrt_f64_neg0( -; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.sqrt.f64(double -0.000000e+00) #[[ATTR18]] +; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.sqrt.f64(double -0.000000e+00) #[[ATTR12]] ; CHECK-NEXT: ret double [[VAL]] ; %val = call double @llvm.amdgcn.sqrt.f64(double -0.0) nounwind readnone @@ -817,7 +817,7 @@ define i1 @test_class_isnan_f32(float %x) nounwind { define i1 @test_class_isnan_f32_strict(float %x) nounwind strictfp { ; CHECK-LABEL: @test_class_isnan_f32_strict( -; CHECK-NEXT: [[VAL:%.*]] = call i1 @llvm.is.fpclass.f32(float [[X:%.*]], /* (nan) */ i32 3) #[[ATTR20:[0-9]+]] +; CHECK-NEXT: [[VAL:%.*]] = call i1 @llvm.is.fpclass.f32(float [[X:%.*]], /* (nan) */ i32 3) #[[ATTR10:[0-9]+]] ; CHECK-NEXT: ret i1 [[VAL]] ; %val = call i1 @llvm.amdgcn.class.f32(float %x, i32 3) strictfp @@ -835,7 +835,7 @@ define i1 @test_class_is_p0_n0_f32(float %x) nounwind { define i1 @test_class_is_p0_n0_f32_strict(float %x) nounwind strictfp { ; CHECK-LABEL: @test_class_is_p0_n0_f32_strict( -; CHECK-NEXT: [[VAL:%.*]] = call i1 @llvm.is.fpclass.f32(float [[X:%.*]], /* (zero) */ i32 96) #[[ATTR20]] +; CHECK-NEXT: [[VAL:%.*]] = call i1 @llvm.is.fpclass.f32(float [[X:%.*]], /* (zero) */ i32 96) #[[ATTR10]] ; CHECK-NEXT: ret i1 [[VAL]] ; %val = call i1 @llvm.amdgcn.class.f32(float %x, i32 96) strictfp @@ -1068,6 +1068,7 @@ define i1 @test_class_is_nan_other_nnan_src(float %x) { ; llvm.amdgcn.cos ; -------------------------------------------------------------------- declare float @llvm.amdgcn.cos.f32(float) nounwind readnone +declare bfloat @llvm.amdgcn.cos.bf16(bfloat) nounwind readnone declare float @llvm.fabs.f32(float) nounwind readnone define float @cos_fneg_f32(float %x) { @@ -1122,11 +1123,19 @@ define float @cos_fabs_unary_fneg_f32(float %x) { ret float %cos } +define bfloat @cos_bf16_constant_fold() { +; CHECK-LABEL: @cos_bf16_constant_fold( +; CHECK-NEXT: ret bfloat -1.000000e+00 +; + %sin = call bfloat @llvm.amdgcn.cos.bf16(bfloat 0.5) + ret bfloat %sin +} ; -------------------------------------------------------------------- ; llvm.amdgcn.sin ; -------------------------------------------------------------------- declare float @llvm.amdgcn.sin.f32(float) nounwind readnone +declare bfloat @llvm.amdgcn.sin.bf16(bfloat) nounwind readnone define float @sin_fneg_f32(float %x) { ; CHECK-LABEL: @sin_fneg_f32( @@ -1176,6 +1185,14 @@ define float @sin_fabs_fneg_fast_f32(float %x) { ret float %sin } +define bfloat @sin_bf16_constant_fold() { +; CHECK-LABEL: @sin_bf16_constant_fold( +; CHECK-NEXT: ret bfloat 0.000000e+00 +; + %sin = call bfloat @llvm.amdgcn.sin.bf16(bfloat 0.5) + ret bfloat %sin +} + ; -------------------------------------------------------------------- ; llvm.amdgcn.cvt.pkrtz ; -------------------------------------------------------------------- @@ -5843,7 +5860,7 @@ define double @trig_preop_constfold_neg32_segment() { define double @trig_preop_constfold_strictfp() strictfp { ; CHECK-LABEL: @trig_preop_constfold_strictfp( -; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.trig.preop.f64(double 3.454350e+02, i32 5) #[[ATTR20]] +; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.trig.preop.f64(double 3.454350e+02, i32 5) #[[ATTR10]] ; CHECK-NEXT: ret double [[VAL]] ; %val = call double @llvm.amdgcn.trig.preop.f64(double 3.454350e+02, i32 5) strictfp @@ -6137,7 +6154,7 @@ define double @trig_preop_strip_copysign(double %mag, double %sign, i32 %idx) { define double @trig_preop_strip_fabs_strictfp(double %val, i32 %idx) strictfp { ; CHECK-LABEL: @trig_preop_strip_fabs_strictfp( -; CHECK-NEXT: [[RESULT:%.*]] = call double @llvm.amdgcn.trig.preop.f64(double [[VAL:%.*]], i32 [[IDX:%.*]]) #[[ATTR20]] +; CHECK-NEXT: [[RESULT:%.*]] = call double @llvm.amdgcn.trig.preop.f64(double [[VAL:%.*]], i32 [[IDX:%.*]]) #[[ATTR10]] ; CHECK-NEXT: ret double [[RESULT]] ; %fabs = call double @llvm.fabs.f64(double %val) @@ -6276,7 +6293,7 @@ define half @test_constant_fold_log_f16_neg10() { define float @test_constant_fold_log_f32_qnan_strictfp() strictfp { ; CHECK-LABEL: @test_constant_fold_log_f32_qnan_strictfp( -; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.log.f32(float +qnan) #[[ATTR20]] +; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.log.f32(float +qnan) #[[ATTR10]] ; CHECK-NEXT: ret float [[VAL]] ; %val = call float @llvm.amdgcn.log.f32(float 0x7FF8000000000000) strictfp @@ -6285,7 +6302,7 @@ define float @test_constant_fold_log_f32_qnan_strictfp() strictfp { define float @test_constant_fold_log_f32_0_strictfp() strictfp { ; CHECK-LABEL: @test_constant_fold_log_f32_0_strictfp( -; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.log.f32(float 0.000000e+00) #[[ATTR20]] +; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.log.f32(float 0.000000e+00) #[[ATTR10]] ; CHECK-NEXT: ret float [[VAL]] ; %val = call float @llvm.amdgcn.log.f32(float 0.0) strictfp @@ -6294,7 +6311,7 @@ define float @test_constant_fold_log_f32_0_strictfp() strictfp { define float @test_constant_fold_log_f32_neg0_strictfp() strictfp { ; CHECK-LABEL: @test_constant_fold_log_f32_neg0_strictfp( -; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.log.f32(float -0.000000e+00) #[[ATTR20]] +; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.log.f32(float -0.000000e+00) #[[ATTR10]] ; CHECK-NEXT: ret float [[VAL]] ; %val = call float @llvm.amdgcn.log.f32(float -0.0) strictfp @@ -6303,7 +6320,7 @@ define float @test_constant_fold_log_f32_neg0_strictfp() strictfp { define float @test_constant_fold_log_f32_neg_strictfp() strictfp { ; CHECK-LABEL: @test_constant_fold_log_f32_neg_strictfp( -; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.log.f32(float -1.000000e+01) #[[ATTR20]] +; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.log.f32(float -1.000000e+01) #[[ATTR10]] ; CHECK-NEXT: ret float [[VAL]] ; %val = call float @llvm.amdgcn.log.f32(float -10.0) strictfp @@ -6320,7 +6337,7 @@ define float @test_constant_fold_log_f32_pinf_strictfp() strictfp { define float @test_constant_fold_log_f32_ninf_strictfp() strictfp { ; CHECK-LABEL: @test_constant_fold_log_f32_ninf_strictfp( -; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.log.f32(float -inf) #[[ATTR20]] +; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.log.f32(float -inf) #[[ATTR10]] ; CHECK-NEXT: ret float [[VAL]] ; %val = call float @llvm.amdgcn.log.f32(float 0xFFF0000000000000) strictfp @@ -6522,7 +6539,7 @@ define half @test_constant_fold_exp2_f16_neg10() { define float @test_constant_fold_exp2_f32_qnan_strictfp() strictfp { ; CHECK-LABEL: @test_constant_fold_exp2_f32_qnan_strictfp( -; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float +qnan) #[[ATTR20]] +; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float +qnan) #[[ATTR10]] ; CHECK-NEXT: ret float [[VAL]] ; %val = call float @llvm.amdgcn.exp2.f32(float 0x7FF8000000000000) strictfp @@ -6531,7 +6548,7 @@ define float @test_constant_fold_exp2_f32_qnan_strictfp() strictfp { define float @test_constant_fold_exp2_f32_0_strictfp() strictfp { ; CHECK-LABEL: @test_constant_fold_exp2_f32_0_strictfp( -; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float 0.000000e+00) #[[ATTR20]] +; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float 0.000000e+00) #[[ATTR10]] ; CHECK-NEXT: ret float [[VAL]] ; %val = call float @llvm.amdgcn.exp2.f32(float 0.0) strictfp @@ -6540,7 +6557,7 @@ define float @test_constant_fold_exp2_f32_0_strictfp() strictfp { define float @test_constant_fold_exp2_f32_neg0_strictfp() strictfp { ; CHECK-LABEL: @test_constant_fold_exp2_f32_neg0_strictfp( -; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float -0.000000e+00) #[[ATTR20]] +; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float -0.000000e+00) #[[ATTR10]] ; CHECK-NEXT: ret float [[VAL]] ; %val = call float @llvm.amdgcn.exp2.f32(float -0.0) strictfp @@ -6549,7 +6566,7 @@ define float @test_constant_fold_exp2_f32_neg0_strictfp() strictfp { define float @test_constant_fold_exp2_f32_1_strictfp() strictfp { ; CHECK-LABEL: @test_constant_fold_exp2_f32_1_strictfp( -; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float 1.000000e+00) #[[ATTR20]] +; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float 1.000000e+00) #[[ATTR10]] ; CHECK-NEXT: ret float [[VAL]] ; %val = call float @llvm.amdgcn.exp2.f32(float 1.0) strictfp @@ -6558,7 +6575,7 @@ define float @test_constant_fold_exp2_f32_1_strictfp() strictfp { define float @test_constant_fold_exp2_f32_neg1_strictfp() strictfp { ; CHECK-LABEL: @test_constant_fold_exp2_f32_neg1_strictfp( -; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float -1.000000e+00) #[[ATTR20]] +; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float -1.000000e+00) #[[ATTR10]] ; CHECK-NEXT: ret float [[VAL]] ; %val = call float @llvm.amdgcn.exp2.f32(float -1.0) strictfp @@ -6567,7 +6584,7 @@ define float @test_constant_fold_exp2_f32_neg1_strictfp() strictfp { define float @test_constant_fold_exp2_f32_2_strictfp() strictfp { ; CHECK-LABEL: @test_constant_fold_exp2_f32_2_strictfp( -; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float 2.000000e+00) #[[ATTR20]] +; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float 2.000000e+00) #[[ATTR10]] ; CHECK-NEXT: ret float [[VAL]] ; %val = call float @llvm.amdgcn.exp2.f32(float 2.0) strictfp @@ -6576,7 +6593,7 @@ define float @test_constant_fold_exp2_f32_2_strictfp() strictfp { define float @test_constant_fold_exp2_f32_neg2_strictfp() strictfp { ; CHECK-LABEL: @test_constant_fold_exp2_f32_neg2_strictfp( -; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float -2.000000e+00) #[[ATTR20]] +; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float -2.000000e+00) #[[ATTR10]] ; CHECK-NEXT: ret float [[VAL]] ; %val = call float @llvm.amdgcn.exp2.f32(float -2.0) strictfp @@ -6585,7 +6602,7 @@ define float @test_constant_fold_exp2_f32_neg2_strictfp() strictfp { define float @test_constant_fold_exp2_f32_neg_strictfp() strictfp { ; CHECK-LABEL: @test_constant_fold_exp2_f32_neg_strictfp( -; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float -1.000000e+01) #[[ATTR20]] +; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float -1.000000e+01) #[[ATTR10]] ; CHECK-NEXT: ret float [[VAL]] ; %val = call float @llvm.amdgcn.exp2.f32(float -10.0) strictfp diff --git a/llvm/test/Transforms/InstCombine/cos-sin-intrinsic.ll b/llvm/test/Transforms/InstCombine/cos-sin-intrinsic.ll index d96d225a2dbc0..7051e8f0f7ff5 100644 --- a/llvm/test/Transforms/InstCombine/cos-sin-intrinsic.ll +++ b/llvm/test/Transforms/InstCombine/cos-sin-intrinsic.ll @@ -3,7 +3,9 @@ declare double @llvm.cos.f64(double %Val) declare float @llvm.cos.f32(float %Val) +declare bfloat @llvm.cos.bf16(bfloat %Val) declare <2 x float> @llvm.cos.v2f32(<2 x float> %Val) +declare <2 x bfloat> @llvm.cos.v2bf16(<2 x bfloat> %Val) declare float @llvm.fabs.f32(float %Val) declare <2 x float> @llvm.fabs.v2f32(<2 x float> %Val) @@ -144,9 +146,27 @@ define <2 x float> @fabs_unary_fneg_v2f32(<2 x float> %x) { ret <2 x float> %cos } +define bfloat @constant_fold_cos_bf16() { +; CHECK-LABEL: @constant_fold_cos_bf16( +; CHECK-NEXT: ret bfloat 8.789060e-01 +; + %r = call bfloat @llvm.cos.bf16(bfloat 0.5) + ret bfloat %r +} + +define <2 x bfloat> @constant_fold_cos_v2bf16() { +; CHECK-LABEL: @constant_fold_cos_v2bf16( +; CHECK-NEXT: ret <2 x bfloat> splat (bfloat 8.789060e-01) +; + %r = call <2 x bfloat> @llvm.cos.v2bf16(<2 x bfloat> <bfloat 0.5, bfloat 0.5>) + ret <2 x bfloat> %r +} + ; Negate is canonicalized after sin. +declare bfloat @llvm.sin.bf16(bfloat) declare <2 x float> @llvm.sin.v2f32(<2 x float>) +declare <2 x bfloat> @llvm.sin.v2bf16(<2 x bfloat> %Val) define <2 x float> @fneg_sin(<2 x float> %x){ ; CHECK-LABEL: @fneg_sin( @@ -193,3 +213,19 @@ define <2 x float> @unary_fneg_sin_fmf(<2 x float> %x){ %r = call nnan arcp afn <2 x float> @llvm.sin.v2f32(<2 x float> %negx) ret <2 x float> %r } + +define bfloat @constant_fold_sin_bf16() { +; CHECK-LABEL: @constant_fold_sin_bf16( +; CHECK-NEXT: ret bfloat 4.785160e-01 +; + %r = call bfloat @llvm.sin.bf16(bfloat 0.5) + ret bfloat %r +} + +define <2 x bfloat> @constant_fold_sin_v2bf16() { +; CHECK-LABEL: @constant_fold_sin_v2bf16( +; CHECK-NEXT: ret <2 x bfloat> splat (bfloat 4.785160e-01) +; + %r = call <2 x bfloat> @llvm.sin.v2bf16(<2 x bfloat> <bfloat 0.5, bfloat 0.5>) + ret <2 x bfloat> %r +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
