llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-llvm-ir Author: Joseph Huber (jhuber6) <details> <summary>Changes</summary> Summary: Some recent support made usage of `__nvvm_reflect` more consistent. We should expose it as an intrinsic rather than forcing users to externally define the function. --- Full diff: https://github.com/llvm/llvm-project/pull/81277.diff 3 Files Affected: - (modified) clang/include/clang/Basic/BuiltinsNVPTX.def (+1) - (modified) clang/test/CodeGen/builtins-nvptx.c (+8) - (modified) llvm/include/llvm/IR/IntrinsicsNVVM.td (+3-2) ``````````diff diff --git a/clang/include/clang/Basic/BuiltinsNVPTX.def b/clang/include/clang/Basic/BuiltinsNVPTX.def index 7819e71d7fe2aa..8d3c5e69d55cf4 100644 --- a/clang/include/clang/Basic/BuiltinsNVPTX.def +++ b/clang/include/clang/Basic/BuiltinsNVPTX.def @@ -159,6 +159,7 @@ BUILTIN(__nvvm_read_ptx_sreg_pm3, "i", "n") BUILTIN(__nvvm_prmt, "UiUiUiUi", "") BUILTIN(__nvvm_exit, "v", "r") +BUILTIN(__nvvm_reflect, "UicC*", "r") TARGET_BUILTIN(__nvvm_nanosleep, "vUi", "n", AND(SM_70, PTX63)) // Min Max diff --git a/clang/test/CodeGen/builtins-nvptx.c b/clang/test/CodeGen/builtins-nvptx.c index ad7c27f2d60d26..4dba7670b5c43e 100644 --- a/clang/test/CodeGen/builtins-nvptx.c +++ b/clang/test/CodeGen/builtins-nvptx.c @@ -44,6 +44,14 @@ __device__ int read_tid() { } +__device__ bool reflect() { + +// CHECK: call i32 @llvm.nvvm.reflect(ptr {{.*}}) + + unsigned x = __nvvm_reflect("__CUDA_ARCH"); + return x >= 700; +} + __device__ int read_ntid() { // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td index d825dc82156432..f7b0fe926959b1 100644 --- a/llvm/include/llvm/IR/IntrinsicsNVVM.td +++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td @@ -1624,8 +1624,9 @@ def int_nvvm_compiler_error : def int_nvvm_compiler_warn : Intrinsic<[], [llvm_anyptr_ty], [], "llvm.nvvm.compiler.warn">; -def int_nvvm_reflect : - Intrinsic<[llvm_i32_ty], [llvm_anyptr_ty], [IntrNoMem], "llvm.nvvm.reflect">; +def int_nvvm_reflect : + Intrinsic<[llvm_i32_ty], [llvm_ptr_ty], [IntrNoMem], "llvm.nvvm.reflect">, + ClangBuiltin<"__nvvm_reflect">; // isspacep.{const, global, local, shared} def int_nvvm_isspacep_const `````````` </details> https://github.com/llvm/llvm-project/pull/81277 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
