1 files changed, 115 insertions, 0 deletions
diff --git a/llvm/test/CodeGen/NVPTX/i8x2-instructions.ll b/llvm/test/CodeGen/NVPTX/i8x2-instructions.ll
index f4053d84593a..db19495b5a4b 100644
--- a/llvm/test/CodeGen/NVPTX/i8x2-instructions.ll
+++ b/llvm/test/CodeGen/NVPTX/i8x2-instructions.ll
@@ -132,5 +132,120 @@ define <2 x float> @test_uitofp_2xi8(<2 x i8> %a) {
   %1 = uitofp <2 x i8> %a to <2 x float>
   ret <2 x float> %1
 }
+
+define void @test_store_i8x2_unaligned(ptr %ptr, <2 x i8> %a) {
+; O0-LABEL: test_store_i8x2_unaligned(
+; O0:       {
+; O0-NEXT:    .reg .b16 %rs<3>;
+; O0-NEXT:    .reg .b32 %r<2>;
+; O0-NEXT:    .reg .b64 %rd<2>;
+; O0-EMPTY:
+; O0-NEXT:  // %bb.0:
+; O0-NEXT:    ld.param.b64 %rd1, [test_store_i8x2_unaligned_param_0];
+; O0-NEXT:    ld.param.v2.b8 {%rs1, %rs2}, [test_store_i8x2_unaligned_param_1];
+; O0-NEXT:    mov.b32 %r1, {%rs1, %rs2};
+; O0-NEXT:    st.b8 [%rd1+1], %rs2;
+; O0-NEXT:    st.b8 [%rd1], %rs1;
+; O0-NEXT:    ret;
+;
+; O3-LABEL: test_store_i8x2_unaligned(
+; O3:       {
+; O3-NEXT:    .reg .b16 %rs<3>;
+; O3-NEXT:    .reg .b64 %rd<2>;
+; O3-EMPTY:
+; O3-NEXT:  // %bb.0:
+; O3-NEXT:    ld.param.b64 %rd1, [test_store_i8x2_unaligned_param_0];
+; O3-NEXT:    ld.param.v2.b8 {%rs1, %rs2}, [test_store_i8x2_unaligned_param_1];
+; O3-NEXT:    st.b8 [%rd1+1], %rs2;
+; O3-NEXT:    st.b8 [%rd1], %rs1;
+; O3-NEXT:    ret;
+  store <2 x i8> %a, ptr %ptr, align 1
+  ret void
+}
+
+define void @test_store_i8x2_unaligned_immediate(ptr %ptr) {
+; O0-LABEL: test_store_i8x2_unaligned_immediate(
+; O0:       {
+; O0-NEXT:    .reg .b64 %rd<2>;
+; O0-EMPTY:
+; O0-NEXT:  // %bb.0:
+; O0-NEXT:    ld.param.b64 %rd1, [test_store_i8x2_unaligned_immediate_param_0];
+; O0-NEXT:    st.b8 [%rd1+1], 2;
+; O0-NEXT:    st.b8 [%rd1], 1;
+; O0-NEXT:    ret;
+;
+; O3-LABEL: test_store_i8x2_unaligned_immediate(
+; O3:       {
+; O3-NEXT:    .reg .b64 %rd<2>;
+; O3-EMPTY:
+; O3-NEXT:  // %bb.0:
+; O3-NEXT:    ld.param.b64 %rd1, [test_store_i8x2_unaligned_immediate_param_0];
+; O3-NEXT:    st.b8 [%rd1+1], 2;
+; O3-NEXT:    st.b8 [%rd1], 1;
+; O3-NEXT:    ret;
+  store <2 x i8> <i8 1, i8 2>, ptr %ptr, align 1
+  ret void
+}
+
+define i32 @test_zext_load_i8x2_unaligned(ptr %ptr) {
+; O0-LABEL: test_zext_load_i8x2_unaligned(
+; O0:       {
+; O0-NEXT:    .reg .b16 %rs<3>;
+; O0-NEXT:    .reg .b64 %rd<2>;
+; O0-EMPTY:
+; O0-NEXT:  // %bb.0:
+; O0-NEXT:    ld.param.b64 %rd1, [test_zext_load_i8x2_unaligned_param_0];
+; O0-NEXT:    ld.b8 %rs1, [%rd1+1];
+; O0-NEXT:    ld.b8 %rs2, [%rd1];
+; O0-NEXT:    st.param.v2.b16 [func_retval0], {%rs2, %rs1};
+; O0-NEXT:    ret;
+;
+; O3-LABEL: test_zext_load_i8x2_unaligned(
+; O3:       {
+; O3-NEXT:    .reg .b16 %rs<3>;
+; O3-NEXT:    .reg .b64 %rd<2>;
+; O3-EMPTY:
+; O3-NEXT:  // %bb.0:
+; O3-NEXT:    ld.param.b64 %rd1, [test_zext_load_i8x2_unaligned_param_0];
+; O3-NEXT:    ld.b8 %rs1, [%rd1+1];
+; O3-NEXT:    ld.b8 %rs2, [%rd1];
+; O3-NEXT:    st.param.v2.b16 [func_retval0], {%rs2, %rs1};
+; O3-NEXT:    ret;
+  %a = load <2 x i8>, ptr %ptr, align 1
+  %b = zext <2 x i8> %a to <2 x i16>
+  %c = bitcast <2 x i16> %b to i32
+  ret i32 %c
+}
+
+define i32 @test_sext_load_i8x2_unaligned(ptr %ptr) {
+; O0-LABEL: test_sext_load_i8x2_unaligned(
+; O0:       {
+; O0-NEXT:    .reg .b16 %rs<3>;
+; O0-NEXT:    .reg .b64 %rd<2>;
+; O0-EMPTY:
+; O0-NEXT:  // %bb.0:
+; O0-NEXT:    ld.param.b64 %rd1, [test_sext_load_i8x2_unaligned_param_0];
+; O0-NEXT:    ld.s8 %rs1, [%rd1+1];
+; O0-NEXT:    ld.s8 %rs2, [%rd1];
+; O0-NEXT:    st.param.v2.b16 [func_retval0], {%rs2, %rs1};
+; O0-NEXT:    ret;
+;
+; O3-LABEL: test_sext_load_i8x2_unaligned(
+; O3:       {
+; O3-NEXT:    .reg .b16 %rs<3>;
+; O3-NEXT:    .reg .b64 %rd<2>;
+; O3-EMPTY:
+; O3-NEXT:  // %bb.0:
+; O3-NEXT:    ld.param.b64 %rd1, [test_sext_load_i8x2_unaligned_param_0];
+; O3-NEXT:    ld.s8 %rs1, [%rd1+1];
+; O3-NEXT:    ld.s8 %rs2, [%rd1];
+; O3-NEXT:    st.param.v2.b16 [func_retval0], {%rs2, %rs1};
+; O3-NEXT:    ret;
+  %a = load <2 x i8>, ptr %ptr, align 1
+  %b = sext <2 x i8> %a to <2 x i16>
+  %c = bitcast <2 x i16> %b to i32
+  ret i32 %c
+}
+
 ;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
 ; COMMON: {{.*}}