factor out intrinsics update change to another PR

Author: modiking

llvm/test/CodeGen/NVPTX/distributed-shared-cluster.ll

@@ -4,56 +4,6 @@
 
 target triple = "nvptx64-nvidia-cuda"
 
-declare ptr addrspace(7) @llvm.nvvm.mapa.shared.cluster(ptr addrspace(3), i32)
-declare i1 @llvm.nvvm.isspacep.shared.cluster(ptr)
-declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
-declare ptr @llvm.nvvm.mapa(ptr, i32)
-
-; Common setup for distributed shared memory cluster addressing
-define i32 @test_distributed_shared_cluster_common(ptr %ptr, ptr addrspace(3) %smem_ptr) local_unnamed_addr {
-; CHECK-LABEL: test_distributed_shared_cluster_common(
-; CHECK:       {
-; CHECK-NEXT:    .reg .pred %p<3>;
-; CHECK-NEXT:    .reg .b32 %r<8>;
-; CHECK-NEXT:    .reg .b64 %rd<5>;
-; CHECK-EMPTY:
-; CHECK-NEXT:  // %bb.0: // %entry
-; CHECK-NEXT:    ld.param.u64 %rd1, [test_distributed_shared_cluster_common_param_0];
-; CHECK-NEXT:    ld.param.u64 %rd2, [test_distributed_shared_cluster_common_param_1];
-; CHECK-NEXT:    mov.u32 %r1, %ctaid.x;
-; CHECK-NEXT:    xor.b32 %r2, %r1, 1;
-; CHECK-NEXT:    isspacep.shared::cluster %p1, %rd1;
-; CHECK-NEXT:    mapa.u64 %rd3, %rd1, %r2;
-; CHECK-NEXT:    isspacep.shared::cluster %p2, %rd3;
-; CHECK-NEXT:    mapa.shared::cluster.u64 %rd4, %rd2, %r2;
-; CHECK-NEXT:    ld.shared::cluster.u32 %r3, [%rd4];
-; CHECK-NEXT:    add.s32 %r4, %r3, 42;
-; CHECK-NEXT:    st.shared::cluster.u32 [%rd4], %r4;
-; CHECK-NEXT:    selp.b32 %r5, 1, 0, %p1;
-; CHECK-NEXT:    selp.b32 %r6, 1, 0, %p2;
-; CHECK-NEXT:    add.s32 %r7, %r5, %r6;
-; CHECK-NEXT:    st.param.b32 [func_retval0], %r7;
-; CHECK-NEXT:    ret;
-entry:
-  %0 = tail call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
-  %1 = xor i32 %0, 1
-  %2 = tail call i1 @llvm.nvvm.isspacep.shared.cluster(ptr %ptr)
-  %3 = tail call ptr @llvm.nvvm.mapa(ptr %ptr, i32 %1)
-  %4 = tail call i1 @llvm.nvvm.isspacep.shared.cluster(ptr %3)
-  %dsmem_ptr = call ptr addrspace(7) @llvm.nvvm.mapa.shared.cluster(ptr addrspace(3) %smem_ptr, i32 %1)
-
-  ; Add load and store to the distributed shared memory cluster
-  %loaded_val = load i32, ptr addrspace(7) %dsmem_ptr
-  %updated_val = add i32 %loaded_val, 42
-  store i32 %updated_val, ptr addrspace(7) %dsmem_ptr
-
-  ; Return value preserves the isspacep test results plus the value operation
-  %5 = zext i1 %2 to i32
-  %6 = zext i1 %4 to i32
-  %ret = add i32 %5, %6
-  ret i32 %ret
-}
-
 ; Floating point atomic operations tests
 define void @test_distributed_shared_cluster_float_atomic(ptr addrspace(7) %dsmem_ptr) local_unnamed_addr {
 ; CHECK-LABEL: test_distributed_shared_cluster_float_atomic(
@@ -207,84 +157,84 @@ define void @test_distributed_shared_cluster_cmpxchg(ptr addrspace(7) %dsmem_ptr
 ; CHECK-NEXT:    shl.b32 %r3, %r37, %r1;
 ; CHECK-NEXT:    ld.shared::cluster.u32 %r38, [%rd1];
 ; CHECK-NEXT:    and.b32 %r48, %r38, %r2;
-; CHECK-NEXT:  $L__BB4_1: // %partword.cmpxchg.loop33
+; CHECK-NEXT:  $L__BB3_1: // %partword.cmpxchg.loop33
 ; CHECK-NEXT:    // =>This Inner Loop Header: Depth=1
 ; CHECK-NEXT:    or.b32 %r39, %r48, %r3;
 ; CHECK-NEXT:    atom.relaxed.shared::cluster.cas.b32 %r6, [%rd1], %r39, %r48;
 ; CHECK-NEXT:    setp.eq.s32 %p1, %r6, %r39;
-; CHECK-NEXT:    @%p1 bra $L__BB4_3;
+; CHECK-NEXT:    @%p1 bra $L__BB3_3;
 ; CHECK-NEXT:  // %bb.2: // %partword.cmpxchg.failure32
-; CHECK-NEXT:    // in Loop: Header=BB4_1 Depth=1
+; CHECK-NEXT:    // in Loop: Header=BB3_1 Depth=1
 ; CHECK-NEXT:    and.b32 %r7, %r6, %r2;
 ; CHECK-NEXT:    setp.ne.s32 %p2, %r48, %r7;
 ; CHECK-NEXT:    mov.b32 %r48, %r7;
-; CHECK-NEXT:    @%p2 bra $L__BB4_1;
-; CHECK-NEXT:  $L__BB4_3: // %partword.cmpxchg.end31
+; CHECK-NEXT:    @%p2 bra $L__BB3_1;
+; CHECK-NEXT:  $L__BB3_3: // %partword.cmpxchg.end31
 ; CHECK-NEXT:    ld.shared::cluster.u32 %r40, [%rd1];
 ; CHECK-NEXT:    and.b32 %r49, %r40, %r2;
-; CHECK-NEXT:  $L__BB4_4: // %partword.cmpxchg.loop23
+; CHECK-NEXT:  $L__BB3_4: // %partword.cmpxchg.loop23
 ; CHECK-NEXT:    // =>This Inner Loop Header: Depth=1
 ; CHECK-NEXT:    or.b32 %r41, %r49, %r3;
 ; CHECK-NEXT:    atom.relaxed.shared::cluster.cas.b32 %r10, [%rd1], %r41, %r49;
 ; CHECK-NEXT:    setp.eq.s32 %p3, %r10, %r41;
-; CHECK-NEXT:    @%p3 bra $L__BB4_6;
+; CHECK-NEXT:    @%p3 bra $L__BB3_6;
 ; CHECK-NEXT:  // %bb.5: // %partword.cmpxchg.failure22
-; CHECK-NEXT:    // in Loop: Header=BB4_4 Depth=1
+; CHECK-NEXT:    // in Loop: Header=BB3_4 Depth=1
 ; CHECK-NEXT:    and.b32 %r11, %r10, %r2;
 ; CHECK-NEXT:    setp.ne.s32 %p4, %r49, %r11;
 ; CHECK-NEXT:    mov.b32 %r49, %r11;
-; CHECK-NEXT:    @%p4 bra $L__BB4_4;
-; CHECK-NEXT:  $L__BB4_6: // %partword.cmpxchg.end21
+; CHECK-NEXT:    @%p4 bra $L__BB3_4;
+; CHECK-NEXT:  $L__BB3_6: // %partword.cmpxchg.end21
 ; CHECK-NEXT:    fence.acq_rel.sys;
 ; CHECK-NEXT:    fence.acq_rel.sys;
 ; CHECK-NEXT:    ld.shared::cluster.u32 %r42, [%rd1];
 ; CHECK-NEXT:    and.b32 %r50, %r42, %r2;
-; CHECK-NEXT:  $L__BB4_7: // %partword.cmpxchg.loop13
+; CHECK-NEXT:  $L__BB3_7: // %partword.cmpxchg.loop13
 ; CHECK-NEXT:    // =>This Inner Loop Header: Depth=1
 ; CHECK-NEXT:    or.b32 %r43, %r50, %r3;
 ; CHECK-NEXT:    atom.relaxed.shared::cluster.cas.b32 %r14, [%rd1], %r43, %r50;
 ; CHECK-NEXT:    setp.eq.s32 %p5, %r14, %r43;
-; CHECK-NEXT:    @%p5 bra $L__BB4_9;
+; CHECK-NEXT:    @%p5 bra $L__BB3_9;
 ; CHECK-NEXT:  // %bb.8: // %partword.cmpxchg.failure12
-; CHECK-NEXT:    // in Loop: Header=BB4_7 Depth=1
+; CHECK-NEXT:    // in Loop: Header=BB3_7 Depth=1
 ; CHECK-NEXT:    and.b32 %r15, %r14, %r2;
 ; CHECK-NEXT:    setp.ne.s32 %p6, %r50, %r15;
 ; CHECK-NEXT:    mov.b32 %r50, %r15;
-; CHECK-NEXT:    @%p6 bra $L__BB4_7;
-; CHECK-NEXT:  $L__BB4_9: // %partword.cmpxchg.end11
+; CHECK-NEXT:    @%p6 bra $L__BB3_7;
+; CHECK-NEXT:  $L__BB3_9: // %partword.cmpxchg.end11
 ; CHECK-NEXT:    fence.acq_rel.sys;
 ; CHECK-NEXT:    ld.shared::cluster.u32 %r44, [%rd1];
 ; CHECK-NEXT:    and.b32 %r51, %r44, %r2;
-; CHECK-NEXT:  $L__BB4_10: // %partword.cmpxchg.loop3
+; CHECK-NEXT:  $L__BB3_10: // %partword.cmpxchg.loop3
 ; CHECK-NEXT:    // =>This Inner Loop Header: Depth=1
 ; CHECK-NEXT:    or.b32 %r45, %r51, %r3;
 ; CHECK-NEXT:    atom.relaxed.shared::cluster.cas.b32 %r18, [%rd1], %r45, %r51;
 ; CHECK-NEXT:    setp.eq.s32 %p7, %r18, %r45;
-; CHECK-NEXT:    @%p7 bra $L__BB4_12;
+; CHECK-NEXT:    @%p7 bra $L__BB3_12;
 ; CHECK-NEXT:  // %bb.11: // %partword.cmpxchg.failure2
-; CHECK-NEXT:    // in Loop: Header=BB4_10 Depth=1
+; CHECK-NEXT:    // in Loop: Header=BB3_10 Depth=1
 ; CHECK-NEXT:    and.b32 %r19, %r18, %r2;
 ; CHECK-NEXT:    setp.ne.s32 %p8, %r51, %r19;
 ; CHECK-NEXT:    mov.b32 %r51, %r19;
-; CHECK-NEXT:    @%p8 bra $L__BB4_10;
-; CHECK-NEXT:  $L__BB4_12: // %partword.cmpxchg.end1
+; CHECK-NEXT:    @%p8 bra $L__BB3_10;
+; CHECK-NEXT:  $L__BB3_12: // %partword.cmpxchg.end1
 ; CHECK-NEXT:    fence.acq_rel.sys;
 ; CHECK-NEXT:    fence.sc.sys;
 ; CHECK-NEXT:    ld.shared::cluster.u32 %r46, [%rd1];
 ; CHECK-NEXT:    and.b32 %r52, %r46, %r2;
-; CHECK-NEXT:  $L__BB4_13: // %partword.cmpxchg.loop
+; CHECK-NEXT:  $L__BB3_13: // %partword.cmpxchg.loop
 ; CHECK-NEXT:    // =>This Inner Loop Header: Depth=1
 ; CHECK-NEXT:    or.b32 %r47, %r52, %r3;
 ; CHECK-NEXT:    atom.relaxed.shared::cluster.cas.b32 %r22, [%rd1], %r47, %r52;
 ; CHECK-NEXT:    setp.eq.s32 %p9, %r22, %r47;
-; CHECK-NEXT:    @%p9 bra $L__BB4_15;
+; CHECK-NEXT:    @%p9 bra $L__BB3_15;
 ; CHECK-NEXT:  // %bb.14: // %partword.cmpxchg.failure
-; CHECK-NEXT:    // in Loop: Header=BB4_13 Depth=1
+; CHECK-NEXT:    // in Loop: Header=BB3_13 Depth=1
 ; CHECK-NEXT:    and.b32 %r23, %r22, %r2;
 ; CHECK-NEXT:    setp.ne.s32 %p10, %r52, %r23;
 ; CHECK-NEXT:    mov.b32 %r52, %r23;
-; CHECK-NEXT:    @%p10 bra $L__BB4_13;
-; CHECK-NEXT:  $L__BB4_15: // %partword.cmpxchg.end
+; CHECK-NEXT:    @%p10 bra $L__BB3_13;
+; CHECK-NEXT:  $L__BB3_15: // %partword.cmpxchg.end
 ; CHECK-NEXT:    fence.acq_rel.sys;
 ; CHECK-NEXT:    ret;
 entry: