[NVPTX] Add clang builtin for __nvvm_reflect intrinsic

[jhuber6]

Summary:
Some recent support made usage of __nvvm_reflect more consistent. We
should expose it as a builtin rather than forcing users to externally
define the function.

[llvmbot]

@llvm/pr-subscribers-clang

@llvm/pr-subscribers-llvm-ir

Author: Joseph Huber (jhuber6)

Changes

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 --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

[Artem-B]

We should expose it as an intrinsic

I think you mean builtin here.

[Artem-B]

LGTM overall.

[Artem-B]

Now that we're exposing it to the end users. We should probably document what it does.
Probably somewhere in https://clang.llvm.org/docs/LanguageExtensions.html

[jhuber6]

It's in the NVPTXUsage.rst, and #81253 mentions it more directly.

[Artem-B]

I vaguely recall that OpenCL folks had to use it with a slightly different signature. I think their pointer argument was in an unusual address space, where OCL keeps their string constants. It would be great to double check that the new builtin does not break them.

[jhuber6]

It seems this broke the NVPTX tests as well. I'll try to check w/ OpenCL, but I know next to nothing about the OpenCL NVPTX support, I was surprised when I leanred that nxptx64-nvidia- can be anything other than CUDA.

[jhuber6]

I'm assuming that was why the anyptr argument was in the intrinsic definition? That caused the frontend to segfault if it was used so I figured it was incorrect.

[jhuber6]

Seems to work fine w/ OpenCL, added a test for it.

[Artem-B]

LGTM