https://github.com/jhuber6 created
https://github.com/llvm/llvm-project/pull/90994
Summary:
Previous patches added support for the LLVM rounding intrinsic
functions. This patch allows them to me emitted using the clang builtins
when targeting AMDGPU.
>From abceb892df93ccfbfe9392fc7de8c93822e85f92 Mon Sep 17 00:00:00 2001
From: Joseph Huber <hube...@outlook.com>
Date: Fri, 3 May 2024 13:55:12 -0500
Subject: [PATCH] [AMDGPU] Allow the `__builtin_flt_rounds` functions on AMDGPU
Summary:
Previous patches added support for the LLVM rounding intrinsic
functions. This patch allows them to me emitted using the clang builtins
when targeting AMDGPU.
---
clang/lib/Sema/SemaChecking.cpp | 16 ++++++++--------
clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 12 ++++++++++++
2 files changed, 20 insertions(+), 8 deletions(-)
diff --git a/clang/lib/Sema/SemaChecking.cpp b/clang/lib/Sema/SemaChecking.cpp
index cf8840c63024d4..f5af0de57b1628 100644
--- a/clang/lib/Sema/SemaChecking.cpp
+++ b/clang/lib/Sema/SemaChecking.cpp
@@ -2535,18 +2535,18 @@ Sema::CheckBuiltinFunctionCall(FunctionDecl *FDecl,
unsigned BuiltinID,
case Builtin::BI_bittestandset64:
case Builtin::BI_interlockedbittestandreset64:
case Builtin::BI_interlockedbittestandset64:
- if (CheckBuiltinTargetInSupported(*this, BuiltinID, TheCall,
- {llvm::Triple::x86_64, llvm::Triple::arm,
- llvm::Triple::thumb,
- llvm::Triple::aarch64}))
+ if (CheckBuiltinTargetInSupported(
+ *this, BuiltinID, TheCall,
+ {llvm::Triple::x86_64, llvm::Triple::arm, llvm::Triple::thumb,
+ llvm::Triple::aarch64, llvm::Triple::amdgcn}))
return ExprError();
break;
case Builtin::BI__builtin_set_flt_rounds:
- if (CheckBuiltinTargetInSupported(*this, BuiltinID, TheCall,
- {llvm::Triple::x86, llvm::Triple::x86_64,
- llvm::Triple::arm, llvm::Triple::thumb,
- llvm::Triple::aarch64}))
+ if (CheckBuiltinTargetInSupported(
+ *this, BuiltinID, TheCall,
+ {llvm::Triple::x86, llvm::Triple::x86_64, llvm::Triple::arm,
+ llvm::Triple::thumb, llvm::Triple::aarch64,
llvm::Triple::amdgcn}))
return ExprError();
break;
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
index bdca97c8878670..338d6bc95655a3 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -839,6 +839,18 @@ unsigned test_wavefrontsize() {
return __builtin_amdgcn_wavefrontsize();
}
+// CHECK-LABEL test_flt_rounds(
+unsigned test_flt_rounds() {
+
+ // CHECK: call i32 @llvm.get.rounding()
+ unsigned mode = __builtin_flt_rounds();
+
+ // CHECK: call void @llvm.set.rounding(i32 %0)
+ __builtin_set_flt_rounds(mode);
+
+ return mode;
+}
+
// CHECK-LABEL test_get_fpenv(
unsigned long test_get_fpenv() {
// CHECK: call i64 @llvm.get.fpenv.i64()
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
