109 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}4 5target datalayout = "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"6 7define <16 x float> @test_v16f32(<16 x float> %a) {8; CHECK-LABEL: test_v16f32(9; CHECK: {10; CHECK-NEXT: .reg .b32 %r<17>;11; CHECK-EMPTY:12; CHECK-NEXT: // %bb.0:13; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_v16f32_param_0];14; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [test_v16f32_param_0+16];15; CHECK-NEXT: ld.param.v4.b32 {%r9, %r10, %r11, %r12}, [test_v16f32_param_0+32];16; CHECK-NEXT: ld.param.v4.b32 {%r13, %r14, %r15, %r16}, [test_v16f32_param_0+48];17; CHECK-NEXT: st.param.v4.b32 [func_retval0+48], {%r13, %r14, %r15, %r16};18; CHECK-NEXT: st.param.v4.b32 [func_retval0+32], {%r9, %r10, %r11, %r12};19; CHECK-NEXT: st.param.v4.b32 [func_retval0+16], {%r5, %r6, %r7, %r8};20; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r1, %r2, %r3, %r4};21; CHECK-NEXT: ret;22 ret <16 x float> %a23}24 25define <8 x float> @test_v8f32(<8 x float> %a) {26; CHECK-LABEL: test_v8f32(27; CHECK: {28; CHECK-NEXT: .reg .b32 %r<9>;29; CHECK-EMPTY:30; CHECK-NEXT: // %bb.0:31; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_v8f32_param_0];32; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [test_v8f32_param_0+16];33; CHECK-NEXT: st.param.v4.b32 [func_retval0+16], {%r5, %r6, %r7, %r8};34; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r1, %r2, %r3, %r4};35; CHECK-NEXT: ret;36 ret <8 x float> %a37}38 39define <4 x float> @test_v4f32(<4 x float> %a) {40; CHECK-LABEL: test_v4f32(41; CHECK: {42; CHECK-NEXT: .reg .b32 %r<5>;43; CHECK-EMPTY:44; CHECK-NEXT: // %bb.0:45; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_v4f32_param_0];46; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r1, %r2, %r3, %r4};47; CHECK-NEXT: ret;48 ret <4 x float> %a49}50 51define <2 x float> @test_v2f32(<2 x float> %a) {52; CHECK-LABEL: test_v2f32(53; CHECK: {54; CHECK-NEXT: .reg .b32 %r<3>;55; CHECK-EMPTY:56; CHECK-NEXT: // %bb.0:57; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_v2f32_param_0];58; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r1, %r2};59; CHECK-NEXT: ret;60 ret <2 x float> %a61}62 63; Oddly shaped vectors should not load any extra elements.64define <3 x float> @test_v3f32(<3 x float> %a) {65; CHECK-LABEL: test_v3f32(66; CHECK: {67; CHECK-NEXT: .reg .b32 %r<4>;68; CHECK-EMPTY:69; CHECK-NEXT: // %bb.0:70; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_v3f32_param_0];71; CHECK-NEXT: ld.param.b32 %r3, [test_v3f32_param_0+8];72; CHECK-NEXT: st.param.b32 [func_retval0+8], %r3;73; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r1, %r2};74; CHECK-NEXT: ret;75 ret <3 x float> %a76}77 78define <8 x i64> @test_v8i64(<8 x i64> %a) {79; CHECK-LABEL: test_v8i64(80; CHECK: {81; CHECK-NEXT: .reg .b64 %rd<9>;82; CHECK-EMPTY:83; CHECK-NEXT: // %bb.0:84; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_v8i64_param_0];85; CHECK-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [test_v8i64_param_0+16];86; CHECK-NEXT: ld.param.v2.b64 {%rd5, %rd6}, [test_v8i64_param_0+32];87; CHECK-NEXT: ld.param.v2.b64 {%rd7, %rd8}, [test_v8i64_param_0+48];88; CHECK-NEXT: st.param.v2.b64 [func_retval0+48], {%rd7, %rd8};89; CHECK-NEXT: st.param.v2.b64 [func_retval0+32], {%rd5, %rd6};90; CHECK-NEXT: st.param.v2.b64 [func_retval0+16], {%rd3, %rd4};91; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd1, %rd2};92; CHECK-NEXT: ret;93 ret <8 x i64> %a94}95 96define <16 x i16> @test_v16i16(<16 x i16> %a) {97; CHECK-LABEL: test_v16i16(98; CHECK: {99; CHECK-NEXT: .reg .b32 %r<9>;100; CHECK-EMPTY:101; CHECK-NEXT: // %bb.0:102; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_v16i16_param_0];103; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [test_v16i16_param_0+16];104; CHECK-NEXT: st.param.v4.b32 [func_retval0+16], {%r5, %r6, %r7, %r8};105; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r1, %r2, %r3, %r4};106; CHECK-NEXT: ret;107 ret <16 x i16> %a108}109