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

[Lewis Crawford via cfe-commits]

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

CUDA semantics specify that clz(0) = bitwidth, so clang should emit clz / ctz intrinsics for NVPTX with zero-is-poison = false.

>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] [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)