[clang] [llvm] [CUDA][NVPTX] Add dp4a/dp2a builtins; define CUDA-13.4's __nv_idp*_impl externs (PR #233434)
Ilya Biryukov via cfe-commits
cfe-commits at lists.llvm.org
Wed Oct 21 03:32:18 PDT 2026
https://github.com/ilya-biryukov updated https://github.com/llvm/llvm-project/pull/233434
>From e7c93e88d02390460ad9bd28c76433d001bc779b Mon Sep 17 00:00:00 2001
From: Ilya Biryukov <ibiryukov at google.com>
Date: Mon, 19 Oct 2026 13:18:11 +0200
Subject: [PATCH 1/3] [CUDA][NVPTX] Add dp4a/dp2a builtins; define CUDA-13.4's
__nv_idp*_impl externs
Since CUDA 13.4 __dp4a/__dp2a_lo/__dp2a_hi from sm_61_intrinsics.hpp
call nvcc-internal device functions and are declared as plain extern "C":
__nv_idp4a_{s_s,u_u}_impl, __nv_idp2a_{lo,hi}_{s_s,u_u}_impl
They were implemented as inline assembly before.
As a result, Clang any use of __dp4a leaves an unresolved .extern in the
PTX and ptxas fails with
ptxas fatal : Unresolved extern function '__nv_idp4a_s_s_impl'
The NVPTX backend already has llvm.nvvm.idp4a.* / llvm.nvvm.idp2a.*. This
change:
* Exposes them as builtins: __nvvm_idp4a_{s,u}_{s,u} and __nvvm_idp2a_{lo,hi}_{s,u}_{s,u}
* Defines the six __nv_idp*_impl functions in __clang_cuda_intrinsics.h on
top of those builtins when CUDA_VERSION >= 13040, next to the existing
__nv_cvta_*_impl shims that serve the same purpose for CUDA 11.
Assited-By: Antigravity with Fable
---
clang/include/clang/Basic/BuiltinsNVPTX.td | 24 +++++++++
clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp | 24 +++++++++
clang/lib/Headers/__clang_cuda_intrinsics.h | 28 +++++++++++
clang/test/CodeGen/builtins-nvptx-sm_61.cu | 56 +++++++++++++++++++++
llvm/include/llvm/IR/IntrinsicsNVVM.td | 4 +-
5 files changed, 135 insertions(+), 1 deletion(-)
create mode 100644 clang/test/CodeGen/builtins-nvptx-sm_61.cu
diff --git a/clang/include/clang/Basic/BuiltinsNVPTX.td b/clang/include/clang/Basic/BuiltinsNVPTX.td
index 8915d89f01be38..cea7a6fe45b3b5 100644
--- a/clang/include/clang/Basic/BuiltinsNVPTX.td
+++ b/clang/include/clang/Basic/BuiltinsNVPTX.td
@@ -203,6 +203,30 @@ let Attributes = [NoThrow] in {
def __nvvm_pm_event_mask : NVPTXBuiltin<"void(_Constant unsigned short)">;
}
+// Integer dot product (dp4a / dp2a).
+//
+// __nvvm_idp4a_<a>_<b>(a, b, c): four-way byte dot product of a and b plus c.
+// __nvvm_idp2a_{lo,hi}_<a>_<b>(a, b, c): two-way dot product of the 16-bit
+// halves of a with the low/high byte pair of b, plus c.
+// <a>, <b> are the signedness ('s' or 'u') of the packed elements of a and b.
+// The result is signed unless both inputs are unsigned.
+let Attributes = [NoThrow] in {
+ def __nvvm_idp4a_s_s : NVPTXBuiltinSM<"int(int, int, int)", SM_61>;
+ def __nvvm_idp4a_s_u : NVPTXBuiltinSM<"int(int, unsigned int, int)", SM_61>;
+ def __nvvm_idp4a_u_s : NVPTXBuiltinSM<"int(unsigned int, int, int)", SM_61>;
+ def __nvvm_idp4a_u_u : NVPTXBuiltinSM<"unsigned int(unsigned int, unsigned int, unsigned int)", SM_61>;
+
+ def __nvvm_idp2a_lo_s_s : NVPTXBuiltinSM<"int(int, int, int)", SM_61>;
+ def __nvvm_idp2a_lo_s_u : NVPTXBuiltinSM<"int(int, unsigned int, int)", SM_61>;
+ def __nvvm_idp2a_lo_u_s : NVPTXBuiltinSM<"int(unsigned int, int, int)", SM_61>;
+ def __nvvm_idp2a_lo_u_u : NVPTXBuiltinSM<"unsigned int(unsigned int, unsigned int, unsigned int)", SM_61>;
+
+ def __nvvm_idp2a_hi_s_s : NVPTXBuiltinSM<"int(int, int, int)", SM_61>;
+ def __nvvm_idp2a_hi_s_u : NVPTXBuiltinSM<"int(int, unsigned int, int)", SM_61>;
+ def __nvvm_idp2a_hi_u_s : NVPTXBuiltinSM<"int(unsigned int, int, int)", SM_61>;
+ def __nvvm_idp2a_hi_u_u : NVPTXBuiltinSM<"unsigned int(unsigned int, unsigned int, unsigned int)", SM_61>;
+}
+
// Min Max
def __nvvm_fmin_f16 : NVPTXBuiltinSMAndPTX<"__fp16(__fp16, __fp16)", SM_80, PTX70>;
diff --git a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
index a513fcc12953ba..6ae564f17ac4ce 100644
--- a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
@@ -449,6 +449,14 @@ static Value *MakeFPArith(unsigned IntrinsicID, APFloat::roundingMode RM,
{CGF.Builder.getInt32(static_cast<int>(RM))});
}
+static Value *MakeIdp2a(unsigned IntrinsicID, bool IsHi, const CallExpr *E,
+ CodeGenFunction &CGF) {
+ return CGF.Builder.CreateCall(
+ CGF.CGM.getIntrinsic(IntrinsicID),
+ {CGF.EmitScalarExpr(E->getArg(0)), CGF.EmitScalarExpr(E->getArg(1)),
+ CGF.Builder.getInt1(IsHi), CGF.EmitScalarExpr(E->getArg(2))});
+}
+
} // namespace
Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID,
@@ -1587,6 +1595,22 @@ Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID,
Intrinsic::nvvm_mbarrier_init, {Ptr->getType()},
{Ptr, EmitScalarExpr(E->getArg(1)), Builder.getInt32(0)});
}
+ case NVPTX::BI__nvvm_idp2a_lo_s_s:
+ return MakeIdp2a(Intrinsic::nvvm_idp2a_s_s, /*IsHi=*/false, E, *this);
+ case NVPTX::BI__nvvm_idp2a_lo_s_u:
+ return MakeIdp2a(Intrinsic::nvvm_idp2a_s_u, /*IsHi=*/false, E, *this);
+ case NVPTX::BI__nvvm_idp2a_lo_u_s:
+ return MakeIdp2a(Intrinsic::nvvm_idp2a_u_s, /*IsHi=*/false, E, *this);
+ case NVPTX::BI__nvvm_idp2a_lo_u_u:
+ return MakeIdp2a(Intrinsic::nvvm_idp2a_u_u, /*IsHi=*/false, E, *this);
+ case NVPTX::BI__nvvm_idp2a_hi_s_s:
+ return MakeIdp2a(Intrinsic::nvvm_idp2a_s_s, /*IsHi=*/true, E, *this);
+ case NVPTX::BI__nvvm_idp2a_hi_s_u:
+ return MakeIdp2a(Intrinsic::nvvm_idp2a_s_u, /*IsHi=*/true, E, *this);
+ case NVPTX::BI__nvvm_idp2a_hi_u_s:
+ return MakeIdp2a(Intrinsic::nvvm_idp2a_u_s, /*IsHi=*/true, E, *this);
+ case NVPTX::BI__nvvm_idp2a_hi_u_u:
+ return MakeIdp2a(Intrinsic::nvvm_idp2a_u_u, /*IsHi=*/true, E, *this);
default:
return nullptr;
}
diff --git a/clang/lib/Headers/__clang_cuda_intrinsics.h b/clang/lib/Headers/__clang_cuda_intrinsics.h
index 2412f9e51ffbfb..11a54cbf29db98 100644
--- a/clang/lib/Headers/__clang_cuda_intrinsics.h
+++ b/clang/lib/Headers/__clang_cuda_intrinsics.h
@@ -843,6 +843,34 @@ __nv_associate_access_property(const void *__ptr, unsigned long long __prop) {
}
#endif // !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 800
+#if CUDA_VERSION >= 13040
+// Since CUDA 13.4 sm_61_intrinsics.hpp depends on `__nv_idp*_impl` being
+// available.
+extern "C" {
+__device__ inline int __nv_idp4a_s_s_impl(int __a, int __b, int __c) {
+ return __nvvm_idp4a_s_s(__a, __b, __c);
+}
+__device__ inline unsigned int
+__nv_idp4a_u_u_impl(unsigned int __a, unsigned int __b, unsigned int __c) {
+ return __nvvm_idp4a_u_u(__a, __b, __c);
+}
+__device__ inline int __nv_idp2a_lo_s_s_impl(int __a, int __b, int __c) {
+ return __nvvm_idp2a_lo_s_s(__a, __b, __c);
+}
+__device__ inline unsigned int
+__nv_idp2a_lo_u_u_impl(unsigned int __a, unsigned int __b, unsigned int __c) {
+ return __nvvm_idp2a_lo_u_u(__a, __b, __c);
+}
+__device__ inline int __nv_idp2a_hi_s_s_impl(int __a, int __b, int __c) {
+ return __nvvm_idp2a_hi_s_s(__a, __b, __c);
+}
+__device__ inline unsigned int
+__nv_idp2a_hi_u_u_impl(unsigned int __a, unsigned int __b, unsigned int __c) {
+ return __nvvm_idp2a_hi_u_u(__a, __b, __c);
+}
+} // extern "C"
+#endif // CUDA_VERSION >= 13040
+
#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 900
__device__ inline unsigned __isCtaShared(const void *ptr) {
return __isShared(ptr);
diff --git a/clang/test/CodeGen/builtins-nvptx-sm_61.cu b/clang/test/CodeGen/builtins-nvptx-sm_61.cu
new file mode 100644
index 00000000000000..3b5686c55327df
--- /dev/null
+++ b/clang/test/CodeGen/builtins-nvptx-sm_61.cu
@@ -0,0 +1,56 @@
+// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_61 \
+// RUN: -fcuda-is-device -target-feature +ptx62 \
+// RUN: -emit-llvm -o - -x cuda %s | FileCheck %s
+// Make sure the builtins still work with newer GPUs.
+// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_90 \
+// RUN: -fcuda-is-device -target-feature +ptx80 \
+// RUN: -emit-llvm -o - -x cuda %s | FileCheck %s
+// RUN: %clang_cc1 -triple nvptx-unknown-unknown -target-cpu sm_60 \
+// RUN: -fcuda-is-device -target-feature +ptx62 \
+// RUN: -S -o /dev/null -x cuda -verify %s
+
+#define __device__ __attribute__((device))
+
+// CHECK-LABEL: @_Z11dot_productPiiiijjj(
+__device__ void dot_product(int *out, int a, int b, int c, unsigned int ua,
+ unsigned int ub, unsigned int uc) {
+ // CHECK: call i32 @llvm.nvvm.idp4a.s.s(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})
+ // expected-error at +1 {{'__nvvm_idp4a_s_s' needs target feature sm_61}}
+ out[0] = __nvvm_idp4a_s_s(a, b, c);
+ // CHECK: call i32 @llvm.nvvm.idp4a.s.u(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})
+ // expected-error at +1 {{'__nvvm_idp4a_s_u' needs target feature sm_61}}
+ out[1] = __nvvm_idp4a_s_u(a, ub, c);
+ // CHECK: call i32 @llvm.nvvm.idp4a.u.s(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})
+ // expected-error at +1 {{'__nvvm_idp4a_u_s' needs target feature sm_61}}
+ out[2] = __nvvm_idp4a_u_s(ua, b, c);
+ // CHECK: call i32 @llvm.nvvm.idp4a.u.u(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})
+ // expected-error at +1 {{'__nvvm_idp4a_u_u' needs target feature sm_61}}
+ out[3] = __nvvm_idp4a_u_u(ua, ub, uc);
+
+ // CHECK: call i32 @llvm.nvvm.idp2a.s.s(i32 %{{.*}}, i32 %{{.*}}, i1 false, i32 %{{.*}})
+ // expected-error at +1 {{'__nvvm_idp2a_lo_s_s' needs target feature sm_61}}
+ out[4] = __nvvm_idp2a_lo_s_s(a, b, c);
+ // CHECK: call i32 @llvm.nvvm.idp2a.s.u(i32 %{{.*}}, i32 %{{.*}}, i1 false, i32 %{{.*}})
+ // expected-error at +1 {{'__nvvm_idp2a_lo_s_u' needs target feature sm_61}}
+ out[5] = __nvvm_idp2a_lo_s_u(a, ub, c);
+ // CHECK: call i32 @llvm.nvvm.idp2a.u.s(i32 %{{.*}}, i32 %{{.*}}, i1 false, i32 %{{.*}})
+ // expected-error at +1 {{'__nvvm_idp2a_lo_u_s' needs target feature sm_61}}
+ out[6] = __nvvm_idp2a_lo_u_s(ua, b, c);
+ // CHECK: call i32 @llvm.nvvm.idp2a.u.u(i32 %{{.*}}, i32 %{{.*}}, i1 false, i32 %{{.*}})
+ // expected-error at +1 {{'__nvvm_idp2a_lo_u_u' needs target feature sm_61}}
+ out[7] = __nvvm_idp2a_lo_u_u(ua, ub, uc);
+
+ // CHECK: call i32 @llvm.nvvm.idp2a.s.s(i32 %{{.*}}, i32 %{{.*}}, i1 true, i32 %{{.*}})
+ // expected-error at +1 {{'__nvvm_idp2a_hi_s_s' needs target feature sm_61}}
+ out[8] = __nvvm_idp2a_hi_s_s(a, b, c);
+ // CHECK: call i32 @llvm.nvvm.idp2a.s.u(i32 %{{.*}}, i32 %{{.*}}, i1 true, i32 %{{.*}})
+ // expected-error at +1 {{'__nvvm_idp2a_hi_s_u' needs target feature sm_61}}
+ out[9] = __nvvm_idp2a_hi_s_u(a, ub, c);
+ // CHECK: call i32 @llvm.nvvm.idp2a.u.s(i32 %{{.*}}, i32 %{{.*}}, i1 true, i32 %{{.*}})
+ // expected-error at +1 {{'__nvvm_idp2a_hi_u_s' needs target feature sm_61}}
+ out[10] = __nvvm_idp2a_hi_u_s(ua, b, c);
+ // CHECK: call i32 @llvm.nvvm.idp2a.u.u(i32 %{{.*}}, i32 %{{.*}}, i1 true, i32 %{{.*}})
+ // expected-error at +1 {{'__nvvm_idp2a_hi_u_u' needs target feature sm_61}}
+ out[11] = __nvvm_idp2a_hi_u_u(ua, ub, uc);
+ // CHECK: ret void
+}
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 4a274fcc73b17b..5bb9de28ff05c4 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -1939,9 +1939,11 @@ let TargetPrefix = "nvvm" in {
//
foreach a_type = ["s", "u"] in {
foreach b_type = ["s", "u"] in {
- def int_nvvm_idp4a_ # a_type # _ # b_type :
+ def int_nvvm_idp4a_ # a_type # _ # b_type : NVVMBuiltin,
PureIntrinsic<[llvm_i32_ty],
[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty]>;
+ // The i1 operand selects the high (true) or low (false) half of the
+ // second operand. Exposed to clang as separate _lo/_hi builtins.
def int_nvvm_idp2a_ # a_type # _ # b_type :
PureIntrinsic<[llvm_i32_ty],
[llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i32_ty],
>From 8dec773cfbd8ae06f5372e4553ca637839f39ea3 Mon Sep 17 00:00:00 2001
From: Ilya Biryukov <ibiryukov at google.com>
Date: Wed, 21 Oct 2026 10:52:07 +0200
Subject: [PATCH 2/3] fixup! [CUDA][NVPTX] Add dp4a/dp2a builtins; define
CUDA-13.4's __nv_idp*_impl externs
---
clang/docs/ReleaseNotes.md | 4 ++++
clang/include/clang/Basic/BuiltinsNVPTX.td | 7 -------
clang/test/CodeGen/builtins-nvptx-sm_61.cu | 10 +++++++++-
llvm/include/llvm/IR/IntrinsicsNVVM.td | 2 --
4 files changed, 13 insertions(+), 10 deletions(-)
diff --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md
index 4ebdfb8c692c4e..2a738d989cda9b 100644
--- a/clang/docs/ReleaseNotes.md
+++ b/clang/docs/ReleaseNotes.md
@@ -1209,6 +1209,10 @@ features cannot lower the translation-unit ABI level;
- External CUDA kernel declarations now use the correct calling convention. (#GH118212)
+- Fixed `__dp4a`, `__dp2a_lo` and `__dp2a_hi` producing unresolved extern
+ functions with CUDA 13.4 and newer. Added `__nvvm_idp4a_*` and
+ `__nvvm_idp2a_*` builtins. (#GH233434)
+
#### AIX Support
#### NetBSD Support
diff --git a/clang/include/clang/Basic/BuiltinsNVPTX.td b/clang/include/clang/Basic/BuiltinsNVPTX.td
index cea7a6fe45b3b5..904cd3c555c30b 100644
--- a/clang/include/clang/Basic/BuiltinsNVPTX.td
+++ b/clang/include/clang/Basic/BuiltinsNVPTX.td
@@ -203,13 +203,6 @@ let Attributes = [NoThrow] in {
def __nvvm_pm_event_mask : NVPTXBuiltin<"void(_Constant unsigned short)">;
}
-// Integer dot product (dp4a / dp2a).
-//
-// __nvvm_idp4a_<a>_<b>(a, b, c): four-way byte dot product of a and b plus c.
-// __nvvm_idp2a_{lo,hi}_<a>_<b>(a, b, c): two-way dot product of the 16-bit
-// halves of a with the low/high byte pair of b, plus c.
-// <a>, <b> are the signedness ('s' or 'u') of the packed elements of a and b.
-// The result is signed unless both inputs are unsigned.
let Attributes = [NoThrow] in {
def __nvvm_idp4a_s_s : NVPTXBuiltinSM<"int(int, int, int)", SM_61>;
def __nvvm_idp4a_s_u : NVPTXBuiltinSM<"int(int, unsigned int, int)", SM_61>;
diff --git a/clang/test/CodeGen/builtins-nvptx-sm_61.cu b/clang/test/CodeGen/builtins-nvptx-sm_61.cu
index 3b5686c55327df..007b574b02cefc 100644
--- a/clang/test/CodeGen/builtins-nvptx-sm_61.cu
+++ b/clang/test/CodeGen/builtins-nvptx-sm_61.cu
@@ -1,10 +1,18 @@
+// The builtin must work with sm_61 and newer architectures (we test on sm_90)
+// specifically. The two runs below look at CHECK lines and ignore
+// expected-error.
+//
// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_61 \
// RUN: -fcuda-is-device -target-feature +ptx62 \
// RUN: -emit-llvm -o - -x cuda %s | FileCheck %s
-// Make sure the builtins still work with newer GPUs.
// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_90 \
// RUN: -fcuda-is-device -target-feature +ptx80 \
// RUN: -emit-llvm -o - -x cuda %s | FileCheck %s
+//
+//
+// The builtin must fail on older GPUs, test on sm_60. The CHECK lines are
+// ignored by this run, we only check expected-error.
+//
// RUN: %clang_cc1 -triple nvptx-unknown-unknown -target-cpu sm_60 \
// RUN: -fcuda-is-device -target-feature +ptx62 \
// RUN: -S -o /dev/null -x cuda -verify %s
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 5bb9de28ff05c4..0c53587994a46a 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -1942,8 +1942,6 @@ let TargetPrefix = "nvvm" in {
def int_nvvm_idp4a_ # a_type # _ # b_type : NVVMBuiltin,
PureIntrinsic<[llvm_i32_ty],
[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty]>;
- // The i1 operand selects the high (true) or low (false) half of the
- // second operand. Exposed to clang as separate _lo/_hi builtins.
def int_nvvm_idp2a_ # a_type # _ # b_type :
PureIntrinsic<[llvm_i32_ty],
[llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i32_ty],
>From bb0dae69fd5f668262e5066f247e405e177db4e5 Mon Sep 17 00:00:00 2001
From: Ilya Biryukov <ibiryukov at google.com>
Date: Wed, 21 Oct 2026 12:29:50 +0200
Subject: [PATCH 3/3] fixup! [CUDA][NVPTX] Add dp4a/dp2a builtins; define
CUDA-13.4's __nv_idp*_impl externs
---
clang/test/CodeGen/builtins-nvptx-sm_61.cu | 9 ++++-----
1 file changed, 4 insertions(+), 5 deletions(-)
diff --git a/clang/test/CodeGen/builtins-nvptx-sm_61.cu b/clang/test/CodeGen/builtins-nvptx-sm_61.cu
index 007b574b02cefc..0796e1eb7a46a4 100644
--- a/clang/test/CodeGen/builtins-nvptx-sm_61.cu
+++ b/clang/test/CodeGen/builtins-nvptx-sm_61.cu
@@ -1,6 +1,5 @@
-// The builtin must work with sm_61 and newer architectures (we test on sm_90)
-// specifically. The two runs below look at CHECK lines and ignore
-// expected-error.
+// The builtins must work on sm_61 and newer (we also test on sm_90). The two
+// runs below only look at the CHECK lines.
//
// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_61 \
// RUN: -fcuda-is-device -target-feature +ptx62 \
@@ -10,8 +9,8 @@
// RUN: -emit-llvm -o - -x cuda %s | FileCheck %s
//
//
-// The builtin must fail on older GPUs, test on sm_60. The CHECK lines are
-// ignored by this run, we only check expected-error.
+// The builtins must be rejected on older GPUs, tested on sm_60 with -verify.
+// The CHECK lines are ignored by this run.
//
// RUN: %clang_cc1 -triple nvptx-unknown-unknown -target-cpu sm_60 \
// RUN: -fcuda-is-device -target-feature +ptx62 \
More information about the cfe-commits
mailing list