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 <[email protected]> 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@+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@+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@+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@+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@+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@+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@+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@+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@+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@+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@+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@+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 <[email protected]> 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 <[email protected]> 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 \ _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
