133 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@value = internal addrspace(1) global i128 0, align 168 9define void @test_b128_input_from_const() {10; CHECK-LABEL: test_b128_input_from_const(11; CHECK: {12; CHECK-NEXT: .reg .b64 %rd<5>;13; CHECK-NEXT: .reg .b128 %rq<2>;14; CHECK-EMPTY:15; CHECK-NEXT: // %bb.0:16; CHECK-NEXT: mov.b64 %rd2, 0;17; CHECK-NEXT: mov.b64 %rd3, 42;18; CHECK-NEXT: mov.b128 %rq1, {%rd3, %rd2};19; CHECK-NEXT: mov.b64 %rd4, value;20; CHECK-NEXT: cvta.global.u64 %rd1, %rd4;21; CHECK-NEXT: // begin inline asm22; CHECK-NEXT: { st.b128 [%rd1], %rq1; }23; CHECK-NEXT: // end inline asm24; CHECK-NEXT: ret;25 tail call void asm sideeffect "{ st.b128 [$0], $1; }", "l,q"(ptr nonnull addrspacecast (ptr addrspace(1) @value to ptr), i128 42)26 ret void27}28 29define void @test_b128_input_from_load(ptr nocapture readonly %data) {30; CHECK-LABEL: test_b128_input_from_load(31; CHECK: {32; CHECK-NEXT: .reg .b64 %rd<7>;33; CHECK-NEXT: .reg .b128 %rq<2>;34; CHECK-EMPTY:35; CHECK-NEXT: // %bb.0:36; CHECK-NEXT: ld.param.b64 %rd2, [test_b128_input_from_load_param_0];37; CHECK-NEXT: cvta.to.global.u64 %rd3, %rd2;38; CHECK-NEXT: ld.global.v2.b64 {%rd4, %rd5}, [%rd3];39; CHECK-NEXT: mov.b64 %rd6, value;40; CHECK-NEXT: cvta.global.u64 %rd1, %rd6;41; CHECK-NEXT: mov.b128 %rq1, {%rd4, %rd5};42; CHECK-NEXT: // begin inline asm43; CHECK-NEXT: { st.b128 [%rd1], %rq1; }44; CHECK-NEXT: // end inline asm45; CHECK-NEXT: ret;46 %1 = addrspacecast ptr %data to ptr addrspace(1)47 %2 = load <2 x i64>, ptr addrspace(1) %1, align 1648 %3 = bitcast <2 x i64> %2 to i12849 tail call void asm sideeffect "{ st.b128 [$0], $1; }", "l,q"(ptr nonnull addrspacecast (ptr addrspace(1) @value to ptr), i128 %3)50 ret void51}52 53define void @test_b128_input_from_select(ptr nocapture readonly %flag) {54; CHECK-LABEL: test_b128_input_from_select(55; CHECK: {56; CHECK-NEXT: .reg .pred %p<2>;57; CHECK-NEXT: .reg .b16 %rs<2>;58; CHECK-NEXT: .reg .b64 %rd<7>;59; CHECK-NEXT: .reg .b128 %rq<2>;60; CHECK-EMPTY:61; CHECK-NEXT: // %bb.0:62; CHECK-NEXT: ld.param.b64 %rd2, [test_b128_input_from_select_param_0];63; CHECK-NEXT: cvta.to.global.u64 %rd3, %rd2;64; CHECK-NEXT: ld.global.b8 %rs1, [%rd3];65; CHECK-NEXT: setp.eq.b16 %p1, %rs1, 0;66; CHECK-NEXT: selp.b64 %rd4, 24, 42, %p1;67; CHECK-NEXT: mov.b64 %rd5, 0;68; CHECK-NEXT: mov.b128 %rq1, {%rd4, %rd5};69; CHECK-NEXT: mov.b64 %rd6, value;70; CHECK-NEXT: cvta.global.u64 %rd1, %rd6;71; CHECK-NEXT: // begin inline asm72; CHECK-NEXT: { st.b128 [%rd1], %rq1; }73; CHECK-NEXT: // end inline asm74; CHECK-NEXT: ret;75 %1 = addrspacecast ptr %flag to ptr addrspace(1)76 %2 = load i8, ptr addrspace(1) %1, align 177 %3 = icmp eq i8 %2, 078 %4 = select i1 %3, i128 24, i128 4279 tail call void asm sideeffect "{ st.b128 [$0], $1; }", "l,q"(ptr nonnull addrspacecast (ptr addrspace(1) @value to ptr), i128 %4)80 ret void81}82 83define void @test_store_b128_output() {84; CHECK-LABEL: test_store_b128_output(85; CHECK: {86; CHECK-NEXT: .reg .b64 %rd<5>;87; CHECK-NEXT: .reg .b128 %rq<2>;88; CHECK-EMPTY:89; CHECK-NEXT: // %bb.0:90; CHECK-NEXT: // begin inline asm91; CHECK-NEXT: { mov.b128 %rq1, 41; }92; CHECK-NEXT: // end inline asm93; CHECK-NEXT: mov.b128 {%rd1, %rd2}, %rq1;94; CHECK-NEXT: add.cc.s64 %rd3, %rd1, 1;95; CHECK-NEXT: addc.cc.s64 %rd4, %rd2, 0;96; CHECK-NEXT: st.global.v2.b64 [value], {%rd3, %rd4};97; CHECK-NEXT: ret;98 %1 = tail call i128 asm "{ mov.b128 $0, 41; }", "=q"()99 %add = add nsw i128 %1, 1100 %2 = bitcast i128 %add to <2 x i64>101 store <2 x i64> %2, ptr addrspace(1) @value, align 16102 ret void103}104 105define void @test_use_of_b128_output(ptr nocapture readonly %data) {106; CHECK-LABEL: test_use_of_b128_output(107; CHECK: {108; CHECK-NEXT: .reg .b64 %rd<9>;109; CHECK-NEXT: .reg .b128 %rq<3>;110; CHECK-EMPTY:111; CHECK-NEXT: // %bb.0:112; CHECK-NEXT: ld.param.b64 %rd1, [test_use_of_b128_output_param_0];113; CHECK-NEXT: cvta.to.global.u64 %rd2, %rd1;114; CHECK-NEXT: ld.global.v2.b64 {%rd3, %rd4}, [%rd2];115; CHECK-NEXT: mov.b128 %rq2, {%rd3, %rd4};116; CHECK-NEXT: // begin inline asm117; CHECK-NEXT: { mov.b128 %rq1, %rq2; }118; CHECK-NEXT: // end inline asm119; CHECK-NEXT: mov.b128 {%rd5, %rd6}, %rq1;120; CHECK-NEXT: add.cc.s64 %rd7, %rd5, 1;121; CHECK-NEXT: addc.cc.s64 %rd8, %rd6, 0;122; CHECK-NEXT: st.global.v2.b64 [value], {%rd7, %rd8};123; CHECK-NEXT: ret;124 %1 = addrspacecast ptr %data to ptr addrspace(1)125 %2 = load <2 x i64>, ptr addrspace(1) %1, align 16126 %3 = bitcast <2 x i64> %2 to i128127 %4 = tail call i128 asm "{ mov.b128 $0, $1; }", "=q,q"(i128 %3)128 %add = add nsw i128 %4, 1129 %5 = bitcast i128 %add to <2 x i64>130 store <2 x i64> %5, ptr addrspace(1) @value, align 16131 ret void132}133