llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-nvptx
Author: Ilya Biryukov (ilya-biryukov)
<details>
<summary>Changes</summary>
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"` in the
corresponding header, namely ` __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 leaves any use of `__dp4a` as 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 `__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
---
Full diff: https://github.com/llvm/llvm-project/pull/233434.diff
5 Files Affected:
- (modified) clang/include/clang/Basic/BuiltinsNVPTX.td (+24)
- (modified) clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp (+24)
- (modified) clang/lib/Headers/__clang_cuda_intrinsics.h (+28)
- (added) clang/test/CodeGen/builtins-nvptx-sm_61.cu (+56)
- (modified) llvm/include/llvm/IR/IntrinsicsNVVM.td (+3-1)
``````````diff
diff --git a/clang/include/clang/Basic/BuiltinsNVPTX.td
b/clang/include/clang/Basic/BuiltinsNVPTX.td
index 8915d89f01be3..cea7a6fe45b3b 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 a513fcc12953b..6ae564f17ac4c 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 2412f9e51ffbf..11a54cbf29db9 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 0000000000000..3b5686c55327d
--- /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 4a274fcc73b17..5bb9de28ff05c 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],
``````````
</details>
https://github.com/llvm/llvm-project/pull/233434
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
