| ; 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=+ptx86 -combiner-disabled -O0 | FileCheck %s |
| ; RUN: %if ptxas-sm_100 && ptxas-isa-8.6 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx86 -combiner-disabled -O0 | %ptxas-verify -arch=sm_100 %} |
| |
| define void @test_insertelement_int32x2() { |
| ; CHECK-LABEL: test_insertelement_int32x2( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<3>; |
| ; CHECK-NEXT: .reg .b64 %rd<3>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: mov.b32 %r1, 0; |
| ; CHECK-NEXT: cvt.u64.u32 %rd2, %r1; |
| ; CHECK-NEXT: mov.b64 {%r2, _}, %rd2; |
| ; CHECK-NEXT: mov.b64 %rd1, {%r2, %r1}; |
| ; CHECK-NEXT: bra.uni $L__BB0_1; |
| ; CHECK-NEXT: $L__BB0_1: |
| ; CHECK-NEXT: ret; |
| %1 = insertelement <2 x i32> zeroinitializer, i32 0, i32 0 |
| br label %2 |
| 2: ; preds = %0 |
| %3 = extractelement <2 x i32> %1, i32 0 |
| %4 = bitcast <2 x i32> %1 to i64 |
| ret void |
| } |
| |
| define void @test_insertelement_float32x2() { |
| ; CHECK-LABEL: test_insertelement_float32x2( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<3>; |
| ; CHECK-NEXT: .reg .b64 %rd<3>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: mov.b32 %r1, 0f00000000; |
| ; CHECK-NEXT: mov.b64 %rd2, {%r1, 0}; |
| ; CHECK-NEXT: mov.b64 {%r2, _}, %rd2; |
| ; CHECK-NEXT: mov.b64 %rd1, {%r2, %r1}; |
| ; CHECK-NEXT: bra.uni $L__BB1_1; |
| ; CHECK-NEXT: $L__BB1_1: |
| ; CHECK-NEXT: ret; |
| %1 = insertelement <2 x float> zeroinitializer, float 0.0, i32 0 |
| br label %2 |
| 2: ; preds = %0 |
| %3 = extractelement <2 x float> %1, i32 0 |
| %4 = bitcast <2 x float> %1 to i64 |
| ret void |
| } |
| |
| define void @test_insertelement_float16x2() { |
| ; CHECK-LABEL: test_insertelement_float16x2( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b16 %rs<4>; |
| ; CHECK-NEXT: .reg .b32 %r<4>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: mov.b32 %r2, 0; |
| ; CHECK-NEXT: mov.b32 {_, %rs1}, %r2; |
| ; CHECK-NEXT: mov.b16 %rs2, 0x0000; |
| ; CHECK-NEXT: mov.b32 %r3, {%rs2, 0}; |
| ; CHECK-NEXT: mov.b32 {%rs3, _}, %r3; |
| ; CHECK-NEXT: mov.b32 %r1, {%rs3, %rs1}; |
| ; CHECK-NEXT: bra.uni $L__BB2_1; |
| ; CHECK-NEXT: $L__BB2_1: |
| ; CHECK-NEXT: ret; |
| %1 = insertelement <2 x half> zeroinitializer, half 0.0, i32 0 |
| br label %2 |
| 2: ; preds = %0 |
| %3 = extractelement <2 x half> %1, i32 0 |
| %4 = bitcast <2 x half> %1 to i32 |
| ret void |
| } |