Skip to content

Commit 0ad2f80

Browse files
dakersnardakersnar
committed
Change pragma printing to use hex
1 parent d71cfc9 commit 0ad2f80

File tree

2 files changed

+18
-18
lines changed

2 files changed

+18
-18
lines changed

‎llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -401,7 +401,7 @@ void NVPTXInstPrinter::printUsedBytesMaskPragma(const MCInst *MI, int OpNum,
401401
assert(Op.isImm() && "Invalid operand");
402402
uint32_t Imm = (uint32_t)Op.getImm();
403403
if (Imm != UINT32_MAX) {
404-
O << ".pragma \"used_bytes_mask " << Imm << "\";\n\t";
404+
O << ".pragma \"used_bytes_mask " << format_hex(Imm, 1) << "\";\n\t";
405405
}
406406
}
407407

‎llvm/test/CodeGen/NVPTX/masked-load-vectors.ll‎

Lines changed: 17 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -18,9 +18,9 @@ define void @global_8xi32(ptr addrspace(1) %a, ptr addrspace(1) %b) {
1818
; SM90-EMPTY:
1919
; SM90-NEXT: // %bb.0:
2020
; SM90-NEXT: ld.param.b64 %rd1, [global_8xi32_param_0];
21-
; SM90-NEXT: .pragma "used_bytes_mask 61440";
21+
; SM90-NEXT: .pragma "used_bytes_mask 0xf000";
2222
; SM90-NEXT: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];
23-
; SM90-NEXT: .pragma "used_bytes_mask 3855";
23+
; SM90-NEXT: .pragma "used_bytes_mask 0xf0f";
2424
; SM90-NEXT: ld.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];
2525
; SM90-NEXT: ld.param.b64 %rd2, [global_8xi32_param_1];
2626
; SM90-NEXT: st.global.b32 [%rd2], %r5;
@@ -35,7 +35,7 @@ define void @global_8xi32(ptr addrspace(1) %a, ptr addrspace(1) %b) {
3535
; SM100-EMPTY:
3636
; SM100-NEXT: // %bb.0:
3737
; SM100-NEXT: ld.param.b64 %rd1, [global_8xi32_param_0];
38-
; SM100-NEXT: .pragma "used_bytes_mask 4026535695";
38+
; SM100-NEXT: .pragma "used_bytes_mask 0xf0000f0f";
3939
; SM100-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];
4040
; SM100-NEXT: ld.param.b64 %rd2, [global_8xi32_param_1];
4141
; SM100-NEXT: st.global.v8.b32 [%rd2], {%r1, _, %r3, _, _, _, _, %r8};
@@ -56,10 +56,10 @@ define void @global_16xi16(ptr addrspace(1) %a, ptr addrspace(1) %b) {
5656
; SM90-EMPTY:
5757
; SM90-NEXT: // %bb.0:
5858
; SM90-NEXT: ld.param.b64 %rd1, [global_16xi16_param_0];
59-
; SM90-NEXT: .pragma "used_bytes_mask 61440";
59+
; SM90-NEXT: .pragma "used_bytes_mask 0xf000";
6060
; SM90-NEXT: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];
6161
; SM90-NEXT: mov.b32 {%rs1, %rs2}, %r4;
62-
; SM90-NEXT: .pragma "used_bytes_mask 3855";
62+
; SM90-NEXT: .pragma "used_bytes_mask 0xf0f";
6363
; SM90-NEXT: ld.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];
6464
; SM90-NEXT: mov.b32 {%rs3, %rs4}, %r7;
6565
; SM90-NEXT: mov.b32 {%rs5, %rs6}, %r5;
@@ -80,7 +80,7 @@ define void @global_16xi16(ptr addrspace(1) %a, ptr addrspace(1) %b) {
8080
; SM100-EMPTY:
8181
; SM100-NEXT: // %bb.0:
8282
; SM100-NEXT: ld.param.b64 %rd1, [global_16xi16_param_0];
83-
; SM100-NEXT: .pragma "used_bytes_mask 4026535695";
83+
; SM100-NEXT: .pragma "used_bytes_mask 0xf0000f0f";
8484
; SM100-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];
8585
; SM100-NEXT: mov.b32 {%rs1, %rs2}, %r8;
8686
; SM100-NEXT: mov.b32 {%rs3, %rs4}, %r3;
@@ -128,9 +128,9 @@ define void @global_8xi32_invariant(ptr addrspace(1) %a, ptr addrspace(1) %b) {
128128
; SM90-EMPTY:
129129
; SM90-NEXT: // %bb.0:
130130
; SM90-NEXT: ld.param.b64 %rd1, [global_8xi32_invariant_param_0];
131-
; SM90-NEXT: .pragma "used_bytes_mask 61440";
131+
; SM90-NEXT: .pragma "used_bytes_mask 0xf000";
132132
; SM90-NEXT: ld.global.nc.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];
133-
; SM90-NEXT: .pragma "used_bytes_mask 3855";
133+
; SM90-NEXT: .pragma "used_bytes_mask 0xf0f";
134134
; SM90-NEXT: ld.global.nc.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];
135135
; SM90-NEXT: ld.param.b64 %rd2, [global_8xi32_invariant_param_1];
136136
; SM90-NEXT: st.global.b32 [%rd2], %r5;
@@ -145,7 +145,7 @@ define void @global_8xi32_invariant(ptr addrspace(1) %a, ptr addrspace(1) %b) {
145145
; SM100-EMPTY:
146146
; SM100-NEXT: // %bb.0:
147147
; SM100-NEXT: ld.param.b64 %rd1, [global_8xi32_invariant_param_0];
148-
; SM100-NEXT: .pragma "used_bytes_mask 4026535695";
148+
; SM100-NEXT: .pragma "used_bytes_mask 0xf0000f0f";
149149
; SM100-NEXT: ld.global.nc.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];
150150
; SM100-NEXT: ld.param.b64 %rd2, [global_8xi32_invariant_param_1];
151151
; SM100-NEXT: st.global.v8.b32 [%rd2], {%r1, _, %r3, _, _, _, _, %r8};
@@ -164,7 +164,7 @@ define void @global_2xi16(ptr addrspace(1) %a, ptr addrspace(1) %b) {
164164
; CHECK-EMPTY:
165165
; CHECK-NEXT: // %bb.0:
166166
; CHECK-NEXT: ld.param.b64 %rd1, [global_2xi16_param_0];
167-
; CHECK-NEXT: .pragma "used_bytes_mask 3";
167+
; CHECK-NEXT: .pragma "used_bytes_mask 0x3";
168168
; CHECK-NEXT: ld.global.b32 %r1, [%rd1];
169169
; CHECK-NEXT: ld.param.b64 %rd2, [global_2xi16_param_1];
170170
; CHECK-NEXT: mov.b32 {%rs1, _}, %r1;
@@ -184,7 +184,7 @@ define void @global_2xi16_invariant(ptr addrspace(1) %a, ptr addrspace(1) %b) {
184184
; CHECK-EMPTY:
185185
; CHECK-NEXT: // %bb.0:
186186
; CHECK-NEXT: ld.param.b64 %rd1, [global_2xi16_invariant_param_0];
187-
; CHECK-NEXT: .pragma "used_bytes_mask 3";
187+
; CHECK-NEXT: .pragma "used_bytes_mask 0x3";
188188
; CHECK-NEXT: ld.global.nc.b32 %r1, [%rd1];
189189
; CHECK-NEXT: ld.param.b64 %rd2, [global_2xi16_invariant_param_1];
190190
; CHECK-NEXT: mov.b32 {%rs1, _}, %r1;
@@ -220,7 +220,7 @@ define void @global_4xi8(ptr addrspace(1) %a, ptr addrspace(1) %b) {
220220
; CHECK-EMPTY:
221221
; CHECK-NEXT: // %bb.0:
222222
; CHECK-NEXT: ld.param.b64 %rd1, [global_4xi8_param_0];
223-
; CHECK-NEXT: .pragma "used_bytes_mask 5";
223+
; CHECK-NEXT: .pragma "used_bytes_mask 0x5";
224224
; CHECK-NEXT: ld.global.b32 %r1, [%rd1];
225225
; CHECK-NEXT: ld.param.b64 %rd2, [global_4xi8_param_1];
226226
; CHECK-NEXT: st.global.b8 [%rd2], %r1;
@@ -240,7 +240,7 @@ define void @global_4xi8_invariant(ptr addrspace(1) %a, ptr addrspace(1) %b) {
240240
; CHECK-EMPTY:
241241
; CHECK-NEXT: // %bb.0:
242242
; CHECK-NEXT: ld.param.b64 %rd1, [global_4xi8_invariant_param_0];
243-
; CHECK-NEXT: .pragma "used_bytes_mask 5";
243+
; CHECK-NEXT: .pragma "used_bytes_mask 0x5";
244244
; CHECK-NEXT: ld.global.nc.b32 %r1, [%rd1];
245245
; CHECK-NEXT: ld.param.b64 %rd2, [global_4xi8_invariant_param_1];
246246
; CHECK-NEXT: st.global.b8 [%rd2], %r1;
@@ -280,7 +280,7 @@ define void @global_2xf32(ptr addrspace(1) %a, ptr addrspace(1) %b) {
280280
; SM90-EMPTY:
281281
; SM90-NEXT: // %bb.0:
282282
; SM90-NEXT: ld.param.b64 %rd1, [global_2xf32_param_0];
283-
; SM90-NEXT: .pragma "used_bytes_mask 15";
283+
; SM90-NEXT: .pragma "used_bytes_mask 0xf";
284284
; SM90-NEXT: ld.global.v2.b32 {%r1, %r2}, [%rd1];
285285
; SM90-NEXT: ld.param.b64 %rd2, [global_2xf32_param_1];
286286
; SM90-NEXT: st.global.b32 [%rd2], %r1;
@@ -293,7 +293,7 @@ define void @global_2xf32(ptr addrspace(1) %a, ptr addrspace(1) %b) {
293293
; SM100-EMPTY:
294294
; SM100-NEXT: // %bb.0:
295295
; SM100-NEXT: ld.param.b64 %rd1, [global_2xf32_param_0];
296-
; SM100-NEXT: .pragma "used_bytes_mask 15";
296+
; SM100-NEXT: .pragma "used_bytes_mask 0xf";
297297
; SM100-NEXT: ld.global.b64 %rd2, [%rd1];
298298
; SM100-NEXT: ld.param.b64 %rd3, [global_2xf32_param_1];
299299
; SM100-NEXT: mov.b64 {%r1, _}, %rd2;
@@ -312,7 +312,7 @@ define void @global_2xf32_invariant(ptr addrspace(1) %a, ptr addrspace(1) %b) {
312312
; SM90-EMPTY:
313313
; SM90-NEXT: // %bb.0:
314314
; SM90-NEXT: ld.param.b64 %rd1, [global_2xf32_invariant_param_0];
315-
; SM90-NEXT: .pragma "used_bytes_mask 15";
315+
; SM90-NEXT: .pragma "used_bytes_mask 0xf";
316316
; SM90-NEXT: ld.global.nc.v2.b32 {%r1, %r2}, [%rd1];
317317
; SM90-NEXT: ld.param.b64 %rd2, [global_2xf32_invariant_param_1];
318318
; SM90-NEXT: st.global.b32 [%rd2], %r1;
@@ -325,7 +325,7 @@ define void @global_2xf32_invariant(ptr addrspace(1) %a, ptr addrspace(1) %b) {
325325
; SM100-EMPTY:
326326
; SM100-NEXT: // %bb.0:
327327
; SM100-NEXT: ld.param.b64 %rd1, [global_2xf32_invariant_param_0];
328-
; SM100-NEXT: .pragma "used_bytes_mask 15";
328+
; SM100-NEXT: .pragma "used_bytes_mask 0xf";
329329
; SM100-NEXT: ld.global.nc.b64 %rd2, [%rd1];
330330
; SM100-NEXT: ld.param.b64 %rd3, [global_2xf32_invariant_param_1];
331331
; SM100-NEXT: mov.b64 {%r1, _}, %rd2;

0 commit comments

Comments
 (0)