1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 2; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s 3 4target triple = "nvptx-nvidia-cuda" 5 6define <3 x i64> @long3() { 7; CHECK-LABEL: long3( 8; CHECK: { 9; CHECK-NEXT: .reg .b64 %rd<2>; 10; CHECK-EMPTY: 11; CHECK-NEXT: // %bb.0: 12; CHECK-NEXT: mov.b64 %rd1, 0; 13; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd1, %rd1}; 14; CHECK-NEXT: st.param.b64 [func_retval0+16], %rd1; 15; CHECK-NEXT: ret; 16 ret <3 x i64> zeroinitializer 17} 18 19define <2 x i64> @long2() { 20; CHECK-LABEL: long2( 21; CHECK: { 22; CHECK-NEXT: .reg .b64 %rd<2>; 23; CHECK-EMPTY: 24; CHECK-NEXT: // %bb.0: 25; CHECK-NEXT: mov.b64 %rd1, 0; 26; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd1, %rd1}; 27; CHECK-NEXT: ret; 28 ret <2 x i64> zeroinitializer 29} 30 31define <1 x i64> @long1() { 32; CHECK-LABEL: long1( 33; CHECK: { 34; CHECK-NEXT: .reg .b64 %rd<2>; 35; CHECK-EMPTY: 36; CHECK-NEXT: // %bb.0: 37; CHECK-NEXT: mov.b64 %rd1, 0; 38; CHECK-NEXT: st.param.b64 [func_retval0], %rd1; 39; CHECK-NEXT: ret; 40 ret <1 x i64> zeroinitializer 41} 42 43define <5 x i32> @int5() { 44; CHECK-LABEL: int5( 45; CHECK: { 46; CHECK-NEXT: .reg .b32 %r<2>; 47; CHECK-EMPTY: 48; CHECK-NEXT: // %bb.0: 49; CHECK-NEXT: mov.b32 %r1, 0; 50; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r1, %r1, %r1, %r1}; 51; CHECK-NEXT: st.param.b32 [func_retval0+16], %r1; 52; CHECK-NEXT: ret; 53 ret <5 x i32> zeroinitializer 54} 55 56define <4 x i32> @int4() { 57; CHECK-LABEL: int4( 58; CHECK: { 59; CHECK-NEXT: .reg .b32 %r<2>; 60; CHECK-EMPTY: 61; CHECK-NEXT: // %bb.0: 62; CHECK-NEXT: mov.b32 %r1, 0; 63; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r1, %r1, %r1, %r1}; 64; CHECK-NEXT: ret; 65 ret <4 x i32> zeroinitializer 66} 67 68define <3 x i32> @int3() { 69; CHECK-LABEL: int3( 70; CHECK: { 71; CHECK-NEXT: .reg .b32 %r<2>; 72; CHECK-EMPTY: 73; CHECK-NEXT: // %bb.0: 74; CHECK-NEXT: mov.b32 %r1, 0; 75; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r1, %r1}; 76; CHECK-NEXT: st.param.b32 [func_retval0+8], %r1; 77; CHECK-NEXT: ret; 78 ret <3 x i32> zeroinitializer 79} 80 81define <2 x i32> @int2() { 82; CHECK-LABEL: int2( 83; CHECK: { 84; CHECK-NEXT: .reg .b32 %r<2>; 85; CHECK-EMPTY: 86; CHECK-NEXT: // %bb.0: 87; CHECK-NEXT: mov.b32 %r1, 0; 88; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r1, %r1}; 89; CHECK-NEXT: ret; 90 ret <2 x i32> zeroinitializer 91} 92 93define <1 x i32> @int1() { 94; CHECK-LABEL: int1( 95; CHECK: { 96; CHECK-NEXT: .reg .b32 %r<2>; 97; CHECK-EMPTY: 98; CHECK-NEXT: // %bb.0: 99; CHECK-NEXT: mov.b32 %r1, 0; 100; CHECK-NEXT: st.param.b32 [func_retval0], %r1; 101; CHECK-NEXT: ret; 102 ret <1 x i32> zeroinitializer 103} 104 105define <9 x i16> @short9() { 106; CHECK-LABEL: short9( 107; CHECK: { 108; CHECK-NEXT: .reg .b16 %rs<2>; 109; CHECK-EMPTY: 110; CHECK-NEXT: // %bb.0: 111; CHECK-NEXT: mov.b16 %rs1, 0; 112; CHECK-NEXT: st.param.v4.b16 [func_retval0], {%rs1, %rs1, %rs1, %rs1}; 113; CHECK-NEXT: st.param.v4.b16 [func_retval0+8], {%rs1, %rs1, %rs1, %rs1}; 114; CHECK-NEXT: st.param.b16 [func_retval0+16], %rs1; 115; CHECK-NEXT: ret; 116 ret <9 x i16> zeroinitializer 117} 118 119define <8 x i16> @short8() { 120; CHECK-LABEL: short8( 121; CHECK: { 122; CHECK-NEXT: .reg .b32 %r<2>; 123; CHECK-EMPTY: 124; CHECK-NEXT: // %bb.0: 125; CHECK-NEXT: mov.b32 %r1, 0; 126; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r1, %r1, %r1, %r1}; 127; CHECK-NEXT: ret; 128 ret <8 x i16> zeroinitializer 129} 130 131define <7 x i16> @short7() { 132; CHECK-LABEL: short7( 133; CHECK: { 134; CHECK-NEXT: .reg .b16 %rs<2>; 135; CHECK-EMPTY: 136; CHECK-NEXT: // %bb.0: 137; CHECK-NEXT: mov.b16 %rs1, 0; 138; CHECK-NEXT: st.param.v4.b16 [func_retval0], {%rs1, %rs1, %rs1, %rs1}; 139; CHECK-NEXT: st.param.v2.b16 [func_retval0+8], {%rs1, %rs1}; 140; CHECK-NEXT: st.param.b16 [func_retval0+12], %rs1; 141; CHECK-NEXT: ret; 142 ret <7 x i16> zeroinitializer 143} 144 145define <5 x i16> @short5() { 146; CHECK-LABEL: short5( 147; CHECK: { 148; CHECK-NEXT: .reg .b16 %rs<2>; 149; CHECK-EMPTY: 150; CHECK-NEXT: // %bb.0: 151; CHECK-NEXT: mov.b16 %rs1, 0; 152; CHECK-NEXT: st.param.v4.b16 [func_retval0], {%rs1, %rs1, %rs1, %rs1}; 153; CHECK-NEXT: st.param.b16 [func_retval0+8], %rs1; 154; CHECK-NEXT: ret; 155 ret <5 x i16> zeroinitializer 156} 157 158define <4 x i16> @short4() { 159; CHECK-LABEL: short4( 160; CHECK: { 161; CHECK-NEXT: .reg .b32 %r<2>; 162; CHECK-EMPTY: 163; CHECK-NEXT: // %bb.0: 164; CHECK-NEXT: mov.b32 %r1, 0; 165; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r1, %r1}; 166; CHECK-NEXT: ret; 167 ret <4 x i16> zeroinitializer 168} 169 170define <3 x i16> @short3() { 171; CHECK-LABEL: short3( 172; CHECK: { 173; CHECK-NEXT: .reg .b16 %rs<2>; 174; CHECK-EMPTY: 175; CHECK-NEXT: // %bb.0: 176; CHECK-NEXT: mov.b16 %rs1, 0; 177; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs1, %rs1}; 178; CHECK-NEXT: st.param.b16 [func_retval0+4], %rs1; 179; CHECK-NEXT: ret; 180 ret <3 x i16> zeroinitializer 181} 182 183define <2 x i16> @short2() { 184; CHECK-LABEL: short2( 185; CHECK: { 186; CHECK-NEXT: .reg .b32 %r<2>; 187; CHECK-EMPTY: 188; CHECK-NEXT: // %bb.0: 189; CHECK-NEXT: mov.b32 %r1, 0; 190; CHECK-NEXT: st.param.b32 [func_retval0], %r1; 191; CHECK-NEXT: ret; 192 ret <2 x i16> zeroinitializer 193} 194 195define <1 x i16> @short1() { 196; CHECK-LABEL: short1( 197; CHECK: { 198; CHECK-NEXT: .reg .b16 %rs<2>; 199; CHECK-EMPTY: 200; CHECK-NEXT: // %bb.0: 201; CHECK-NEXT: mov.b16 %rs1, 0; 202; CHECK-NEXT: st.param.b16 [func_retval0], %rs1; 203; CHECK-NEXT: ret; 204 ret <1 x i16> zeroinitializer 205} 206 207define <17 x i8> @byte17() { 208; CHECK-LABEL: byte17( 209; CHECK: { 210; CHECK-NEXT: .reg .b16 %rs<2>; 211; CHECK-EMPTY: 212; CHECK-NEXT: // %bb.0: 213; CHECK-NEXT: mov.b16 %rs1, 0; 214; CHECK-NEXT: st.param.v4.b8 [func_retval0], {%rs1, %rs1, %rs1, %rs1}; 215; CHECK-NEXT: st.param.v4.b8 [func_retval0+4], {%rs1, %rs1, %rs1, %rs1}; 216; CHECK-NEXT: st.param.v4.b8 [func_retval0+8], {%rs1, %rs1, %rs1, %rs1}; 217; CHECK-NEXT: st.param.v4.b8 [func_retval0+12], {%rs1, %rs1, %rs1, %rs1}; 218; CHECK-NEXT: st.param.b8 [func_retval0+16], %rs1; 219; CHECK-NEXT: ret; 220 ret <17 x i8> zeroinitializer 221} 222 223define <16 x i8> @byte16() { 224; CHECK-LABEL: byte16( 225; CHECK: { 226; CHECK-NEXT: .reg .b32 %r<2>; 227; CHECK-EMPTY: 228; CHECK-NEXT: // %bb.0: 229; CHECK-NEXT: mov.b32 %r1, 0; 230; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r1, %r1, %r1, %r1}; 231; CHECK-NEXT: ret; 232 ret <16 x i8> zeroinitializer 233} 234 235define <15 x i8> @byte15() { 236; CHECK-LABEL: byte15( 237; CHECK: { 238; CHECK-NEXT: .reg .b16 %rs<2>; 239; CHECK-EMPTY: 240; CHECK-NEXT: // %bb.0: 241; CHECK-NEXT: mov.b16 %rs1, 0; 242; CHECK-NEXT: st.param.v4.b8 [func_retval0], {%rs1, %rs1, %rs1, %rs1}; 243; CHECK-NEXT: st.param.v4.b8 [func_retval0+4], {%rs1, %rs1, %rs1, %rs1}; 244; CHECK-NEXT: st.param.v4.b8 [func_retval0+8], {%rs1, %rs1, %rs1, %rs1}; 245; CHECK-NEXT: st.param.v2.b8 [func_retval0+12], {%rs1, %rs1}; 246; CHECK-NEXT: st.param.b8 [func_retval0+14], %rs1; 247; CHECK-NEXT: ret; 248 ret <15 x i8> zeroinitializer 249} 250 251define <9 x i8> @byte9() { 252; CHECK-LABEL: byte9( 253; CHECK: { 254; CHECK-NEXT: .reg .b16 %rs<2>; 255; CHECK-EMPTY: 256; CHECK-NEXT: // %bb.0: 257; CHECK-NEXT: mov.b16 %rs1, 0; 258; CHECK-NEXT: st.param.v4.b8 [func_retval0], {%rs1, %rs1, %rs1, %rs1}; 259; CHECK-NEXT: st.param.v4.b8 [func_retval0+4], {%rs1, %rs1, %rs1, %rs1}; 260; CHECK-NEXT: st.param.b8 [func_retval0+8], %rs1; 261; CHECK-NEXT: ret; 262 ret <9 x i8> zeroinitializer 263} 264 265define <8 x i8> @byte8() { 266; CHECK-LABEL: byte8( 267; CHECK: { 268; CHECK-NEXT: .reg .b32 %r<2>; 269; CHECK-EMPTY: 270; CHECK-NEXT: // %bb.0: 271; CHECK-NEXT: mov.b32 %r1, 0; 272; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r1, %r1}; 273; CHECK-NEXT: ret; 274 ret <8 x i8> zeroinitializer 275} 276 277define <7 x i8> @byte7() { 278; CHECK-LABEL: byte7( 279; CHECK: { 280; CHECK-NEXT: .reg .b16 %rs<2>; 281; CHECK-EMPTY: 282; CHECK-NEXT: // %bb.0: 283; CHECK-NEXT: mov.b16 %rs1, 0; 284; CHECK-NEXT: st.param.v4.b8 [func_retval0], {%rs1, %rs1, %rs1, %rs1}; 285; CHECK-NEXT: st.param.v2.b8 [func_retval0+4], {%rs1, %rs1}; 286; CHECK-NEXT: st.param.b8 [func_retval0+6], %rs1; 287; CHECK-NEXT: ret; 288 ret <7 x i8> zeroinitializer 289} 290 291define <5 x i8> @byte5() { 292; CHECK-LABEL: byte5( 293; CHECK: { 294; CHECK-NEXT: .reg .b16 %rs<2>; 295; CHECK-EMPTY: 296; CHECK-NEXT: // %bb.0: 297; CHECK-NEXT: mov.b16 %rs1, 0; 298; CHECK-NEXT: st.param.v4.b8 [func_retval0], {%rs1, %rs1, %rs1, %rs1}; 299; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs1; 300; CHECK-NEXT: ret; 301 ret <5 x i8> zeroinitializer 302} 303 304define <4 x i8> @byte4() { 305; CHECK-LABEL: byte4( 306; CHECK: { 307; CHECK-NEXT: .reg .b32 %r<2>; 308; CHECK-EMPTY: 309; CHECK-NEXT: // %bb.0: 310; CHECK-NEXT: mov.b32 %r1, 0; 311; CHECK-NEXT: st.param.b32 [func_retval0], %r1; 312; CHECK-NEXT: ret; 313 ret <4 x i8> zeroinitializer 314} 315 316define <3 x i8> @byte3() { 317; CHECK-LABEL: byte3( 318; CHECK: { 319; CHECK-NEXT: .reg .b32 %r<2>; 320; CHECK-EMPTY: 321; CHECK-NEXT: // %bb.0: 322; CHECK-NEXT: mov.b32 %r1, 0; 323; CHECK-NEXT: st.param.b32 [func_retval0], %r1; 324; CHECK-NEXT: ret; 325 ret <3 x i8> zeroinitializer 326} 327 328define <2 x i8> @byte2() { 329; CHECK-LABEL: byte2( 330; CHECK: { 331; CHECK-NEXT: .reg .b32 %r<2>; 332; CHECK-EMPTY: 333; CHECK-NEXT: // %bb.0: 334; CHECK-NEXT: mov.b32 %r1, 0; 335; CHECK-NEXT: st.param.b32 [func_retval0], %r1; 336; CHECK-NEXT: ret; 337 ret <2 x i8> zeroinitializer 338} 339 340define <1 x i8> @byte1() { 341; CHECK-LABEL: byte1( 342; CHECK: { 343; CHECK-NEXT: .reg .b16 %rs<2>; 344; CHECK-EMPTY: 345; CHECK-NEXT: // %bb.0: 346; CHECK-NEXT: mov.b16 %rs1, 0; 347; CHECK-NEXT: st.param.b8 [func_retval0], %rs1; 348; CHECK-NEXT: ret; 349 ret <1 x i8> zeroinitializer 350} 351 352define <17 x i1> @bit17() { 353; CHECK-LABEL: bit17( 354; CHECK: { 355; CHECK-NEXT: .reg .b16 %rs<2>; 356; CHECK-EMPTY: 357; CHECK-NEXT: // %bb.0: 358; CHECK-NEXT: mov.b16 %rs1, 0; 359; CHECK-NEXT: st.param.v4.b8 [func_retval0], {%rs1, %rs1, %rs1, %rs1}; 360; CHECK-NEXT: st.param.v4.b8 [func_retval0+4], {%rs1, %rs1, %rs1, %rs1}; 361; CHECK-NEXT: st.param.v4.b8 [func_retval0+8], {%rs1, %rs1, %rs1, %rs1}; 362; CHECK-NEXT: st.param.v4.b8 [func_retval0+12], {%rs1, %rs1, %rs1, %rs1}; 363; CHECK-NEXT: st.param.b8 [func_retval0+16], %rs1; 364; CHECK-NEXT: ret; 365 ret <17 x i1> zeroinitializer 366} 367 368define <16 x i1> @bit16() { 369; CHECK-LABEL: bit16( 370; CHECK: { 371; CHECK-NEXT: .reg .b16 %rs<2>; 372; CHECK-EMPTY: 373; CHECK-NEXT: // %bb.0: 374; CHECK-NEXT: mov.b16 %rs1, 0; 375; CHECK-NEXT: st.param.v2.b8 [func_retval0], {%rs1, %rs1}; 376; CHECK-NEXT: st.param.v2.b8 [func_retval0+2], {%rs1, %rs1}; 377; CHECK-NEXT: st.param.v2.b8 [func_retval0+4], {%rs1, %rs1}; 378; CHECK-NEXT: st.param.v2.b8 [func_retval0+6], {%rs1, %rs1}; 379; CHECK-NEXT: st.param.v2.b8 [func_retval0+8], {%rs1, %rs1}; 380; CHECK-NEXT: st.param.v2.b8 [func_retval0+10], {%rs1, %rs1}; 381; CHECK-NEXT: st.param.v2.b8 [func_retval0+12], {%rs1, %rs1}; 382; CHECK-NEXT: st.param.v2.b8 [func_retval0+14], {%rs1, %rs1}; 383; CHECK-NEXT: ret; 384 ret <16 x i1> zeroinitializer 385} 386 387define <15 x i1> @bit15() { 388; CHECK-LABEL: bit15( 389; CHECK: { 390; CHECK-NEXT: .reg .b16 %rs<2>; 391; CHECK-EMPTY: 392; CHECK-NEXT: // %bb.0: 393; CHECK-NEXT: mov.b16 %rs1, 0; 394; CHECK-NEXT: st.param.v2.b8 [func_retval0], {%rs1, %rs1}; 395; CHECK-NEXT: st.param.v2.b8 [func_retval0+2], {%rs1, %rs1}; 396; CHECK-NEXT: st.param.v2.b8 [func_retval0+4], {%rs1, %rs1}; 397; CHECK-NEXT: st.param.v2.b8 [func_retval0+6], {%rs1, %rs1}; 398; CHECK-NEXT: st.param.v2.b8 [func_retval0+8], {%rs1, %rs1}; 399; CHECK-NEXT: st.param.v2.b8 [func_retval0+10], {%rs1, %rs1}; 400; CHECK-NEXT: st.param.v2.b8 [func_retval0+12], {%rs1, %rs1}; 401; CHECK-NEXT: st.param.b8 [func_retval0+14], %rs1; 402; CHECK-NEXT: ret; 403 ret <15 x i1> zeroinitializer 404} 405 406define <9 x i1> @bit9() { 407; CHECK-LABEL: bit9( 408; CHECK: { 409; CHECK-NEXT: .reg .b16 %rs<2>; 410; CHECK-EMPTY: 411; CHECK-NEXT: // %bb.0: 412; CHECK-NEXT: mov.b16 %rs1, 0; 413; CHECK-NEXT: st.param.v2.b8 [func_retval0], {%rs1, %rs1}; 414; CHECK-NEXT: st.param.v2.b8 [func_retval0+2], {%rs1, %rs1}; 415; CHECK-NEXT: st.param.v2.b8 [func_retval0+4], {%rs1, %rs1}; 416; CHECK-NEXT: st.param.v2.b8 [func_retval0+6], {%rs1, %rs1}; 417; CHECK-NEXT: st.param.b8 [func_retval0+8], %rs1; 418; CHECK-NEXT: ret; 419 ret <9 x i1> zeroinitializer 420} 421 422define <8 x i1> @bit8() { 423; CHECK-LABEL: bit8( 424; CHECK: { 425; CHECK-NEXT: .reg .b16 %rs<2>; 426; CHECK-EMPTY: 427; CHECK-NEXT: // %bb.0: 428; CHECK-NEXT: mov.b16 %rs1, 0; 429; CHECK-NEXT: st.param.b8 [func_retval0], %rs1; 430; CHECK-NEXT: st.param.b8 [func_retval0+1], %rs1; 431; CHECK-NEXT: st.param.b8 [func_retval0+2], %rs1; 432; CHECK-NEXT: st.param.b8 [func_retval0+3], %rs1; 433; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs1; 434; CHECK-NEXT: st.param.b8 [func_retval0+5], %rs1; 435; CHECK-NEXT: st.param.b8 [func_retval0+6], %rs1; 436; CHECK-NEXT: st.param.b8 [func_retval0+7], %rs1; 437; CHECK-NEXT: ret; 438 ret <8 x i1> zeroinitializer 439} 440 441define <7 x i1> @bit7() { 442; CHECK-LABEL: bit7( 443; CHECK: { 444; CHECK-NEXT: .reg .b16 %rs<2>; 445; CHECK-EMPTY: 446; CHECK-NEXT: // %bb.0: 447; CHECK-NEXT: mov.b16 %rs1, 0; 448; CHECK-NEXT: st.param.b8 [func_retval0], %rs1; 449; CHECK-NEXT: st.param.b8 [func_retval0+1], %rs1; 450; CHECK-NEXT: st.param.b8 [func_retval0+2], %rs1; 451; CHECK-NEXT: st.param.b8 [func_retval0+3], %rs1; 452; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs1; 453; CHECK-NEXT: st.param.b8 [func_retval0+5], %rs1; 454; CHECK-NEXT: st.param.b8 [func_retval0+6], %rs1; 455; CHECK-NEXT: ret; 456 ret <7 x i1> zeroinitializer 457} 458 459define <5 x i1> @bit5() { 460; CHECK-LABEL: bit5( 461; CHECK: { 462; CHECK-NEXT: .reg .b16 %rs<2>; 463; CHECK-EMPTY: 464; CHECK-NEXT: // %bb.0: 465; CHECK-NEXT: mov.b16 %rs1, 0; 466; CHECK-NEXT: st.param.b8 [func_retval0], %rs1; 467; CHECK-NEXT: st.param.b8 [func_retval0+1], %rs1; 468; CHECK-NEXT: st.param.b8 [func_retval0+2], %rs1; 469; CHECK-NEXT: st.param.b8 [func_retval0+3], %rs1; 470; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs1; 471; CHECK-NEXT: ret; 472 ret <5 x i1> zeroinitializer 473} 474 475define <4 x i1> @bit4() { 476; CHECK-LABEL: bit4( 477; CHECK: { 478; CHECK-NEXT: .reg .b16 %rs<2>; 479; CHECK-EMPTY: 480; CHECK-NEXT: // %bb.0: 481; CHECK-NEXT: mov.b16 %rs1, 0; 482; CHECK-NEXT: st.param.b8 [func_retval0], %rs1; 483; CHECK-NEXT: st.param.b8 [func_retval0+1], %rs1; 484; CHECK-NEXT: st.param.b8 [func_retval0+2], %rs1; 485; CHECK-NEXT: st.param.b8 [func_retval0+3], %rs1; 486; CHECK-NEXT: ret; 487 ret <4 x i1> zeroinitializer 488} 489 490define <3 x i1> @bit3() { 491; CHECK-LABEL: bit3( 492; CHECK: { 493; CHECK-NEXT: .reg .b16 %rs<2>; 494; CHECK-EMPTY: 495; CHECK-NEXT: // %bb.0: 496; CHECK-NEXT: mov.b16 %rs1, 0; 497; CHECK-NEXT: st.param.b8 [func_retval0], %rs1; 498; CHECK-NEXT: st.param.b8 [func_retval0+1], %rs1; 499; CHECK-NEXT: st.param.b8 [func_retval0+2], %rs1; 500; CHECK-NEXT: ret; 501 ret <3 x i1> zeroinitializer 502} 503 504define <2 x i1> @bit2() { 505; CHECK-LABEL: bit2( 506; CHECK: { 507; CHECK-NEXT: .reg .b16 %rs<2>; 508; CHECK-EMPTY: 509; CHECK-NEXT: // %bb.0: 510; CHECK-NEXT: mov.b16 %rs1, 0; 511; CHECK-NEXT: st.param.b8 [func_retval0], %rs1; 512; CHECK-NEXT: st.param.b8 [func_retval0+1], %rs1; 513; CHECK-NEXT: ret; 514 ret <2 x i1> zeroinitializer 515} 516 517define <1 x i1> @bit1() { 518; CHECK-LABEL: bit1( 519; CHECK: { 520; CHECK-NEXT: .reg .b16 %rs<2>; 521; CHECK-EMPTY: 522; CHECK-NEXT: // %bb.0: 523; CHECK-NEXT: mov.b16 %rs1, 0; 524; CHECK-NEXT: st.param.b8 [func_retval0], %rs1; 525; CHECK-NEXT: ret; 526 ret <1 x i1> zeroinitializer 527} 528