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 <hube...@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; } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits