[clang] [NVPTX][clang] Ensure CLZ(0) is defined on NVPTX (PR #185630)

[Lewis Crawford via cfe-commits]

https://github.com/LewisCrawford updated https://github.com/llvm/llvm-project/pull/185630

>From 9b343b77de39504a62c6d1b4fd909e2b4f5c16fc Mon Sep 17 00:00:00 2001
From: Lewis Crawford <lcrawford at nvidia.com>
Date: Tue, 10 Mar 2026 11:56:50 +0000
Subject: [PATCH 1/2] [NVPTX][clang] Ensure CLZ(0) is defined on NVPTX

CUDA semantics specify that clz(0) = bitwidth, so
clang should emit clz / ctz intrinsics for NVPTX
with zero-is-poison = false.
---
 clang/lib/Basic/Targets/NVPTX.h                     |  2 ++
 clang/test/CodeGenCUDA/builtin-count-zeros-nvptx.cu | 12 ++++++++++++
 2 files changed, 14 insertions(+)
 create mode 100644 clang/test/CodeGenCUDA/builtin-count-zeros-nvptx.cu

diff --git a/clang/lib/Basic/Targets/NVPTX.h b/clang/lib/Basic/Targets/NVPTX.h
index 6f8df323f379c..7921a042e9e9b 100644
--- a/clang/lib/Basic/Targets/NVPTX.h
+++ b/clang/lib/Basic/Targets/NVPTX.h
@@ -83,6 +83,8 @@ class LLVM_LIBRARY_VISIBILITY NVPTXTargetInfo : public TargetInfo {
 
   bool useFP16ConversionIntrinsics() const override { return false; }
 
+  bool isCLZForZeroUndef() const override { return false; }
+
   bool
   initFeatureMap(llvm::StringMap<bool> &Features, DiagnosticsEngine &Diags,
                  StringRef CPU,
diff --git a/clang/test/CodeGenCUDA/builtin-count-zeros-nvptx.cu b/clang/test/CodeGenCUDA/builtin-count-zeros-nvptx.cu
new file mode 100644
index 0000000000000..f003b32ca73b2
--- /dev/null
+++ b/clang/test/CodeGenCUDA/builtin-count-zeros-nvptx.cu
@@ -0,0 +1,12 @@
+// REQUIRES: nvptx-registered-target
+// RUN: %clang_cc1 -x cuda -triple nvptx64-unknown-unknown -fcuda-is-device -emit-llvm %s -o - | FileCheck %s
+//
+// Ensure NVPTX uses isCLZForZeroUndef() = false (CUDA semantics: CLZ(i32 0) == 32).
+
+#include "Inputs/cuda.h"
+
+__device__ int f(int x) {
+  return __builtin_ctz(x) + __builtin_clz(x);
+}
+// CHECK: call i32 @llvm.cttz.i32({{.*}}, i1 false)
+// CHECK: call i32 @llvm.ctlz.i32({{.*}}, i1 false)

>From 0bddd8d61ac90c6d996a294fdbdd91243eeb0933 Mon Sep 17 00:00:00 2001
From: Lewis Crawford <lcrawford at nvidia.com>
Date: Tue, 10 Mar 2026 12:37:02 +0000
Subject: [PATCH 2/2] Update Headers/gpuintrin.c test

---
 clang/test/Headers/gpuintrin.c | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/clang/test/Headers/gpuintrin.c b/clang/test/Headers/gpuintrin.c
index 04b50acc4a049..c6a20dec210bb 100644
--- a/clang/test/Headers/gpuintrin.c
+++ b/clang/test/Headers/gpuintrin.c
@@ -1109,7 +1109,7 @@ __gpu_kernel void foo() {
 // NVPTX-NEXT:    br i1 [[TOBOOL]], label %[[COND_TRUE:.*]], label %[[COND_FALSE:.*]]
 // NVPTX:       [[COND_TRUE]]:
 // NVPTX-NEXT:    [[TMP3:%.*]] = load i64, ptr [[__BELOW]], align 8
-// NVPTX-NEXT:    [[TMP4:%.*]] = call i64 @llvm.ctlz.i64(i64 [[TMP3]], i1 true)
+// NVPTX-NEXT:    [[TMP4:%.*]] = call i64 @llvm.ctlz.i64(i64 [[TMP3]], i1 false)
 // NVPTX-NEXT:    [[CAST:%.*]] = trunc i64 [[TMP4]] to i32
 // NVPTX-NEXT:    [[SUB2:%.*]] = sub nsw i32 63, [[CAST]]
 // NVPTX-NEXT:    br label %[[COND_END:.*]]