[llvm] [NVPTX] Add clmad instruction selection (PR #216412)
via llvm-commits
llvm-commits at lists.llvm.org
Fri Aug 14 14:51:20 PDT 2026
https://github.com/robertvirany created https://github.com/llvm/llvm-project/pull/216412
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.
>From 75c8d845a5b001908ce09311c4441cc98b0665c8 Mon Sep 17 00:00:00 2001
From: Robert Virany <robertvirany at gmail.com>
Date: Fri, 14 Aug 2026 11:41:09 -0600
Subject: [PATCH 1/2] [NVPTX] Precommit tests for clmad code generation
---
llvm/test/CodeGen/NVPTX/clmad.ll | 257 +++++++++++++++++++++++++++++++
1 file changed, 257 insertions(+)
create mode 100644 llvm/test/CodeGen/NVPTX/clmad.ll
diff --git a/llvm/test/CodeGen/NVPTX/clmad.ll b/llvm/test/CodeGen/NVPTX/clmad.ll
new file mode 100644
index 0000000000000..a8099936b36b3
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/clmad.ll
@@ -0,0 +1,257 @@
+; 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<48>;
+; 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: and.b64 %rd3, %rd2, 1229782938247303441;
+; CHECK-NEXT: and.b64 %rd4, %rd1, 2459565876494606882;
+; CHECK-NEXT: mul.lo.s64 %rd5, %rd4, %rd3;
+; CHECK-NEXT: and.b64 %rd6, %rd2, 2459565876494606882;
+; CHECK-NEXT: and.b64 %rd7, %rd1, 1229782938247303441;
+; CHECK-NEXT: mul.lo.s64 %rd8, %rd7, %rd6;
+; CHECK-NEXT: xor.b64 %rd9, %rd8, %rd5;
+; CHECK-NEXT: and.b64 %rd10, %rd2, -8608480567731124088;
+; CHECK-NEXT: and.b64 %rd11, %rd1, 4919131752989213764;
+; CHECK-NEXT: mul.lo.s64 %rd12, %rd11, %rd10;
+; CHECK-NEXT: xor.b64 %rd13, %rd9, %rd12;
+; CHECK-NEXT: and.b64 %rd14, %rd2, 4919131752989213764;
+; CHECK-NEXT: and.b64 %rd15, %rd1, -8608480567731124088;
+; CHECK-NEXT: mul.lo.s64 %rd16, %rd15, %rd14;
+; CHECK-NEXT: xor.b64 %rd17, %rd13, %rd16;
+; CHECK-NEXT: and.b64 %rd18, %rd17, 2459565876494606882;
+; CHECK-NEXT: mul.lo.s64 %rd19, %rd4, %rd10;
+; CHECK-NEXT: mul.lo.s64 %rd20, %rd7, %rd3;
+; CHECK-NEXT: xor.b64 %rd21, %rd20, %rd19;
+; CHECK-NEXT: mul.lo.s64 %rd22, %rd11, %rd14;
+; CHECK-NEXT: xor.b64 %rd23, %rd21, %rd22;
+; CHECK-NEXT: mul.lo.s64 %rd24, %rd15, %rd6;
+; CHECK-NEXT: xor.b64 %rd25, %rd23, %rd24;
+; CHECK-NEXT: and.b64 %rd26, %rd25, 1229782938247303441;
+; CHECK-NEXT: or.b64 %rd27, %rd26, %rd18;
+; CHECK-NEXT: mul.lo.s64 %rd28, %rd4, %rd6;
+; CHECK-NEXT: mul.lo.s64 %rd29, %rd7, %rd14;
+; CHECK-NEXT: xor.b64 %rd30, %rd29, %rd28;
+; CHECK-NEXT: mul.lo.s64 %rd31, %rd11, %rd3;
+; CHECK-NEXT: xor.b64 %rd32, %rd30, %rd31;
+; CHECK-NEXT: mul.lo.s64 %rd33, %rd15, %rd10;
+; CHECK-NEXT: xor.b64 %rd34, %rd32, %rd33;
+; CHECK-NEXT: and.b64 %rd35, %rd34, 4919131752989213764;
+; CHECK-NEXT: or.b64 %rd36, %rd27, %rd35;
+; CHECK-NEXT: mul.lo.s64 %rd37, %rd4, %rd14;
+; CHECK-NEXT: mul.lo.s64 %rd38, %rd7, %rd10;
+; CHECK-NEXT: xor.b64 %rd39, %rd38, %rd37;
+; CHECK-NEXT: mul.lo.s64 %rd40, %rd11, %rd6;
+; CHECK-NEXT: xor.b64 %rd41, %rd39, %rd40;
+; CHECK-NEXT: mul.lo.s64 %rd42, %rd15, %rd3;
+; CHECK-NEXT: xor.b64 %rd43, %rd41, %rd42;
+; CHECK-NEXT: and.b64 %rd44, %rd43, -8608480567731124088;
+; CHECK-NEXT: or.b64 %rd45, %rd36, %rd44;
+; CHECK-NEXT: ld.param::func.b64 %rd46, [clmad_lo_param_2];
+; CHECK-NEXT: xor.b64 %rd47, %rd45, %rd46;
+; CHECK-NEXT: st.param::func.b64 [func_retval0], %rd47;
+; 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<52>;
+; 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: brev.b64 %rd4, %rd2;
+; CHECK-NEXT: and.b64 %rd5, %rd4, 1229782938247303441;
+; CHECK-NEXT: brev.b64 %rd6, %rd1;
+; CHECK-NEXT: and.b64 %rd7, %rd6, 2459565876494606882;
+; CHECK-NEXT: mul.lo.s64 %rd8, %rd7, %rd5;
+; CHECK-NEXT: and.b64 %rd9, %rd4, 2459565876494606882;
+; CHECK-NEXT: and.b64 %rd10, %rd6, 1229782938247303441;
+; CHECK-NEXT: mul.lo.s64 %rd11, %rd10, %rd9;
+; CHECK-NEXT: xor.b64 %rd12, %rd11, %rd8;
+; CHECK-NEXT: and.b64 %rd13, %rd4, -8608480567731124088;
+; CHECK-NEXT: and.b64 %rd14, %rd6, 4919131752989213764;
+; CHECK-NEXT: mul.lo.s64 %rd15, %rd14, %rd13;
+; CHECK-NEXT: xor.b64 %rd16, %rd12, %rd15;
+; CHECK-NEXT: and.b64 %rd17, %rd4, 4919131752989213764;
+; CHECK-NEXT: and.b64 %rd18, %rd6, -8608480567731124088;
+; CHECK-NEXT: mul.lo.s64 %rd19, %rd18, %rd17;
+; CHECK-NEXT: xor.b64 %rd20, %rd16, %rd19;
+; CHECK-NEXT: and.b64 %rd21, %rd20, 2459565876494606882;
+; CHECK-NEXT: mul.lo.s64 %rd22, %rd7, %rd13;
+; CHECK-NEXT: mul.lo.s64 %rd23, %rd10, %rd5;
+; CHECK-NEXT: xor.b64 %rd24, %rd23, %rd22;
+; CHECK-NEXT: mul.lo.s64 %rd25, %rd14, %rd17;
+; CHECK-NEXT: xor.b64 %rd26, %rd24, %rd25;
+; CHECK-NEXT: mul.lo.s64 %rd27, %rd18, %rd9;
+; CHECK-NEXT: xor.b64 %rd28, %rd26, %rd27;
+; CHECK-NEXT: and.b64 %rd29, %rd28, 1229782938247303441;
+; CHECK-NEXT: or.b64 %rd30, %rd29, %rd21;
+; CHECK-NEXT: mul.lo.s64 %rd31, %rd7, %rd9;
+; CHECK-NEXT: mul.lo.s64 %rd32, %rd10, %rd17;
+; CHECK-NEXT: xor.b64 %rd33, %rd32, %rd31;
+; CHECK-NEXT: mul.lo.s64 %rd34, %rd14, %rd5;
+; CHECK-NEXT: xor.b64 %rd35, %rd33, %rd34;
+; CHECK-NEXT: mul.lo.s64 %rd36, %rd18, %rd13;
+; CHECK-NEXT: xor.b64 %rd37, %rd35, %rd36;
+; CHECK-NEXT: and.b64 %rd38, %rd37, 4919131752989213764;
+; CHECK-NEXT: or.b64 %rd39, %rd30, %rd38;
+; CHECK-NEXT: mul.lo.s64 %rd40, %rd7, %rd17;
+; CHECK-NEXT: mul.lo.s64 %rd41, %rd10, %rd13;
+; CHECK-NEXT: xor.b64 %rd42, %rd41, %rd40;
+; CHECK-NEXT: mul.lo.s64 %rd43, %rd14, %rd9;
+; CHECK-NEXT: xor.b64 %rd44, %rd42, %rd43;
+; CHECK-NEXT: mul.lo.s64 %rd45, %rd18, %rd5;
+; CHECK-NEXT: xor.b64 %rd46, %rd44, %rd45;
+; CHECK-NEXT: and.b64 %rd47, %rd46, 614891469123651720;
+; CHECK-NEXT: or.b64 %rd48, %rd39, %rd47;
+; CHECK-NEXT: brev.b64 %rd49, %rd48;
+; CHECK-NEXT: shr.u64 %rd50, %rd49, 1;
+; CHECK-NEXT: xor.b64 %rd51, %rd50, %rd3;
+; CHECK-NEXT: st.param::func.b64 [func_retval0], %rd51;
+; 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<46>;
+; 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: and.b64 %rd3, %rd2, 1229782938247303441;
+; CHECK-NEXT: and.b64 %rd4, %rd1, 2459565876494606882;
+; CHECK-NEXT: mul.lo.s64 %rd5, %rd4, %rd3;
+; CHECK-NEXT: and.b64 %rd6, %rd2, 2459565876494606882;
+; CHECK-NEXT: and.b64 %rd7, %rd1, 1229782938247303441;
+; CHECK-NEXT: mul.lo.s64 %rd8, %rd7, %rd6;
+; CHECK-NEXT: xor.b64 %rd9, %rd8, %rd5;
+; CHECK-NEXT: and.b64 %rd10, %rd2, -8608480567731124088;
+; CHECK-NEXT: and.b64 %rd11, %rd1, 4919131752989213764;
+; CHECK-NEXT: mul.lo.s64 %rd12, %rd11, %rd10;
+; CHECK-NEXT: xor.b64 %rd13, %rd9, %rd12;
+; CHECK-NEXT: and.b64 %rd14, %rd2, 4919131752989213764;
+; CHECK-NEXT: and.b64 %rd15, %rd1, -8608480567731124088;
+; CHECK-NEXT: mul.lo.s64 %rd16, %rd15, %rd14;
+; CHECK-NEXT: xor.b64 %rd17, %rd13, %rd16;
+; CHECK-NEXT: and.b64 %rd18, %rd17, 2459565876494606882;
+; CHECK-NEXT: mul.lo.s64 %rd19, %rd4, %rd10;
+; CHECK-NEXT: mul.lo.s64 %rd20, %rd7, %rd3;
+; CHECK-NEXT: xor.b64 %rd21, %rd20, %rd19;
+; CHECK-NEXT: mul.lo.s64 %rd22, %rd11, %rd14;
+; CHECK-NEXT: xor.b64 %rd23, %rd21, %rd22;
+; CHECK-NEXT: mul.lo.s64 %rd24, %rd15, %rd6;
+; CHECK-NEXT: xor.b64 %rd25, %rd23, %rd24;
+; CHECK-NEXT: and.b64 %rd26, %rd25, 1229782938247303441;
+; CHECK-NEXT: or.b64 %rd27, %rd26, %rd18;
+; CHECK-NEXT: mul.lo.s64 %rd28, %rd4, %rd6;
+; CHECK-NEXT: mul.lo.s64 %rd29, %rd7, %rd14;
+; CHECK-NEXT: xor.b64 %rd30, %rd29, %rd28;
+; CHECK-NEXT: mul.lo.s64 %rd31, %rd11, %rd3;
+; CHECK-NEXT: xor.b64 %rd32, %rd30, %rd31;
+; CHECK-NEXT: mul.lo.s64 %rd33, %rd15, %rd10;
+; CHECK-NEXT: xor.b64 %rd34, %rd32, %rd33;
+; CHECK-NEXT: and.b64 %rd35, %rd34, 4919131752989213764;
+; CHECK-NEXT: or.b64 %rd36, %rd27, %rd35;
+; CHECK-NEXT: mul.lo.s64 %rd37, %rd4, %rd14;
+; CHECK-NEXT: mul.lo.s64 %rd38, %rd7, %rd10;
+; CHECK-NEXT: xor.b64 %rd39, %rd38, %rd37;
+; CHECK-NEXT: mul.lo.s64 %rd40, %rd11, %rd6;
+; CHECK-NEXT: xor.b64 %rd41, %rd39, %rd40;
+; CHECK-NEXT: mul.lo.s64 %rd42, %rd15, %rd3;
+; CHECK-NEXT: xor.b64 %rd43, %rd41, %rd42;
+; CHECK-NEXT: and.b64 %rd44, %rd43, -8608480567731124088;
+; CHECK-NEXT: or.b64 %rd45, %rd36, %rd44;
+; CHECK-NEXT: st.param::func.b64 [func_retval0], %rd45;
+; 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<50>;
+; 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: brev.b64 %rd3, %rd2;
+; CHECK-NEXT: and.b64 %rd4, %rd3, 1229782938247303441;
+; CHECK-NEXT: brev.b64 %rd5, %rd1;
+; CHECK-NEXT: and.b64 %rd6, %rd5, 2459565876494606882;
+; CHECK-NEXT: mul.lo.s64 %rd7, %rd6, %rd4;
+; CHECK-NEXT: and.b64 %rd8, %rd3, 2459565876494606882;
+; CHECK-NEXT: and.b64 %rd9, %rd5, 1229782938247303441;
+; CHECK-NEXT: mul.lo.s64 %rd10, %rd9, %rd8;
+; CHECK-NEXT: xor.b64 %rd11, %rd10, %rd7;
+; CHECK-NEXT: and.b64 %rd12, %rd3, -8608480567731124088;
+; CHECK-NEXT: and.b64 %rd13, %rd5, 4919131752989213764;
+; CHECK-NEXT: mul.lo.s64 %rd14, %rd13, %rd12;
+; CHECK-NEXT: xor.b64 %rd15, %rd11, %rd14;
+; CHECK-NEXT: and.b64 %rd16, %rd3, 4919131752989213764;
+; CHECK-NEXT: and.b64 %rd17, %rd5, -8608480567731124088;
+; CHECK-NEXT: mul.lo.s64 %rd18, %rd17, %rd16;
+; CHECK-NEXT: xor.b64 %rd19, %rd15, %rd18;
+; CHECK-NEXT: and.b64 %rd20, %rd19, 2459565876494606882;
+; CHECK-NEXT: mul.lo.s64 %rd21, %rd6, %rd12;
+; CHECK-NEXT: mul.lo.s64 %rd22, %rd9, %rd4;
+; CHECK-NEXT: xor.b64 %rd23, %rd22, %rd21;
+; CHECK-NEXT: mul.lo.s64 %rd24, %rd13, %rd16;
+; CHECK-NEXT: xor.b64 %rd25, %rd23, %rd24;
+; CHECK-NEXT: mul.lo.s64 %rd26, %rd17, %rd8;
+; CHECK-NEXT: xor.b64 %rd27, %rd25, %rd26;
+; CHECK-NEXT: and.b64 %rd28, %rd27, 1229782938247303441;
+; CHECK-NEXT: or.b64 %rd29, %rd28, %rd20;
+; CHECK-NEXT: mul.lo.s64 %rd30, %rd6, %rd8;
+; CHECK-NEXT: mul.lo.s64 %rd31, %rd9, %rd16;
+; CHECK-NEXT: xor.b64 %rd32, %rd31, %rd30;
+; CHECK-NEXT: mul.lo.s64 %rd33, %rd13, %rd4;
+; CHECK-NEXT: xor.b64 %rd34, %rd32, %rd33;
+; CHECK-NEXT: mul.lo.s64 %rd35, %rd17, %rd12;
+; CHECK-NEXT: xor.b64 %rd36, %rd34, %rd35;
+; CHECK-NEXT: and.b64 %rd37, %rd36, 4919131752989213764;
+; CHECK-NEXT: or.b64 %rd38, %rd29, %rd37;
+; CHECK-NEXT: mul.lo.s64 %rd39, %rd6, %rd16;
+; CHECK-NEXT: mul.lo.s64 %rd40, %rd9, %rd12;
+; CHECK-NEXT: xor.b64 %rd41, %rd40, %rd39;
+; CHECK-NEXT: mul.lo.s64 %rd42, %rd13, %rd8;
+; CHECK-NEXT: xor.b64 %rd43, %rd41, %rd42;
+; CHECK-NEXT: mul.lo.s64 %rd44, %rd17, %rd4;
+; CHECK-NEXT: xor.b64 %rd45, %rd43, %rd44;
+; CHECK-NEXT: and.b64 %rd46, %rd45, 614891469123651720;
+; CHECK-NEXT: or.b64 %rd47, %rd38, %rd46;
+; CHECK-NEXT: brev.b64 %rd48, %rd47;
+; CHECK-NEXT: shr.u64 %rd49, %rd48, 1;
+; CHECK-NEXT: st.param::func.b64 [func_retval0], %rd49;
+; 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)
>From 8147575a5c829ad04cf3b2cd07ee5cda1630a734 Mon Sep 17 00:00:00 2001
From: Robert Virany <robertvirany at gmail.com>
Date: Fri, 14 Aug 2026 15:33:01 -0600
Subject: [PATCH 2/2] [NVPTX] Select clmad for carryless multiply-add
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 2 +
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 45 +++++
llvm/test/CodeGen/NVPTX/clmad.ll | 204 ++------------------
3 files changed, 60 insertions(+), 191 deletions(-)
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
index a8099936b36b3..4d8699f982b1b 100644
--- a/llvm/test/CodeGen/NVPTX/clmad.ll
+++ b/llvm/test/CodeGen/NVPTX/clmad.ll
@@ -6,57 +6,14 @@ 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<48>;
+; 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: and.b64 %rd3, %rd2, 1229782938247303441;
-; CHECK-NEXT: and.b64 %rd4, %rd1, 2459565876494606882;
-; CHECK-NEXT: mul.lo.s64 %rd5, %rd4, %rd3;
-; CHECK-NEXT: and.b64 %rd6, %rd2, 2459565876494606882;
-; CHECK-NEXT: and.b64 %rd7, %rd1, 1229782938247303441;
-; CHECK-NEXT: mul.lo.s64 %rd8, %rd7, %rd6;
-; CHECK-NEXT: xor.b64 %rd9, %rd8, %rd5;
-; CHECK-NEXT: and.b64 %rd10, %rd2, -8608480567731124088;
-; CHECK-NEXT: and.b64 %rd11, %rd1, 4919131752989213764;
-; CHECK-NEXT: mul.lo.s64 %rd12, %rd11, %rd10;
-; CHECK-NEXT: xor.b64 %rd13, %rd9, %rd12;
-; CHECK-NEXT: and.b64 %rd14, %rd2, 4919131752989213764;
-; CHECK-NEXT: and.b64 %rd15, %rd1, -8608480567731124088;
-; CHECK-NEXT: mul.lo.s64 %rd16, %rd15, %rd14;
-; CHECK-NEXT: xor.b64 %rd17, %rd13, %rd16;
-; CHECK-NEXT: and.b64 %rd18, %rd17, 2459565876494606882;
-; CHECK-NEXT: mul.lo.s64 %rd19, %rd4, %rd10;
-; CHECK-NEXT: mul.lo.s64 %rd20, %rd7, %rd3;
-; CHECK-NEXT: xor.b64 %rd21, %rd20, %rd19;
-; CHECK-NEXT: mul.lo.s64 %rd22, %rd11, %rd14;
-; CHECK-NEXT: xor.b64 %rd23, %rd21, %rd22;
-; CHECK-NEXT: mul.lo.s64 %rd24, %rd15, %rd6;
-; CHECK-NEXT: xor.b64 %rd25, %rd23, %rd24;
-; CHECK-NEXT: and.b64 %rd26, %rd25, 1229782938247303441;
-; CHECK-NEXT: or.b64 %rd27, %rd26, %rd18;
-; CHECK-NEXT: mul.lo.s64 %rd28, %rd4, %rd6;
-; CHECK-NEXT: mul.lo.s64 %rd29, %rd7, %rd14;
-; CHECK-NEXT: xor.b64 %rd30, %rd29, %rd28;
-; CHECK-NEXT: mul.lo.s64 %rd31, %rd11, %rd3;
-; CHECK-NEXT: xor.b64 %rd32, %rd30, %rd31;
-; CHECK-NEXT: mul.lo.s64 %rd33, %rd15, %rd10;
-; CHECK-NEXT: xor.b64 %rd34, %rd32, %rd33;
-; CHECK-NEXT: and.b64 %rd35, %rd34, 4919131752989213764;
-; CHECK-NEXT: or.b64 %rd36, %rd27, %rd35;
-; CHECK-NEXT: mul.lo.s64 %rd37, %rd4, %rd14;
-; CHECK-NEXT: mul.lo.s64 %rd38, %rd7, %rd10;
-; CHECK-NEXT: xor.b64 %rd39, %rd38, %rd37;
-; CHECK-NEXT: mul.lo.s64 %rd40, %rd11, %rd6;
-; CHECK-NEXT: xor.b64 %rd41, %rd39, %rd40;
-; CHECK-NEXT: mul.lo.s64 %rd42, %rd15, %rd3;
-; CHECK-NEXT: xor.b64 %rd43, %rd41, %rd42;
-; CHECK-NEXT: and.b64 %rd44, %rd43, -8608480567731124088;
-; CHECK-NEXT: or.b64 %rd45, %rd36, %rd44;
-; CHECK-NEXT: ld.param::func.b64 %rd46, [clmad_lo_param_2];
-; CHECK-NEXT: xor.b64 %rd47, %rd45, %rd46;
-; CHECK-NEXT: st.param::func.b64 [func_retval0], %rd47;
+; 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
@@ -66,61 +23,14 @@ define i64 @clmad_lo(i64 %a, i64 %b, i64 %c) {
define i64 @clmad_hi(i64 %a, i64 %b, i64 %c) {
; CHECK-LABEL: clmad_hi(
; CHECK: {
-; CHECK-NEXT: .reg .b64 %rd<52>;
+; 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: brev.b64 %rd4, %rd2;
-; CHECK-NEXT: and.b64 %rd5, %rd4, 1229782938247303441;
-; CHECK-NEXT: brev.b64 %rd6, %rd1;
-; CHECK-NEXT: and.b64 %rd7, %rd6, 2459565876494606882;
-; CHECK-NEXT: mul.lo.s64 %rd8, %rd7, %rd5;
-; CHECK-NEXT: and.b64 %rd9, %rd4, 2459565876494606882;
-; CHECK-NEXT: and.b64 %rd10, %rd6, 1229782938247303441;
-; CHECK-NEXT: mul.lo.s64 %rd11, %rd10, %rd9;
-; CHECK-NEXT: xor.b64 %rd12, %rd11, %rd8;
-; CHECK-NEXT: and.b64 %rd13, %rd4, -8608480567731124088;
-; CHECK-NEXT: and.b64 %rd14, %rd6, 4919131752989213764;
-; CHECK-NEXT: mul.lo.s64 %rd15, %rd14, %rd13;
-; CHECK-NEXT: xor.b64 %rd16, %rd12, %rd15;
-; CHECK-NEXT: and.b64 %rd17, %rd4, 4919131752989213764;
-; CHECK-NEXT: and.b64 %rd18, %rd6, -8608480567731124088;
-; CHECK-NEXT: mul.lo.s64 %rd19, %rd18, %rd17;
-; CHECK-NEXT: xor.b64 %rd20, %rd16, %rd19;
-; CHECK-NEXT: and.b64 %rd21, %rd20, 2459565876494606882;
-; CHECK-NEXT: mul.lo.s64 %rd22, %rd7, %rd13;
-; CHECK-NEXT: mul.lo.s64 %rd23, %rd10, %rd5;
-; CHECK-NEXT: xor.b64 %rd24, %rd23, %rd22;
-; CHECK-NEXT: mul.lo.s64 %rd25, %rd14, %rd17;
-; CHECK-NEXT: xor.b64 %rd26, %rd24, %rd25;
-; CHECK-NEXT: mul.lo.s64 %rd27, %rd18, %rd9;
-; CHECK-NEXT: xor.b64 %rd28, %rd26, %rd27;
-; CHECK-NEXT: and.b64 %rd29, %rd28, 1229782938247303441;
-; CHECK-NEXT: or.b64 %rd30, %rd29, %rd21;
-; CHECK-NEXT: mul.lo.s64 %rd31, %rd7, %rd9;
-; CHECK-NEXT: mul.lo.s64 %rd32, %rd10, %rd17;
-; CHECK-NEXT: xor.b64 %rd33, %rd32, %rd31;
-; CHECK-NEXT: mul.lo.s64 %rd34, %rd14, %rd5;
-; CHECK-NEXT: xor.b64 %rd35, %rd33, %rd34;
-; CHECK-NEXT: mul.lo.s64 %rd36, %rd18, %rd13;
-; CHECK-NEXT: xor.b64 %rd37, %rd35, %rd36;
-; CHECK-NEXT: and.b64 %rd38, %rd37, 4919131752989213764;
-; CHECK-NEXT: or.b64 %rd39, %rd30, %rd38;
-; CHECK-NEXT: mul.lo.s64 %rd40, %rd7, %rd17;
-; CHECK-NEXT: mul.lo.s64 %rd41, %rd10, %rd13;
-; CHECK-NEXT: xor.b64 %rd42, %rd41, %rd40;
-; CHECK-NEXT: mul.lo.s64 %rd43, %rd14, %rd9;
-; CHECK-NEXT: xor.b64 %rd44, %rd42, %rd43;
-; CHECK-NEXT: mul.lo.s64 %rd45, %rd18, %rd5;
-; CHECK-NEXT: xor.b64 %rd46, %rd44, %rd45;
-; CHECK-NEXT: and.b64 %rd47, %rd46, 614891469123651720;
-; CHECK-NEXT: or.b64 %rd48, %rd39, %rd47;
-; CHECK-NEXT: brev.b64 %rd49, %rd48;
-; CHECK-NEXT: shr.u64 %rd50, %rd49, 1;
-; CHECK-NEXT: xor.b64 %rd51, %rd50, %rd3;
-; CHECK-NEXT: st.param::func.b64 [func_retval0], %rd51;
+; 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
@@ -134,55 +44,13 @@ define i64 @clmad_hi(i64 %a, i64 %b, i64 %c) {
define i64 @clmul_lo(i64 %a, i64 %b) {
; CHECK-LABEL: clmul_lo(
; CHECK: {
-; CHECK-NEXT: .reg .b64 %rd<46>;
+; 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: and.b64 %rd3, %rd2, 1229782938247303441;
-; CHECK-NEXT: and.b64 %rd4, %rd1, 2459565876494606882;
-; CHECK-NEXT: mul.lo.s64 %rd5, %rd4, %rd3;
-; CHECK-NEXT: and.b64 %rd6, %rd2, 2459565876494606882;
-; CHECK-NEXT: and.b64 %rd7, %rd1, 1229782938247303441;
-; CHECK-NEXT: mul.lo.s64 %rd8, %rd7, %rd6;
-; CHECK-NEXT: xor.b64 %rd9, %rd8, %rd5;
-; CHECK-NEXT: and.b64 %rd10, %rd2, -8608480567731124088;
-; CHECK-NEXT: and.b64 %rd11, %rd1, 4919131752989213764;
-; CHECK-NEXT: mul.lo.s64 %rd12, %rd11, %rd10;
-; CHECK-NEXT: xor.b64 %rd13, %rd9, %rd12;
-; CHECK-NEXT: and.b64 %rd14, %rd2, 4919131752989213764;
-; CHECK-NEXT: and.b64 %rd15, %rd1, -8608480567731124088;
-; CHECK-NEXT: mul.lo.s64 %rd16, %rd15, %rd14;
-; CHECK-NEXT: xor.b64 %rd17, %rd13, %rd16;
-; CHECK-NEXT: and.b64 %rd18, %rd17, 2459565876494606882;
-; CHECK-NEXT: mul.lo.s64 %rd19, %rd4, %rd10;
-; CHECK-NEXT: mul.lo.s64 %rd20, %rd7, %rd3;
-; CHECK-NEXT: xor.b64 %rd21, %rd20, %rd19;
-; CHECK-NEXT: mul.lo.s64 %rd22, %rd11, %rd14;
-; CHECK-NEXT: xor.b64 %rd23, %rd21, %rd22;
-; CHECK-NEXT: mul.lo.s64 %rd24, %rd15, %rd6;
-; CHECK-NEXT: xor.b64 %rd25, %rd23, %rd24;
-; CHECK-NEXT: and.b64 %rd26, %rd25, 1229782938247303441;
-; CHECK-NEXT: or.b64 %rd27, %rd26, %rd18;
-; CHECK-NEXT: mul.lo.s64 %rd28, %rd4, %rd6;
-; CHECK-NEXT: mul.lo.s64 %rd29, %rd7, %rd14;
-; CHECK-NEXT: xor.b64 %rd30, %rd29, %rd28;
-; CHECK-NEXT: mul.lo.s64 %rd31, %rd11, %rd3;
-; CHECK-NEXT: xor.b64 %rd32, %rd30, %rd31;
-; CHECK-NEXT: mul.lo.s64 %rd33, %rd15, %rd10;
-; CHECK-NEXT: xor.b64 %rd34, %rd32, %rd33;
-; CHECK-NEXT: and.b64 %rd35, %rd34, 4919131752989213764;
-; CHECK-NEXT: or.b64 %rd36, %rd27, %rd35;
-; CHECK-NEXT: mul.lo.s64 %rd37, %rd4, %rd14;
-; CHECK-NEXT: mul.lo.s64 %rd38, %rd7, %rd10;
-; CHECK-NEXT: xor.b64 %rd39, %rd38, %rd37;
-; CHECK-NEXT: mul.lo.s64 %rd40, %rd11, %rd6;
-; CHECK-NEXT: xor.b64 %rd41, %rd39, %rd40;
-; CHECK-NEXT: mul.lo.s64 %rd42, %rd15, %rd3;
-; CHECK-NEXT: xor.b64 %rd43, %rd41, %rd42;
-; CHECK-NEXT: and.b64 %rd44, %rd43, -8608480567731124088;
-; CHECK-NEXT: or.b64 %rd45, %rd36, %rd44;
-; CHECK-NEXT: st.param::func.b64 [func_retval0], %rd45;
+; 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
@@ -191,59 +59,13 @@ define i64 @clmul_lo(i64 %a, i64 %b) {
define i64 @clmul_hi(i64 %a, i64 %b) {
; CHECK-LABEL: clmul_hi(
; CHECK: {
-; CHECK-NEXT: .reg .b64 %rd<50>;
+; 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: brev.b64 %rd3, %rd2;
-; CHECK-NEXT: and.b64 %rd4, %rd3, 1229782938247303441;
-; CHECK-NEXT: brev.b64 %rd5, %rd1;
-; CHECK-NEXT: and.b64 %rd6, %rd5, 2459565876494606882;
-; CHECK-NEXT: mul.lo.s64 %rd7, %rd6, %rd4;
-; CHECK-NEXT: and.b64 %rd8, %rd3, 2459565876494606882;
-; CHECK-NEXT: and.b64 %rd9, %rd5, 1229782938247303441;
-; CHECK-NEXT: mul.lo.s64 %rd10, %rd9, %rd8;
-; CHECK-NEXT: xor.b64 %rd11, %rd10, %rd7;
-; CHECK-NEXT: and.b64 %rd12, %rd3, -8608480567731124088;
-; CHECK-NEXT: and.b64 %rd13, %rd5, 4919131752989213764;
-; CHECK-NEXT: mul.lo.s64 %rd14, %rd13, %rd12;
-; CHECK-NEXT: xor.b64 %rd15, %rd11, %rd14;
-; CHECK-NEXT: and.b64 %rd16, %rd3, 4919131752989213764;
-; CHECK-NEXT: and.b64 %rd17, %rd5, -8608480567731124088;
-; CHECK-NEXT: mul.lo.s64 %rd18, %rd17, %rd16;
-; CHECK-NEXT: xor.b64 %rd19, %rd15, %rd18;
-; CHECK-NEXT: and.b64 %rd20, %rd19, 2459565876494606882;
-; CHECK-NEXT: mul.lo.s64 %rd21, %rd6, %rd12;
-; CHECK-NEXT: mul.lo.s64 %rd22, %rd9, %rd4;
-; CHECK-NEXT: xor.b64 %rd23, %rd22, %rd21;
-; CHECK-NEXT: mul.lo.s64 %rd24, %rd13, %rd16;
-; CHECK-NEXT: xor.b64 %rd25, %rd23, %rd24;
-; CHECK-NEXT: mul.lo.s64 %rd26, %rd17, %rd8;
-; CHECK-NEXT: xor.b64 %rd27, %rd25, %rd26;
-; CHECK-NEXT: and.b64 %rd28, %rd27, 1229782938247303441;
-; CHECK-NEXT: or.b64 %rd29, %rd28, %rd20;
-; CHECK-NEXT: mul.lo.s64 %rd30, %rd6, %rd8;
-; CHECK-NEXT: mul.lo.s64 %rd31, %rd9, %rd16;
-; CHECK-NEXT: xor.b64 %rd32, %rd31, %rd30;
-; CHECK-NEXT: mul.lo.s64 %rd33, %rd13, %rd4;
-; CHECK-NEXT: xor.b64 %rd34, %rd32, %rd33;
-; CHECK-NEXT: mul.lo.s64 %rd35, %rd17, %rd12;
-; CHECK-NEXT: xor.b64 %rd36, %rd34, %rd35;
-; CHECK-NEXT: and.b64 %rd37, %rd36, 4919131752989213764;
-; CHECK-NEXT: or.b64 %rd38, %rd29, %rd37;
-; CHECK-NEXT: mul.lo.s64 %rd39, %rd6, %rd16;
-; CHECK-NEXT: mul.lo.s64 %rd40, %rd9, %rd12;
-; CHECK-NEXT: xor.b64 %rd41, %rd40, %rd39;
-; CHECK-NEXT: mul.lo.s64 %rd42, %rd13, %rd8;
-; CHECK-NEXT: xor.b64 %rd43, %rd41, %rd42;
-; CHECK-NEXT: mul.lo.s64 %rd44, %rd17, %rd4;
-; CHECK-NEXT: xor.b64 %rd45, %rd43, %rd44;
-; CHECK-NEXT: and.b64 %rd46, %rd45, 614891469123651720;
-; CHECK-NEXT: or.b64 %rd47, %rd38, %rd46;
-; CHECK-NEXT: brev.b64 %rd48, %rd47;
-; CHECK-NEXT: shr.u64 %rd49, %rd48, 1;
-; CHECK-NEXT: st.param::func.b64 [func_retval0], %rd49;
+; 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
More information about the llvm-commits
mailing list