[llvm] [NVPTX] Add clmad instruction selection (PR #216412)
via llvm-commits
llvm-commits at lists.llvm.org
Fri Aug 14 14:52:00 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-nvptx
Author: robertvirany
<details>
<summary>Changes</summary>
Add instruction selection for the PTX 9.3 `clmad.{lo,hi}.u64` instructions on SM80 and newer targets.
Select `clmad` for `i64` carryless multiply-add expressions represented as an XOR of `ISD::CLMUL` or `ISD::CLMULH` with an accumulator. Standalone carryless multiplies are selected as `clmad` with an immediate zero-accumulator.
The first commit precommits tests showing the existing SelectionDAG expansion.
---
Full diff: https://github.com/llvm/llvm-project/pull/216412.diff
3 Files Affected:
- (modified) llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp (+2)
- (modified) llvm/lib/Target/NVPTX/NVPTXInstrInfo.td (+45)
- (added) llvm/test/CodeGen/NVPTX/clmad.ll (+79)
``````````diff
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index e788b0e44041f..ceafc51fa1925 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -711,6 +711,8 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM,
setOperationAction(ISD::SRA_PARTS, MVT::i64 , Custom);
setOperationAction(ISD::SRL_PARTS, MVT::i64 , Custom);
+ if (STI.hasFeature(NVPTX::PTX93) && STI.hasFeature(NVPTX::SM80))
+ setOperationAction({ISD::CLMUL, ISD::CLMULH}, MVT::i64, Legal);
setOperationAction(ISD::BITREVERSE, MVT::i32, Legal);
setOperationAction(ISD::BITREVERSE, MVT::i64, Legal);
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index f20c3ed9fedf5..9456fd935fd26 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -1163,6 +1163,51 @@ let Predicates = [hasOptEnabled] in {
}
}
+//
+// Carryless multiply-add
+//
+multiclass CLMADInst<string mode, SDPatternOperator op> {
+ def rrr:
+ BasicNVPTXInst<(outs I64RT.RC:$dst),
+ (ins I64RT.RC:$a, I64RT.RC:$b, I64RT.RC:$c),
+ "clmad." # mode # ".u64",
+ [(set I64RT.Ty:$dst, (xor
+ (OneUse2<op> I64RT.Ty:$a, I64RT.Ty:$b), I64RT.Ty:$c))]>;
+ def rri:
+ BasicNVPTXInst<(outs I64RT.RC:$dst),
+ (ins I64RT.RC:$a, I64RT.RC:$b, I64RT.Imm:$c),
+ "clmad." # mode # ".u64",
+ [(set I64RT.Ty:$dst, (xor
+ (OneUse2<op> I64RT.Ty:$a, I64RT.Ty:$b), I64RT.ImmNode:$c))]>;
+ def rir:
+ BasicNVPTXInst<(outs I64RT.RC:$dst),
+ (ins I64RT.RC:$a, I64RT.Imm:$b, I64RT.RC:$c),
+ "clmad." # mode # ".u64",
+ [(set I64RT.Ty:$dst, (xor
+ (OneUse2<op> I64RT.Ty:$a, I64RT.ImmNode:$b), I64RT.Ty:$c))]>;
+ def rii:
+ BasicNVPTXInst<(outs I64RT.RC:$dst),
+ (ins I64RT.RC:$a, I64RT.Imm:$b, I64RT.Imm:$c),
+ "clmad." # mode # ".u64",
+ [(set I64RT.Ty:$dst, (xor
+ (OneUse2<op> I64RT.Ty:$a, I64RT.ImmNode:$b), I64RT.ImmNode:$c))]>;
+}
+
+let Predicates = [PTX93, SM80] in {
+ defm CLMAD_HI : CLMADInst<"hi", clmulh>;
+ defm CLMAD_LO : CLMADInst<"lo", clmul>;
+
+ def : Pat<(clmul I64RT.Ty:$a, I64RT.Ty:$b),
+ (CLMAD_LOrri $a, $b, 0)>;
+ def : Pat<(clmulh I64RT.Ty:$a, I64RT.Ty:$b),
+ (CLMAD_HIrri $a, $b, 0)>;
+
+ def : Pat<(clmul I64RT.Ty:$a, I64RT.ImmNode:$b),
+ (CLMAD_LOrii $a, $b, 0)>;
+ def : Pat<(clmulh I64RT.Ty:$a, I64RT.ImmNode:$b),
+ (CLMAD_HIrii $a, $b, 0)>;
+}
+
//-----------------------------------
// Floating Point Arithmetic
//-----------------------------------
diff --git a/llvm/test/CodeGen/NVPTX/clmad.ll b/llvm/test/CodeGen/NVPTX/clmad.ll
new file mode 100644
index 0000000000000..4d8699f982b1b
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/clmad.ll
@@ -0,0 +1,79 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx93 | FileCheck %s
+; RUN: %if ptxas-sm_100 && ptxas-isa-9.3 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx93 | %ptxas-verify -arch=sm_100 %}
+target triple = "nvptx64-nvidia-cuda"
+
+define i64 @clmad_lo(i64 %a, i64 %b, i64 %c) {
+; CHECK-LABEL: clmad_lo(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [clmad_lo_param_0];
+; CHECK-NEXT: ld.param::func.b64 %rd2, [clmad_lo_param_1];
+; CHECK-NEXT: ld.param::func.b64 %rd3, [clmad_lo_param_2];
+; CHECK-NEXT: clmad.lo.u64 %rd4, %rd1, %rd2, %rd3;
+; CHECK-NEXT: st.param::func.b64 [func_retval0], %rd4;
+; CHECK-NEXT: ret;
+ %product = call i64 @llvm.clmul.i64(i64 %a, i64 %b)
+ %result = xor i64 %product, %c
+ ret i64 %result
+}
+
+define i64 @clmad_hi(i64 %a, i64 %b, i64 %c) {
+; CHECK-LABEL: clmad_hi(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [clmad_hi_param_0];
+; CHECK-NEXT: ld.param::func.b64 %rd2, [clmad_hi_param_1];
+; CHECK-NEXT: ld.param::func.b64 %rd3, [clmad_hi_param_2];
+; CHECK-NEXT: clmad.hi.u64 %rd4, %rd1, %rd2, %rd3;
+; CHECK-NEXT: st.param::func.b64 [func_retval0], %rd4;
+; CHECK-NEXT: ret;
+ %a.wide = zext i64 %a to i128
+ %b.wide = zext i64 %b to i128
+ %product = call i128 @llvm.clmul.i128(i128 %a.wide, i128 %b.wide)
+ %shifted = lshr i128 %product, 64
+ %high = trunc i128 %shifted to i64
+ %result = xor i64 %high, %c
+ ret i64 %result
+}
+
+define i64 @clmul_lo(i64 %a, i64 %b) {
+; CHECK-LABEL: clmul_lo(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [clmul_lo_param_0];
+; CHECK-NEXT: ld.param::func.b64 %rd2, [clmul_lo_param_1];
+; CHECK-NEXT: clmad.lo.u64 %rd3, %rd1, %rd2, 0;
+; CHECK-NEXT: st.param::func.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ret;
+ %product = call i64 @llvm.clmul.i64(i64 %a, i64 %b)
+ ret i64 %product
+}
+
+define i64 @clmul_hi(i64 %a, i64 %b) {
+; CHECK-LABEL: clmul_hi(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [clmul_hi_param_0];
+; CHECK-NEXT: ld.param::func.b64 %rd2, [clmul_hi_param_1];
+; CHECK-NEXT: clmad.hi.u64 %rd3, %rd1, %rd2, 0;
+; CHECK-NEXT: st.param::func.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ret;
+ %a.wide = zext i64 %a to i128
+ %b.wide = zext i64 %b to i128
+ %product = call i128 @llvm.clmul.i128(i128 %a.wide, i128 %b.wide)
+ %shifted = lshr i128 %product, 64
+ %high = trunc i128 %shifted to i64
+ ret i64 %high
+}
+
+declare i64 @llvm.clmul.i64(i64, i64)
+declare i128 @llvm.clmul.i128(i128, i128)
``````````
</details>
https://github.com/llvm/llvm-project/pull/216412
More information about the llvm-commits
mailing list