[clang] [NVPTX] Add a clang builtin for the `warpsize` intrinsic (PR #110316)

[Joseph Huber via cfe-commits]

https://github.com/jhuber6 created https://github.com/llvm/llvm-project/pull/110316

Summary:
There's an intrinsic for the warp size, we want to expose this to make
the interface proposed in
https://github.com/llvm/llvm-project/pull/110179 more generic.


>From 63d45843ee15c940680e4d6a3ea87138ebfc5b69 Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Fri, 27 Sep 2024 14:08:51 -0500
Subject: [PATCH] [NVPTX] Add a clang builtin for the `warpsize` intrinsic

Summary:
There's an intrinsic for the warp size, we want to expose this to make
the interface proposed in
https://github.com/llvm/llvm-project/pull/110179 more generic.
---
 clang/include/clang/Basic/BuiltinsNVPTX.def | 1 +
 clang/test/CodeGen/builtins-nvptx.c         | 4 +++-
 2 files changed, 4 insertions(+), 1 deletion(-)

diff --git a/clang/include/clang/Basic/BuiltinsNVPTX.def b/clang/include/clang/Basic/BuiltinsNVPTX.def
index 6fff562165080a..6b7bce5bc00d4f 100644
--- a/clang/include/clang/Basic/BuiltinsNVPTX.def
+++ b/clang/include/clang/Basic/BuiltinsNVPTX.def
@@ -139,6 +139,7 @@ TARGET_BUILTIN(__nvvm_is_explicit_cluster, "b", "nc", AND(SM_90, PTX78))
 BUILTIN(__nvvm_read_ptx_sreg_laneid, "i", "nc")
 BUILTIN(__nvvm_read_ptx_sreg_warpid, "i", "nc")
 BUILTIN(__nvvm_read_ptx_sreg_nwarpid, "i", "nc")
+BUILTIN(__nvvm_read_ptx_sreg_warpsize, "i", "nc")
 
 BUILTIN(__nvvm_read_ptx_sreg_smid, "i", "nc")
 BUILTIN(__nvvm_read_ptx_sreg_nsmid, "i", "nc")
diff --git a/clang/test/CodeGen/builtins-nvptx.c b/clang/test/CodeGen/builtins-nvptx.c
index bfa72e8bd69454..0d0e3ecdb90c9e 100644
--- a/clang/test/CodeGen/builtins-nvptx.c
+++ b/clang/test/CodeGen/builtins-nvptx.c
@@ -114,6 +114,7 @@ __device__ int read_ids() {
 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.smid()
 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nsmid()
 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.gridid()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
 
   int a = __nvvm_read_ptx_sreg_laneid();
   int b = __nvvm_read_ptx_sreg_warpid();
@@ -121,8 +122,9 @@ __device__ int read_ids() {
   int d = __nvvm_read_ptx_sreg_smid();
   int e = __nvvm_read_ptx_sreg_nsmid();
   int f = __nvvm_read_ptx_sreg_gridid();
+  int g = __nvvm_read_ptx_sreg_warpsize();
 
-  return a + b + c + d + e + f;
+  return a + b + c + d + e + f + g;
 
 }