66 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --extra_scrub --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_70 -mattr=+ptx83 | FileCheck %s3; RUN: %if ptxas-sm_70 && ptxas-isa-8.3 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_70 -mattr=+ptx83 | %ptxas-verify -arch=sm_70 %}4 5target triple = "nvptx64-nvidia-cuda"6 7@size = internal addrspace(1) global i32 0, align 48@x = internal addrspace(1) global i128 0, align 169 10define void @test_b128_in_loop() {11; CHECK-LABEL: test_b128_in_loop(12; CHECK: {13; CHECK-NEXT: .reg .pred %p<3>;14; CHECK-NEXT: .reg .b64 %rd<5>;15; CHECK-NEXT: .reg .b128 %rq<2>;16; CHECK-EMPTY:17; CHECK-NEXT: // %bb.0:18; CHECK-NEXT: ld.global.s32 %rd1, [size];19; CHECK-NEXT: setp.eq.b64 %p1, %rd1, 0;20; CHECK-NEXT: @%p1 bra $L__BB0_3;21; CHECK-NEXT: // %bb.1: // %BB122; CHECK-NEXT: ld.global.v2.b64 {%rd2, %rd3}, [x];23; CHECK-NEXT: mov.b64 %rd4, 0;24; CHECK-NEXT: $L__BB0_2: // %BB225; CHECK-NEXT: // =>This Inner Loop Header: Depth=126; CHECK-NEXT: mov.b128 %rq1, {%rd2, %rd3};27; CHECK-NEXT: // begin inline asm28; CHECK-NEXT: {29; CHECK-NEXT: .reg .b64 lo;30; CHECK-NEXT: .reg .b64 hi;31; CHECK-NEXT: mov.b128 {lo, hi}, %rq1;32; CHECK-NEXT: add.cc.u64 lo, lo, %rd4;33; CHECK-NEXT: mov.b128 %rq1, {lo, hi};34; CHECK-NEXT: }35; CHECK-NEXT: // end inline asm36; CHECK-NEXT: mov.b128 {%rd2, %rd3}, %rq1;37; CHECK-NEXT: st.global.v2.b64 [x], {%rd2, %rd3};38; CHECK-NEXT: add.s64 %rd4, %rd4, 1;39; CHECK-NEXT: setp.ne.b64 %p2, %rd1, %rd4;40; CHECK-NEXT: @%p2 bra $L__BB0_2;41; CHECK-NEXT: $L__BB0_3: // %BB342; CHECK-NEXT: ret;43 44 %1 = load i32, ptr addrspace(1) @size, align 445 %2 = icmp eq i32 %1, 046 br i1 %2, label %BB3, label %BB147 48BB1: ; preds = %049 %3 = load i128, ptr addrspace(1) @x, align 1650 %4 = sext i32 %1 to i6451 br label %BB252 53BB2: ; preds = %BB2, %BB154 %5 = phi i128 [ %7, %BB2 ], [ %3, %BB1 ]55 %6 = phi i64 [ %9, %BB2 ], [ 0, %BB1 ]56 %7 = tail call i128 asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, $1;\0A\09mov.b128 $0, {lo, hi};\0A\09}", "=q,l,0"(i64 %6, i128 %5)57 %8 = bitcast i128 %7 to <2 x i64>58 store <2 x i64> %8, ptr addrspace(1) @x, align 1659 %9 = add nuw i64 %6, 160 %10 = icmp eq i64 %9, %461 br i1 %10, label %BB3, label %BB262 63BB3: ; preds = %BB2, %064 ret void65}66