address comments

Author: AlexMaclean

llvm/test/CodeGen/NVPTX/bf16-instructions.ll

@@ -1605,5 +1605,23 @@ define <2 x bfloat> @test_maxnum_v2(<2 x bfloat> %a, <2 x bfloat> %b) {
   ret <2 x bfloat> %r
 }
 
+define void @store_bf16(ptr %p1, ptr %p2, bfloat %v) {
+; CHECK-LABEL: store_bf16(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b16 %rs<2>;
+; CHECK-NEXT:    .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b64 %rd1, [store_bf16_param_0];
+; CHECK-NEXT:    ld.param.b16 %rs1, [store_bf16_param_2];
+; CHECK-NEXT:    st.b16 [%rd1], %rs1;
+; CHECK-NEXT:    ld.param.b64 %rd2, [store_bf16_param_1];
+; CHECK-NEXT:    st.b16 [%rd2], 0x3F80;
+; CHECK-NEXT:    ret;
+  store bfloat %v, ptr %p1
+  store bfloat 1.0, ptr %p2
+  ret void
+}
+
 declare bfloat @llvm.maximum.bf16(bfloat, bfloat)
 declare <2 x bfloat> @llvm.maximum.v2bf16(<2 x bfloat>, <2 x bfloat>)

llvm/test/CodeGen/NVPTX/bf16x2-instructions.ll

@@ -723,3 +723,20 @@ define <2 x bfloat> @test_copysign(<2 x bfloat> %a, <2 x bfloat> %b) #0 {
   ret <2 x bfloat> %r
 }
 
+define void @test_store_bf16x2(ptr %p1, ptr %p2, <2 x bfloat> %v) {
+; CHECK-LABEL: test_store_bf16x2(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b32 %r<2>;
+; CHECK-NEXT:    .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b64 %rd1, [test_store_bf16x2_param_0];
+; CHECK-NEXT:    ld.param.b32 %r1, [test_store_bf16x2_param_2];
+; CHECK-NEXT:    st.b32 [%rd1], %r1;
+; CHECK-NEXT:    ld.param.b64 %rd2, [test_store_bf16x2_param_1];
+; CHECK-NEXT:    st.b32 [%rd2], 1065369472;
+; CHECK-NEXT:    ret;
+  store <2 x bfloat> %v, ptr %p1
+  store <2 x bfloat> <bfloat 1.0, bfloat 1.0>, ptr %p2
+  ret void
+}

llvm/test/CodeGen/NVPTX/f16x2-instructions.ll

@@ -2295,5 +2295,25 @@ define <2 x half> @test_uitofp_2xi16_to_2xhalf(<2 x i16> %a) #0 {
   ret <2 x half> %r
 }
 
+define void @test_store_2xhalf(ptr %p1, ptr %p2, <2 x half> %v) {
+; CHECK-LABEL: test_store_2xhalf(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b32 %r<2>;
+; CHECK-NEXT:    .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b32 %r1, [test_store_2xhalf_param_2];
+; CHECK-NEXT:    ld.param.b64 %rd2, [test_store_2xhalf_param_1];
+; CHECK-NEXT:    ld.param.b64 %rd1, [test_store_2xhalf_param_0];
+; CHECK-NEXT:    st.b32 [%rd1], %r1;
+; CHECK-NEXT:    st.b32 [%rd2], 1006648320;
+; CHECK-NEXT:    ret;
+  store <2 x half> %v, ptr %p1
+  store <2 x half> <half 1.0, half 1.0>, ptr %p2
+  ret void
+}
+
+
+
 attributes #0 = { nounwind }
 attributes #1 = { "unsafe-fp-math" = "true" }