[llvm] Reland "[NVPTX] Add support for atomic add for f16 type" (PR #85197)

Adrian Kuegel via llvm-commits llvm-commits at lists.llvm.org
Thu Mar 14 23:50:30 PDT 2024


https://github.com/akuegel updated https://github.com/llvm/llvm-project/pull/85197

>From 73787e40241d83fce40631f1e555be0250b1916d Mon Sep 17 00:00:00 2001
From: Adrian Kuegel <akuegel at google.com>
Date: Thu, 14 Mar 2024 08:10:55 +0000
Subject: [PATCH 1/3] Reapply "[NVPTX] Add support for atomic add for f16 type"
 (#84918)

This reverts commit afd47587039e5a93919eb962cfe3a230cc91c504.
---
 llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp |   3 +
 llvm/lib/Target/NVPTX/NVPTXIntrinsics.td    |  15 +++
 llvm/test/CodeGen/NVPTX/atomics-sm70.ll     | 121 ++++++++++++++++++++
 llvm/test/CodeGen/NVPTX/atomics.ll          |   7 ++
 4 files changed, 146 insertions(+)
 create mode 100644 llvm/test/CodeGen/NVPTX/atomics-sm70.ll

diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index c979c03dc1b835..c411c8ef9528d7 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -6100,6 +6100,9 @@ NVPTXTargetLowering::shouldExpandAtomicRMWInIR(AtomicRMWInst *AI) const {
 
   if (AI->isFloatingPointOperation()) {
     if (AI->getOperation() == AtomicRMWInst::BinOp::FAdd) {
+      if (Ty->isHalfTy() && STI.getSmVersion() >= 70 &&
+          STI.getPTXVersion() >= 63)
+        return AtomicExpansionKind::None;
       if (Ty->isFloatTy())
         return AtomicExpansionKind::None;
       if (Ty->isDoubleTy() && STI.hasAtomAddF64())
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 477789a164ead2..869b13369e87e1 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -1630,6 +1630,13 @@ defm INT_PTX_ATOM_ADD_GEN_64 : F_ATOMIC_2<i64, Int64Regs, "", ".u64", ".add",
 defm INT_PTX_ATOM_ADD_GEN_64_USE_G : F_ATOMIC_2<i64, Int64Regs, ".global", ".u64",
   ".add", atomic_load_add_64_gen, i64imm, imm>;
 
+defm INT_PTX_ATOM_ADD_G_F16 : F_ATOMIC_2<f16, Int16Regs, ".global", ".f16", ".add.noftz",
+  atomic_load_add_g, f16imm, fpimm, [hasSM<70>, hasPTX<63>]>;
+defm INT_PTX_ATOM_ADD_S_F16 : F_ATOMIC_2<f16, Int16Regs, ".shared", ".f16", ".add.noftz",
+  atomic_load_add_s, f16imm, fpimm, [hasSM<70>, hasPTX<63>]>;
+defm INT_PTX_ATOM_ADD_GEN_F16 : F_ATOMIC_2<f16, Int16Regs, "", ".f16", ".add.noftz",
+  atomic_load_add_gen, f16imm, fpimm, [hasSM<70>, hasPTX<63>]>;
+
 defm INT_PTX_ATOM_ADD_G_F32 : F_ATOMIC_2<f32, Float32Regs, ".global", ".f32", ".add",
   atomic_load_add_g, f32imm, fpimm>;
 defm INT_PTX_ATOM_ADD_S_F32 : F_ATOMIC_2<f32, Float32Regs, ".shared", ".f32", ".add",
@@ -2007,6 +2014,9 @@ multiclass ATOM2P_impl<string AsmStr,  Intrinsic Intr,
                        SDNode Imm, ValueType ImmTy,
                        list<Predicate> Preds> {
   let AddedComplexity = 1 in {
+    def : ATOM23_impl<AsmStr, regT, regclass, Preds,
+                      (ins Int16Regs:$src, regclass:$b),
+                      (Intr (i16 Int16Regs:$src), (regT regclass:$b))>;
     def : ATOM23_impl<AsmStr, regT, regclass, Preds,
                       (ins Int32Regs:$src, regclass:$b),
                       (Intr (i32 Int32Regs:$src), (regT regclass:$b))>;
@@ -2017,6 +2027,9 @@ multiclass ATOM2P_impl<string AsmStr,  Intrinsic Intr,
   // tablegen can't infer argument types from Intrinsic (though it can
   // from Instruction) so we have to enforce specific type on
   // immediates via explicit cast to ImmTy.
+  def : ATOM23_impl<AsmStr, regT, regclass, Preds,
+                    (ins Int16Regs:$src, ImmType:$b),
+                    (Intr (i16 Int16Regs:$src), (ImmTy Imm:$b))>;
   def : ATOM23_impl<AsmStr, regT, regclass, Preds,
                     (ins Int32Regs:$src, ImmType:$b),
                     (Intr (i32 Int32Regs:$src), (ImmTy Imm:$b))>;
@@ -2136,6 +2149,8 @@ multiclass ATOM2_add_impl<string OpStr> {
    defm _s32  : ATOM2S_impl<OpStr, "i", "s32", i32, Int32Regs, i32imm, imm, i32, []>;
    defm _u32  : ATOM2S_impl<OpStr, "i", "u32", i32, Int32Regs, i32imm, imm, i32, []>;
    defm _u64  : ATOM2S_impl<OpStr, "i", "u64", i64, Int64Regs, i64imm, imm, i64, []>;
+   defm _f16  : ATOM2S_impl<OpStr, "f", "f16", f16, Int16Regs, f16imm, fpimm, f16,
+                            [hasSM<70>, hasPTX<63>]>;
    defm _f32  : ATOM2S_impl<OpStr, "f", "f32", f32, Float32Regs, f32imm, fpimm, f32,
                             []>;
    defm _f64  : ATOM2S_impl<OpStr, "f", "f64", f64, Float64Regs, f64imm, fpimm, f64,
diff --git a/llvm/test/CodeGen/NVPTX/atomics-sm70.ll b/llvm/test/CodeGen/NVPTX/atomics-sm70.ll
new file mode 100644
index 00000000000000..f68bb2049d0061
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/atomics-sm70.ll
@@ -0,0 +1,121 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 4
+; RUN: llc < %s -march=nvptx -mcpu=sm_70 -mattr=+ptx63 | FileCheck %s --check-prefixes=CHECK
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx63 | FileCheck %s --check-prefixes=CHECK64
+; RUN: llc < %s -march=nvptx -mcpu=sm_70 -mattr=+ptx62 | FileCheck %s --check-prefixes=CHECKPTX62
+; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_70 -mattr=+ptx63 | %ptxas-verify -arch=sm_70 %}
+; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx63 | %ptxas-verify -arch=sm_70 %}
+; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_70 -mattr=+ptx62 | %ptxas-verify -arch=sm_70 %}
+
+target triple = "nvptx64-nvidia-cuda"
+
+define void @test(ptr %dp0, ptr addrspace(1) %dp1, ptr addrspace(3) %dp3, half %val) {
+; CHECK-LABEL: test(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b16 %rs<5>;
+; CHECK-NEXT:    .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.u32 %r1, [test_param_0];
+; CHECK-NEXT:    ld.param.b16 %rs1, [test_param_3];
+; CHECK-NEXT:    atom.add.noftz.f16 %rs2, [%r1], %rs1;
+; CHECK-NEXT:    ld.param.u32 %r2, [test_param_1];
+; CHECK-NEXT:    atom.global.add.noftz.f16 %rs3, [%r2], %rs1;
+; CHECK-NEXT:    ld.param.u32 %r3, [test_param_2];
+; CHECK-NEXT:    atom.shared.add.noftz.f16 %rs4, [%r3], %rs1;
+; CHECK-NEXT:    ret;
+;
+; CHECK64-LABEL: test(
+; CHECK64:       {
+; CHECK64-NEXT:    .reg .b16 %rs<5>;
+; CHECK64-NEXT:    .reg .b64 %rd<4>;
+; CHECK64-EMPTY:
+; CHECK64-NEXT:  // %bb.0:
+; CHECK64-NEXT:    ld.param.u64 %rd1, [test_param_0];
+; CHECK64-NEXT:    ld.param.b16 %rs1, [test_param_3];
+; CHECK64-NEXT:    atom.add.noftz.f16 %rs2, [%rd1], %rs1;
+; CHECK64-NEXT:    ld.param.u64 %rd2, [test_param_1];
+; CHECK64-NEXT:    atom.global.add.noftz.f16 %rs3, [%rd2], %rs1;
+; CHECK64-NEXT:    ld.param.u64 %rd3, [test_param_2];
+; CHECK64-NEXT:    atom.shared.add.noftz.f16 %rs4, [%rd3], %rs1;
+; CHECK64-NEXT:    ret;
+;
+; CHECKPTX62-LABEL: test(
+; CHECKPTX62:       {
+; CHECKPTX62-NEXT:    .reg .pred %p<4>;
+; CHECKPTX62-NEXT:    .reg .b16 %rs<14>;
+; CHECKPTX62-NEXT:    .reg .b32 %r<49>;
+; CHECKPTX62-EMPTY:
+; CHECKPTX62-NEXT:  // %bb.0:
+; CHECKPTX62-NEXT:    ld.param.b16 %rs1, [test_param_3];
+; CHECKPTX62-NEXT:    ld.param.u32 %r20, [test_param_2];
+; CHECKPTX62-NEXT:    ld.param.u32 %r19, [test_param_1];
+; CHECKPTX62-NEXT:    ld.param.u32 %r21, [test_param_0];
+; CHECKPTX62-NEXT:    and.b32 %r1, %r21, -4;
+; CHECKPTX62-NEXT:    and.b32 %r22, %r21, 3;
+; CHECKPTX62-NEXT:    shl.b32 %r2, %r22, 3;
+; CHECKPTX62-NEXT:    mov.b32 %r23, 65535;
+; CHECKPTX62-NEXT:    shl.b32 %r24, %r23, %r2;
+; CHECKPTX62-NEXT:    not.b32 %r3, %r24;
+; CHECKPTX62-NEXT:    ld.u32 %r46, [%r1];
+; CHECKPTX62-NEXT:  $L__BB0_1: // %atomicrmw.start
+; CHECKPTX62-NEXT:    // =>This Inner Loop Header: Depth=1
+; CHECKPTX62-NEXT:    shr.u32 %r25, %r46, %r2;
+; CHECKPTX62-NEXT:    cvt.u16.u32 %rs2, %r25;
+; CHECKPTX62-NEXT:    add.rn.f16 %rs4, %rs2, %rs1;
+; CHECKPTX62-NEXT:    cvt.u32.u16 %r26, %rs4;
+; CHECKPTX62-NEXT:    shl.b32 %r27, %r26, %r2;
+; CHECKPTX62-NEXT:    and.b32 %r28, %r46, %r3;
+; CHECKPTX62-NEXT:    or.b32 %r29, %r28, %r27;
+; CHECKPTX62-NEXT:    atom.cas.b32 %r6, [%r1], %r46, %r29;
+; CHECKPTX62-NEXT:    setp.ne.s32 %p1, %r6, %r46;
+; CHECKPTX62-NEXT:    mov.u32 %r46, %r6;
+; CHECKPTX62-NEXT:    @%p1 bra $L__BB0_1;
+; CHECKPTX62-NEXT:  // %bb.2: // %atomicrmw.end
+; CHECKPTX62-NEXT:    and.b32 %r7, %r19, -4;
+; CHECKPTX62-NEXT:    shl.b32 %r30, %r19, 3;
+; CHECKPTX62-NEXT:    and.b32 %r8, %r30, 24;
+; CHECKPTX62-NEXT:    shl.b32 %r32, %r23, %r8;
+; CHECKPTX62-NEXT:    not.b32 %r9, %r32;
+; CHECKPTX62-NEXT:    ld.global.u32 %r47, [%r7];
+; CHECKPTX62-NEXT:  $L__BB0_3: // %atomicrmw.start9
+; CHECKPTX62-NEXT:    // =>This Inner Loop Header: Depth=1
+; CHECKPTX62-NEXT:    shr.u32 %r33, %r47, %r8;
+; CHECKPTX62-NEXT:    cvt.u16.u32 %rs6, %r33;
+; CHECKPTX62-NEXT:    add.rn.f16 %rs8, %rs6, %rs1;
+; CHECKPTX62-NEXT:    cvt.u32.u16 %r34, %rs8;
+; CHECKPTX62-NEXT:    shl.b32 %r35, %r34, %r8;
+; CHECKPTX62-NEXT:    and.b32 %r36, %r47, %r9;
+; CHECKPTX62-NEXT:    or.b32 %r37, %r36, %r35;
+; CHECKPTX62-NEXT:    atom.global.cas.b32 %r12, [%r7], %r47, %r37;
+; CHECKPTX62-NEXT:    setp.ne.s32 %p2, %r12, %r47;
+; CHECKPTX62-NEXT:    mov.u32 %r47, %r12;
+; CHECKPTX62-NEXT:    @%p2 bra $L__BB0_3;
+; CHECKPTX62-NEXT:  // %bb.4: // %atomicrmw.end8
+; CHECKPTX62-NEXT:    and.b32 %r13, %r20, -4;
+; CHECKPTX62-NEXT:    shl.b32 %r38, %r20, 3;
+; CHECKPTX62-NEXT:    and.b32 %r14, %r38, 24;
+; CHECKPTX62-NEXT:    shl.b32 %r40, %r23, %r14;
+; CHECKPTX62-NEXT:    not.b32 %r15, %r40;
+; CHECKPTX62-NEXT:    ld.shared.u32 %r48, [%r13];
+; CHECKPTX62-NEXT:  $L__BB0_5: // %atomicrmw.start27
+; CHECKPTX62-NEXT:    // =>This Inner Loop Header: Depth=1
+; CHECKPTX62-NEXT:    shr.u32 %r41, %r48, %r14;
+; CHECKPTX62-NEXT:    cvt.u16.u32 %rs10, %r41;
+; CHECKPTX62-NEXT:    add.rn.f16 %rs12, %rs10, %rs1;
+; CHECKPTX62-NEXT:    cvt.u32.u16 %r42, %rs12;
+; CHECKPTX62-NEXT:    shl.b32 %r43, %r42, %r14;
+; CHECKPTX62-NEXT:    and.b32 %r44, %r48, %r15;
+; CHECKPTX62-NEXT:    or.b32 %r45, %r44, %r43;
+; CHECKPTX62-NEXT:    atom.shared.cas.b32 %r18, [%r13], %r48, %r45;
+; CHECKPTX62-NEXT:    setp.ne.s32 %p3, %r18, %r48;
+; CHECKPTX62-NEXT:    mov.u32 %r48, %r18;
+; CHECKPTX62-NEXT:    @%p3 bra $L__BB0_5;
+; CHECKPTX62-NEXT:  // %bb.6: // %atomicrmw.end26
+; CHECKPTX62-NEXT:    ret;
+  %r1 = atomicrmw fadd ptr %dp0, half %val seq_cst
+  %r2 = atomicrmw fadd ptr addrspace(1) %dp1, half %val seq_cst
+  %ret = atomicrmw fadd ptr addrspace(3) %dp3, half %val seq_cst
+  ret void
+}
+
+attributes #1 = { argmemonly nounwind }
diff --git a/llvm/test/CodeGen/NVPTX/atomics.ll b/llvm/test/CodeGen/NVPTX/atomics.ll
index e99d0fd05e346b..6f2b5dcf47f13b 100644
--- a/llvm/test/CodeGen/NVPTX/atomics.ll
+++ b/llvm/test/CodeGen/NVPTX/atomics.ll
@@ -175,6 +175,13 @@ define float @atomicrmw_add_f32_generic(ptr %addr, float %val) {
   ret float %ret
 }
 
+; CHECK-LABEL: atomicrmw_add_f16_generic
+define half @atomicrmw_add_f16_generic(ptr %addr, half %val) {
+; CHECK: atom.cas
+  %ret = atomicrmw fadd ptr %addr, half %val seq_cst
+  ret half %ret
+}
+
 ; CHECK-LABEL: atomicrmw_add_f32_addrspace1
 define float @atomicrmw_add_f32_addrspace1(ptr addrspace(1) %addr, float %val) {
 ; CHECK: atom.global.add.f32

>From 5e73429a14b640ba89debc718038de013eb35d19 Mon Sep 17 00:00:00 2001
From: Adrian Kuegel <akuegel at google.com>
Date: Thu, 14 Mar 2024 08:20:54 +0000
Subject: [PATCH 2/3] Fix problem with fp16 immediates

---
 llvm/lib/Target/NVPTX/NVPTXIntrinsics.td |   4 +-
 llvm/test/CodeGen/NVPTX/atomics-sm70.ll  | 139 +++++++++++++----------
 2 files changed, 83 insertions(+), 60 deletions(-)

diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 869b13369e87e1..e6a2b67c28f510 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -1520,11 +1520,13 @@ multiclass F_ATOMIC_2_imp<ValueType ptrT, NVPTXRegClass ptrclass,
   def imm : NVPTXInst<(outs regclass:$dst), (ins ptrclass:$addr, IMMType:$b),
     !strconcat("atom", SpaceStr, OpcStr, TypeStr, " \t$dst, [$addr], $b;", ""),
     [(set (regT regclass:$dst), (IntOp (ptrT ptrclass:$addr), IMM:$b))]>,
-  Requires<Pred>;
+  Requires<!listconcat(Pred, [!if(!eq(TypeStr, ".f16"), Predicate<"false">, Predicate<"true">)])>;
 }
 multiclass F_ATOMIC_2<ValueType regT, NVPTXRegClass regclass, string SpaceStr, string TypeStr,
   string OpcStr, PatFrag IntOp, Operand IMMType, SDNode IMM,
   list<Predicate> Pred = []> {
+  defm p16 : F_ATOMIC_2_imp<i16, Int16Regs, regT, regclass, SpaceStr, TypeStr, OpcStr,
+    IntOp, IMMType, IMM, Pred>;
   defm p32 : F_ATOMIC_2_imp<i32, Int32Regs, regT, regclass, SpaceStr, TypeStr, OpcStr,
     IntOp, IMMType, IMM, Pred>;
   defm p64 : F_ATOMIC_2_imp<i64, Int64Regs, regT, regclass, SpaceStr, TypeStr, OpcStr,
diff --git a/llvm/test/CodeGen/NVPTX/atomics-sm70.ll b/llvm/test/CodeGen/NVPTX/atomics-sm70.ll
index f68bb2049d0061..9cc45fbe313b7e 100644
--- a/llvm/test/CodeGen/NVPTX/atomics-sm70.ll
+++ b/llvm/test/CodeGen/NVPTX/atomics-sm70.ll
@@ -11,7 +11,7 @@ target triple = "nvptx64-nvidia-cuda"
 define void @test(ptr %dp0, ptr addrspace(1) %dp1, ptr addrspace(3) %dp3, half %val) {
 ; CHECK-LABEL: test(
 ; CHECK:       {
-; CHECK-NEXT:    .reg .b16 %rs<5>;
+; CHECK-NEXT:    .reg .b16 %rs<7>;
 ; CHECK-NEXT:    .reg .b32 %r<4>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
@@ -19,14 +19,16 @@ define void @test(ptr %dp0, ptr addrspace(1) %dp1, ptr addrspace(3) %dp3, half %
 ; CHECK-NEXT:    ld.param.b16 %rs1, [test_param_3];
 ; CHECK-NEXT:    atom.add.noftz.f16 %rs2, [%r1], %rs1;
 ; CHECK-NEXT:    ld.param.u32 %r2, [test_param_1];
-; CHECK-NEXT:    atom.global.add.noftz.f16 %rs3, [%r2], %rs1;
+; CHECK-NEXT:    mov.b16 %rs3, 0x3C00;
+; CHECK-NEXT:    atom.add.noftz.f16 %rs4, [%r1], %rs3;
 ; CHECK-NEXT:    ld.param.u32 %r3, [test_param_2];
-; CHECK-NEXT:    atom.shared.add.noftz.f16 %rs4, [%r3], %rs1;
+; CHECK-NEXT:    atom.global.add.noftz.f16 %rs5, [%r2], %rs1;
+; CHECK-NEXT:    atom.shared.add.noftz.f16 %rs6, [%r3], %rs1;
 ; CHECK-NEXT:    ret;
 ;
 ; CHECK64-LABEL: test(
 ; CHECK64:       {
-; CHECK64-NEXT:    .reg .b16 %rs<5>;
+; CHECK64-NEXT:    .reg .b16 %rs<7>;
 ; CHECK64-NEXT:    .reg .b64 %rd<4>;
 ; CHECK64-EMPTY:
 ; CHECK64-NEXT:  // %bb.0:
@@ -34,87 +36,106 @@ define void @test(ptr %dp0, ptr addrspace(1) %dp1, ptr addrspace(3) %dp3, half %
 ; CHECK64-NEXT:    ld.param.b16 %rs1, [test_param_3];
 ; CHECK64-NEXT:    atom.add.noftz.f16 %rs2, [%rd1], %rs1;
 ; CHECK64-NEXT:    ld.param.u64 %rd2, [test_param_1];
-; CHECK64-NEXT:    atom.global.add.noftz.f16 %rs3, [%rd2], %rs1;
+; CHECK64-NEXT:    mov.b16 %rs3, 0x3C00;
+; CHECK64-NEXT:    atom.add.noftz.f16 %rs4, [%rd1], %rs3;
 ; CHECK64-NEXT:    ld.param.u64 %rd3, [test_param_2];
-; CHECK64-NEXT:    atom.shared.add.noftz.f16 %rs4, [%rd3], %rs1;
+; CHECK64-NEXT:    atom.global.add.noftz.f16 %rs5, [%rd2], %rs1;
+; CHECK64-NEXT:    atom.shared.add.noftz.f16 %rs6, [%rd3], %rs1;
 ; CHECK64-NEXT:    ret;
 ;
 ; CHECKPTX62-LABEL: test(
 ; CHECKPTX62:       {
-; CHECKPTX62-NEXT:    .reg .pred %p<4>;
-; CHECKPTX62-NEXT:    .reg .b16 %rs<14>;
-; CHECKPTX62-NEXT:    .reg .b32 %r<49>;
+; CHECKPTX62-NEXT:    .reg .pred %p<5>;
+; CHECKPTX62-NEXT:    .reg .b16 %rs<19>;
+; CHECKPTX62-NEXT:    .reg .b32 %r<58>;
 ; CHECKPTX62-EMPTY:
 ; CHECKPTX62-NEXT:  // %bb.0:
 ; CHECKPTX62-NEXT:    ld.param.b16 %rs1, [test_param_3];
-; CHECKPTX62-NEXT:    ld.param.u32 %r20, [test_param_2];
-; CHECKPTX62-NEXT:    ld.param.u32 %r19, [test_param_1];
-; CHECKPTX62-NEXT:    ld.param.u32 %r21, [test_param_0];
-; CHECKPTX62-NEXT:    and.b32 %r1, %r21, -4;
-; CHECKPTX62-NEXT:    and.b32 %r22, %r21, 3;
-; CHECKPTX62-NEXT:    shl.b32 %r2, %r22, 3;
-; CHECKPTX62-NEXT:    mov.b32 %r23, 65535;
-; CHECKPTX62-NEXT:    shl.b32 %r24, %r23, %r2;
-; CHECKPTX62-NEXT:    not.b32 %r3, %r24;
-; CHECKPTX62-NEXT:    ld.u32 %r46, [%r1];
+; CHECKPTX62-NEXT:    ld.param.u32 %r23, [test_param_2];
+; CHECKPTX62-NEXT:    ld.param.u32 %r22, [test_param_1];
+; CHECKPTX62-NEXT:    ld.param.u32 %r24, [test_param_0];
+; CHECKPTX62-NEXT:    and.b32 %r1, %r24, -4;
+; CHECKPTX62-NEXT:    and.b32 %r25, %r24, 3;
+; CHECKPTX62-NEXT:    shl.b32 %r2, %r25, 3;
+; CHECKPTX62-NEXT:    mov.b32 %r26, 65535;
+; CHECKPTX62-NEXT:    shl.b32 %r27, %r26, %r2;
+; CHECKPTX62-NEXT:    not.b32 %r3, %r27;
+; CHECKPTX62-NEXT:    ld.u32 %r54, [%r1];
 ; CHECKPTX62-NEXT:  $L__BB0_1: // %atomicrmw.start
 ; CHECKPTX62-NEXT:    // =>This Inner Loop Header: Depth=1
-; CHECKPTX62-NEXT:    shr.u32 %r25, %r46, %r2;
-; CHECKPTX62-NEXT:    cvt.u16.u32 %rs2, %r25;
+; CHECKPTX62-NEXT:    shr.u32 %r28, %r54, %r2;
+; CHECKPTX62-NEXT:    cvt.u16.u32 %rs2, %r28;
 ; CHECKPTX62-NEXT:    add.rn.f16 %rs4, %rs2, %rs1;
-; CHECKPTX62-NEXT:    cvt.u32.u16 %r26, %rs4;
-; CHECKPTX62-NEXT:    shl.b32 %r27, %r26, %r2;
-; CHECKPTX62-NEXT:    and.b32 %r28, %r46, %r3;
-; CHECKPTX62-NEXT:    or.b32 %r29, %r28, %r27;
-; CHECKPTX62-NEXT:    atom.cas.b32 %r6, [%r1], %r46, %r29;
-; CHECKPTX62-NEXT:    setp.ne.s32 %p1, %r6, %r46;
-; CHECKPTX62-NEXT:    mov.u32 %r46, %r6;
+; CHECKPTX62-NEXT:    cvt.u32.u16 %r29, %rs4;
+; CHECKPTX62-NEXT:    shl.b32 %r30, %r29, %r2;
+; CHECKPTX62-NEXT:    and.b32 %r31, %r54, %r3;
+; CHECKPTX62-NEXT:    or.b32 %r32, %r31, %r30;
+; CHECKPTX62-NEXT:    atom.cas.b32 %r6, [%r1], %r54, %r32;
+; CHECKPTX62-NEXT:    setp.ne.s32 %p1, %r6, %r54;
+; CHECKPTX62-NEXT:    mov.u32 %r54, %r6;
 ; CHECKPTX62-NEXT:    @%p1 bra $L__BB0_1;
 ; CHECKPTX62-NEXT:  // %bb.2: // %atomicrmw.end
-; CHECKPTX62-NEXT:    and.b32 %r7, %r19, -4;
-; CHECKPTX62-NEXT:    shl.b32 %r30, %r19, 3;
-; CHECKPTX62-NEXT:    and.b32 %r8, %r30, 24;
-; CHECKPTX62-NEXT:    shl.b32 %r32, %r23, %r8;
-; CHECKPTX62-NEXT:    not.b32 %r9, %r32;
-; CHECKPTX62-NEXT:    ld.global.u32 %r47, [%r7];
+; CHECKPTX62-NEXT:    ld.u32 %r55, [%r1];
 ; CHECKPTX62-NEXT:  $L__BB0_3: // %atomicrmw.start9
 ; CHECKPTX62-NEXT:    // =>This Inner Loop Header: Depth=1
-; CHECKPTX62-NEXT:    shr.u32 %r33, %r47, %r8;
+; CHECKPTX62-NEXT:    shr.u32 %r33, %r55, %r2;
 ; CHECKPTX62-NEXT:    cvt.u16.u32 %rs6, %r33;
-; CHECKPTX62-NEXT:    add.rn.f16 %rs8, %rs6, %rs1;
-; CHECKPTX62-NEXT:    cvt.u32.u16 %r34, %rs8;
-; CHECKPTX62-NEXT:    shl.b32 %r35, %r34, %r8;
-; CHECKPTX62-NEXT:    and.b32 %r36, %r47, %r9;
+; CHECKPTX62-NEXT:    mov.b16 %rs8, 0x3C00;
+; CHECKPTX62-NEXT:    add.rn.f16 %rs9, %rs6, %rs8;
+; CHECKPTX62-NEXT:    cvt.u32.u16 %r34, %rs9;
+; CHECKPTX62-NEXT:    shl.b32 %r35, %r34, %r2;
+; CHECKPTX62-NEXT:    and.b32 %r36, %r55, %r3;
 ; CHECKPTX62-NEXT:    or.b32 %r37, %r36, %r35;
-; CHECKPTX62-NEXT:    atom.global.cas.b32 %r12, [%r7], %r47, %r37;
-; CHECKPTX62-NEXT:    setp.ne.s32 %p2, %r12, %r47;
-; CHECKPTX62-NEXT:    mov.u32 %r47, %r12;
+; CHECKPTX62-NEXT:    atom.cas.b32 %r9, [%r1], %r55, %r37;
+; CHECKPTX62-NEXT:    setp.ne.s32 %p2, %r9, %r55;
+; CHECKPTX62-NEXT:    mov.u32 %r55, %r9;
 ; CHECKPTX62-NEXT:    @%p2 bra $L__BB0_3;
 ; CHECKPTX62-NEXT:  // %bb.4: // %atomicrmw.end8
-; CHECKPTX62-NEXT:    and.b32 %r13, %r20, -4;
-; CHECKPTX62-NEXT:    shl.b32 %r38, %r20, 3;
-; CHECKPTX62-NEXT:    and.b32 %r14, %r38, 24;
-; CHECKPTX62-NEXT:    shl.b32 %r40, %r23, %r14;
-; CHECKPTX62-NEXT:    not.b32 %r15, %r40;
-; CHECKPTX62-NEXT:    ld.shared.u32 %r48, [%r13];
+; CHECKPTX62-NEXT:    and.b32 %r10, %r22, -4;
+; CHECKPTX62-NEXT:    shl.b32 %r38, %r22, 3;
+; CHECKPTX62-NEXT:    and.b32 %r11, %r38, 24;
+; CHECKPTX62-NEXT:    shl.b32 %r40, %r26, %r11;
+; CHECKPTX62-NEXT:    not.b32 %r12, %r40;
+; CHECKPTX62-NEXT:    ld.global.u32 %r56, [%r10];
 ; CHECKPTX62-NEXT:  $L__BB0_5: // %atomicrmw.start27
 ; CHECKPTX62-NEXT:    // =>This Inner Loop Header: Depth=1
-; CHECKPTX62-NEXT:    shr.u32 %r41, %r48, %r14;
-; CHECKPTX62-NEXT:    cvt.u16.u32 %rs10, %r41;
-; CHECKPTX62-NEXT:    add.rn.f16 %rs12, %rs10, %rs1;
-; CHECKPTX62-NEXT:    cvt.u32.u16 %r42, %rs12;
-; CHECKPTX62-NEXT:    shl.b32 %r43, %r42, %r14;
-; CHECKPTX62-NEXT:    and.b32 %r44, %r48, %r15;
+; CHECKPTX62-NEXT:    shr.u32 %r41, %r56, %r11;
+; CHECKPTX62-NEXT:    cvt.u16.u32 %rs11, %r41;
+; CHECKPTX62-NEXT:    add.rn.f16 %rs13, %rs11, %rs1;
+; CHECKPTX62-NEXT:    cvt.u32.u16 %r42, %rs13;
+; CHECKPTX62-NEXT:    shl.b32 %r43, %r42, %r11;
+; CHECKPTX62-NEXT:    and.b32 %r44, %r56, %r12;
 ; CHECKPTX62-NEXT:    or.b32 %r45, %r44, %r43;
-; CHECKPTX62-NEXT:    atom.shared.cas.b32 %r18, [%r13], %r48, %r45;
-; CHECKPTX62-NEXT:    setp.ne.s32 %p3, %r18, %r48;
-; CHECKPTX62-NEXT:    mov.u32 %r48, %r18;
+; CHECKPTX62-NEXT:    atom.global.cas.b32 %r15, [%r10], %r56, %r45;
+; CHECKPTX62-NEXT:    setp.ne.s32 %p3, %r15, %r56;
+; CHECKPTX62-NEXT:    mov.u32 %r56, %r15;
 ; CHECKPTX62-NEXT:    @%p3 bra $L__BB0_5;
 ; CHECKPTX62-NEXT:  // %bb.6: // %atomicrmw.end26
+; CHECKPTX62-NEXT:    and.b32 %r16, %r23, -4;
+; CHECKPTX62-NEXT:    shl.b32 %r46, %r23, 3;
+; CHECKPTX62-NEXT:    and.b32 %r17, %r46, 24;
+; CHECKPTX62-NEXT:    shl.b32 %r48, %r26, %r17;
+; CHECKPTX62-NEXT:    not.b32 %r18, %r48;
+; CHECKPTX62-NEXT:    ld.shared.u32 %r57, [%r16];
+; CHECKPTX62-NEXT:  $L__BB0_7: // %atomicrmw.start45
+; CHECKPTX62-NEXT:    // =>This Inner Loop Header: Depth=1
+; CHECKPTX62-NEXT:    shr.u32 %r49, %r57, %r17;
+; CHECKPTX62-NEXT:    cvt.u16.u32 %rs15, %r49;
+; CHECKPTX62-NEXT:    add.rn.f16 %rs17, %rs15, %rs1;
+; CHECKPTX62-NEXT:    cvt.u32.u16 %r50, %rs17;
+; CHECKPTX62-NEXT:    shl.b32 %r51, %r50, %r17;
+; CHECKPTX62-NEXT:    and.b32 %r52, %r57, %r18;
+; CHECKPTX62-NEXT:    or.b32 %r53, %r52, %r51;
+; CHECKPTX62-NEXT:    atom.shared.cas.b32 %r21, [%r16], %r57, %r53;
+; CHECKPTX62-NEXT:    setp.ne.s32 %p4, %r21, %r57;
+; CHECKPTX62-NEXT:    mov.u32 %r57, %r21;
+; CHECKPTX62-NEXT:    @%p4 bra $L__BB0_7;
+; CHECKPTX62-NEXT:  // %bb.8: // %atomicrmw.end44
 ; CHECKPTX62-NEXT:    ret;
   %r1 = atomicrmw fadd ptr %dp0, half %val seq_cst
-  %r2 = atomicrmw fadd ptr addrspace(1) %dp1, half %val seq_cst
-  %ret = atomicrmw fadd ptr addrspace(3) %dp3, half %val seq_cst
+  %r2 = atomicrmw fadd ptr %dp0, half 1.0 seq_cst
+  %r3 = atomicrmw fadd ptr addrspace(1) %dp1, half %val seq_cst
+  %r4 = atomicrmw fadd ptr addrspace(3) %dp3, half %val seq_cst
   ret void
 }
 

>From d09b706e7ecadaf54ad8a1c264b88630eb35e1d8 Mon Sep 17 00:00:00 2001
From: Adrian Kuegel <akuegel at google.com>
Date: Fri, 15 Mar 2024 06:49:39 +0000
Subject: [PATCH 3/3] Simplify according to review comments

---
 llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 4 +---
 1 file changed, 1 insertion(+), 3 deletions(-)

diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index e6a2b67c28f510..c0c53380a13e9b 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -1520,13 +1520,11 @@ multiclass F_ATOMIC_2_imp<ValueType ptrT, NVPTXRegClass ptrclass,
   def imm : NVPTXInst<(outs regclass:$dst), (ins ptrclass:$addr, IMMType:$b),
     !strconcat("atom", SpaceStr, OpcStr, TypeStr, " \t$dst, [$addr], $b;", ""),
     [(set (regT regclass:$dst), (IntOp (ptrT ptrclass:$addr), IMM:$b))]>,
-  Requires<!listconcat(Pred, [!if(!eq(TypeStr, ".f16"), Predicate<"false">, Predicate<"true">)])>;
+  Requires<!if(!eq(TypeStr, ".f16"), [Predicate<"false">], Pred)>;
 }
 multiclass F_ATOMIC_2<ValueType regT, NVPTXRegClass regclass, string SpaceStr, string TypeStr,
   string OpcStr, PatFrag IntOp, Operand IMMType, SDNode IMM,
   list<Predicate> Pred = []> {
-  defm p16 : F_ATOMIC_2_imp<i16, Int16Regs, regT, regclass, SpaceStr, TypeStr, OpcStr,
-    IntOp, IMMType, IMM, Pred>;
   defm p32 : F_ATOMIC_2_imp<i32, Int32Regs, regT, regclass, SpaceStr, TypeStr, OpcStr,
     IntOp, IMMType, IMM, Pred>;
   defm p64 : F_ATOMIC_2_imp<i64, Int64Regs, regT, regclass, SpaceStr, TypeStr, OpcStr,



More information about the llvm-commits mailing list