[llvm] [NVPTX] Support BFloat Store Parameter (PR #137074)

[Justin Fargnoli via llvm-commits]

================
@@ -2000,3 +2000,51 @@ declare void @call_v4_i8(%struct.char4 alignstack(4))
 declare void @call_v4_i16(%struct.short4 alignstack(8))
 declare void @call_v4_i32(%struct.int4 alignstack(16))
 declare void @call_v4_f32(%struct.float4 alignstack(16))
+
+define void @st_param_bfloat() {
+; CHECK-LABEL: st_param_bfloat(
+; CHECK: {
+; CHECK-NEXT:	.reg .b16 	%rs<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:// %bb.0:
+; CHECK-NEXT:	mov.b16 	%rs1, 0x4100;
+; CHECK-NEXT:	{ // callseq 83, 0
+; CHECK-NEXT:	.param .align 2 .b8 param0[2];
+; CHECK-NEXT:	st.param.b16 	[param0], %rs1;
+; CHECK-NEXT:	call.uni
+; CHECK-NEXT:	call_bfloat,
+; CHECK-NEXT:	(
+; CHECK-NEXT:	param0
+; CHECK-NEXT:	);
+; CHECK-NEXT:	} // callseq 83
+; CHECK-NEXT:	ret;
+  %five = bitcast i16 16640 to bfloat
+  call void @call_bfloat(bfloat %five)
+  ret void
+}
+
+define void @st_param_v2bfloat(<2 x bfloat> %val) {
+; CHECK-LABEL: st_param_v2bfloat(
+; CHECK:	.param .align 4 .b8 st_param_v2bfloat_param_0[4]
+; CHECK-NEXT: )
+; CHECK-NEXT: {
+; CHECK-NEXT:		.reg .b32 	%r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:	// %bb.0:
+; CHECK-NEXT:	ld.param.b32 	%r1, [st_param_v2bfloat_param_0];
+; CHECK-NEXT:	{ // callseq 84, 0
+; CHECK-NEXT:	.param .align 4 .b8 param0[4];
+; CHECK-NEXT:	st.param.b32 	[param0], %r1;
+; CHECK-NEXT:	call.uni
+; CHECK-NEXT:	call_v2bfloat,
+; CHECK-NEXT:	(
+; CHECK-NEXT:	param0
+; CHECK-NEXT:	);
+; CHECK-NEXT:	} // callseq 84
+; CHECK-NEXT:	ret;
+  call void @call_v2bfloat(<2 x bfloat> %val)
+  ret void
+}
----------------
justinfargnoli wrote:

It doesn't look like the output of this test has changed ([source](https://godbolt.org/z/raM6dKjvh))

https://github.com/llvm/llvm-project/pull/137074