Skip to content

Commit 101c2d0

Browse files
AlexMacleanAlexMaclean
committed
more tests
1 parent 76be42d commit 101c2d0

File tree

7 files changed

+57
-58
lines changed

7 files changed

+57
-58
lines changed

‎llvm/test/CodeGen/MIR/NVPTX/expected-floating-point-literal.mir

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -12,8 +12,8 @@
1212
---
1313
name: test
1414
registers:
15-
- { id: 0, class: float32regs }
16-
- { id: 1, class: float32regs }
15+
- { id: 0, class: int32regs }
16+
- { id: 1, class: int32regs }
1717
body: |
1818
bb.0.entry:
1919
%0 = LD_f32 0, 4, 1, 2, 32, &test_param_0, 0

‎llvm/test/CodeGen/MIR/NVPTX/floating-point-immediate-operands.mir

Lines changed: 18 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -30,24 +30,24 @@
3030
---
3131
name: test
3232
registers:
33-
- { id: 0, class: float32regs }
34-
- { id: 1, class: float64regs }
33+
- { id: 0, class: int32regs }
34+
- { id: 1, class: int64regs }
3535
- { id: 2, class: int32regs }
36-
- { id: 3, class: float64regs }
37-
- { id: 4, class: float32regs }
38-
- { id: 5, class: float32regs }
39-
- { id: 6, class: float32regs }
40-
- { id: 7, class: float32regs }
36+
- { id: 3, class: int64regs }
37+
- { id: 4, class: int32regs }
38+
- { id: 5, class: int32regs }
39+
- { id: 6, class: int32regs }
40+
- { id: 7, class: int32regs }
4141
body: |
4242
bb.0.entry:
4343
%0 = LD_f32 0, 0, 4, 2, 32, &test_param_0, 0
4444
%1 = CVT_f64_f32 %0, 0
4545
%2 = LD_i32 0, 0, 4, 0, 32, &test_param_1, 0
46-
; CHECK: %3:float64regs = FADD_rnf64ri %1, double 3.250000e+00
46+
; CHECK: %3:int64regs = FADD_rnf64ri %1, double 3.250000e+00
4747
%3 = FADD_rnf64ri %1, double 3.250000e+00
4848
%4 = CVT_f32_f64 %3, 5
4949
%5 = CVT_f32_s32 %2, 5
50-
; CHECK: %6:float32regs = FADD_rnf32ri %5, float 6.250000e+00
50+
; CHECK: %6:int32regs = FADD_rnf32ri %5, float 6.250000e+00
5151
%6 = FADD_rnf32ri %5, float 6.250000e+00
5252
%7 = FMUL_rnf32rr %6, %4
5353
StoreRetvalF32 %7, 0
@@ -56,24 +56,24 @@ body: |
5656
---
5757
name: test2
5858
registers:
59-
- { id: 0, class: float32regs }
60-
- { id: 1, class: float64regs }
59+
- { id: 0, class: int32regs }
60+
- { id: 1, class: int64regs }
6161
- { id: 2, class: int32regs }
62-
- { id: 3, class: float64regs }
63-
- { id: 4, class: float32regs }
64-
- { id: 5, class: float32regs }
65-
- { id: 6, class: float32regs }
66-
- { id: 7, class: float32regs }
62+
- { id: 3, class: int64regs }
63+
- { id: 4, class: int32regs }
64+
- { id: 5, class: int32regs }
65+
- { id: 6, class: int32regs }
66+
- { id: 7, class: int32regs }
6767
body: |
6868
bb.0.entry:
6969
%0 = LD_f32 0, 0, 4, 2, 32, &test2_param_0, 0
7070
%1 = CVT_f64_f32 %0, 0
7171
%2 = LD_i32 0, 0, 4, 0, 32, &test2_param_1, 0
72-
; CHECK: %3:float64regs = FADD_rnf64ri %1, double 0x7FF8000000000000
72+
; CHECK: %3:int64regs = FADD_rnf64ri %1, double 0x7FF8000000000000
7373
%3 = FADD_rnf64ri %1, double 0x7FF8000000000000
7474
%4 = CVT_f32_f64 %3, 5
7575
%5 = CVT_f32_s32 %2, 5
76-
; CHECK: %6:float32regs = FADD_rnf32ri %5, float 0x7FF8000000000000
76+
; CHECK: %6:int32regs = FADD_rnf32ri %5, float 0x7FF8000000000000
7777
%6 = FADD_rnf32ri %5, float 0x7FF8000000000000
7878
%7 = FMUL_rnf32rr %6, %4
7979
StoreRetvalF32 %7, 0

‎llvm/test/CodeGen/MIR/NVPTX/floating-point-invalid-type-error.mir

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -12,8 +12,8 @@
1212
---
1313
name: test
1414
registers:
15-
- { id: 0, class: float32regs }
16-
- { id: 1, class: float32regs }
15+
- { id: 0, class: int32regs }
16+
- { id: 1, class: int32regs }
1717
body: |
1818
bb.0.entry:
1919
%0 = LD_f32 0, 4, 1, 2, 32, &test_param_0, 0

‎llvm/test/DebugInfo/NVPTX/debug-addr-class.ll

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -317,7 +317,7 @@ declare void @llvm.dbg.declare(metadata, metadata, metadata)
317317
; CHECK-NEXT: .b8 5 // DW_AT_location
318318
; CHECK-NEXT: .b8 144
319319
; CHECK-NEXT: .b8 177
320-
; CHECK-NEXT: .b8 204
320+
; CHECK-NEXT: .b8 228
321321
; CHECK-NEXT: .b8 149
322322
; CHECK-NEXT: .b8 1
323323
; CHECK-NEXT: .b8 97 // DW_AT_name
@@ -357,7 +357,7 @@ declare void @llvm.dbg.declare(metadata, metadata, metadata)
357357
; CHECK-NEXT: .b8 2 // DW_AT_address_class
358358
; CHECK-NEXT: .b8 5 // DW_AT_location
359359
; CHECK-NEXT: .b8 144
360-
; CHECK-NEXT: .b8 177
360+
; CHECK-NEXT: .b8 178
361361
; CHECK-NEXT: .b8 228
362362
; CHECK-NEXT: .b8 149
363363
; CHECK-NEXT: .b8 1

‎llvm/test/DebugInfo/NVPTX/debug-info.ll

Lines changed: 11 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -20,8 +20,7 @@
2020
; CHECK: )
2121
; CHECK: {
2222
; CHECK-DAG: .reg .pred %p<2>;
23-
; CHECK-DAG: .reg .b32 %f<5>;
24-
; CHECK-DAG: .reg .b32 %r<6>;
23+
; CHECK-DAG: .reg .b32 %r<10>;
2524
; CHECK-DAG: .reg .b64 %rd<8>;
2625
; CHECK: .loc [[DEBUG_INFO_CU:[0-9]+]] 5 0
2726
; CHECK: ld.param.b32 %r{{.+}}, [{{.+}}];
@@ -41,18 +40,18 @@
4140
; CHECK: setp.ge.s32 %p{{.+}}, %r{{.+}}, %r{{.+}};
4241
; CHECK: .loc [[DEBUG_INFO_CU]] 7 7
4342
; CHECK: @%p{{.+}} bra [[BB:\$L__.+]];
44-
; CHECK: ld.param.b32 %f{{.+}}, [{{.+}}];
43+
; CHECK: ld.param.b32 %r{{.+}}, [{{.+}}];
4544
; CHECK: .loc [[DEBUG_INFO_CU]] 8 13
4645
; CHECK: mul.wide.u32 %rd{{.+}}, %r{{.+}}, 4;
4746
; CHECK: add.s64 %rd{{.+}}, %rd{{.+}}, %rd{{.+}};
48-
; CHECK: ld.global.b32 %f{{.+}}, [%rd{{.+}}];
47+
; CHECK: ld.global.b32 %r{{.+}}, [%rd{{.+}}];
4948
; CHECK: .loc [[DEBUG_INFO_CU]] 8 19
5049
; CHECK: add.s64 %rd{{.+}}, %rd{{.+}}, %rd{{.+}};
51-
; CHECK: ld.global.b32 %f{{.+}}, [%rd{{.+}}];
50+
; CHECK: ld.global.b32 %r{{.+}}, [%rd{{.+}}];
5251
; CHECK: .loc [[DEBUG_INFO_CU]] 3 82
53-
; CHECK: fma.rn.f32 %f{{.+}}, %f{{.+}}, %f{{.+}}, %f{{.+}};
52+
; CHECK: fma.rn.f32 %r{{.+}}, %r{{.+}}, %r{{.+}}, %r{{.+}};
5453
; CHECK: .loc [[DEBUG_INFO_CU]] 3 78
55-
; CHECK: st.global.b32 [%rd{{.+}}], %f{{.+}};
54+
; CHECK: st.global.b32 [%rd{{.+}}], %r{{.+}};
5655
; CHECK: [[BB]]:
5756
; CHECK: .loc [[DEBUG_INFO_CU]] 9 1
5857
; CHECK: ret;
@@ -105,8 +104,8 @@ if.end: ; preds = %if.then, %entry
105104
; CHECK-NEXT: .b8 5 // Loc expr size
106105
; CHECK-NEXT: .b8 0
107106
; CHECK-NEXT: .b8 144 // DW_OP_regx
108-
; CHECK-NEXT: .b8 177 // 2450993
109-
; CHECK-NEXT: .b8 204 //
107+
; CHECK-NEXT: .b8 178 // 2454066
108+
; CHECK-NEXT: .b8 228 //
110109
; CHECK-NEXT: .b8 149 //
111110
; CHECK-NEXT: .b8 1 //
112111
; CHECK-NEXT: .b64 0
@@ -2514,7 +2513,7 @@ if.end: ; preds = %if.then, %entry
25142513
; CHECK-NEXT: .b8 2 // DW_AT_address_class
25152514
; CHECK-NEXT: .b8 5 // DW_AT_location
25162515
; CHECK-NEXT: .b8 144
2517-
; CHECK-NEXT: .b8 178
2516+
; CHECK-NEXT: .b8 179
25182517
; CHECK-NEXT: .b8 228
25192518
; CHECK-NEXT: .b8 149
25202519
; CHECK-NEXT: .b8 1
@@ -2597,8 +2596,8 @@ if.end: ; preds = %if.then, %entry
25972596
; CHECK-NEXT: .b8 2 // DW_AT_address_class
25982597
; CHECK-NEXT: .b8 5 // DW_AT_location
25992598
; CHECK-NEXT: .b8 144
2600-
; CHECK-NEXT: .b8 179
2601-
; CHECK-NEXT: .b8 204
2599+
; CHECK-NEXT: .b8 184
2600+
; CHECK-NEXT: .b8 228
26022601
; CHECK-NEXT: .b8 149
26032602
; CHECK-NEXT: .b8 1
26042603
; CHECK-NEXT: .b32 2079 // DW_AT_abstract_origin

‎llvm/test/Transforms/SeparateConstOffsetFromGEP/NVPTX/split-gep-and-gvn.ll

Lines changed: 18 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -70,10 +70,10 @@ define void @sum_of_array(i32 %x, i32 %y, ptr nocapture %output) {
7070
ret void
7171
}
7272
; PTX-LABEL: sum_of_array(
73-
; PTX-DAG: ld.shared.b32 {{%f[0-9]+}}, [[[BASE_REG:%(rd|r)[0-9]+]]]
74-
; PTX-DAG: ld.shared.b32 {{%f[0-9]+}}, [[[BASE_REG]]+4]
75-
; PTX-DAG: ld.shared.b32 {{%f[0-9]+}}, [[[BASE_REG]]+128]
76-
; PTX-DAG: ld.shared.b32 {{%f[0-9]+}}, [[[BASE_REG]]+132]
73+
; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG:%(rd|r)[0-9]+]]]
74+
; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG]]+4]
75+
; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG]]+128]
76+
; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG]]+132]
7777

7878
; TODO: GVN is unable to preserve the "inbounds" keyword on the first GEP. Need
7979
; some infrastructure changes to enable such optimizations.
@@ -134,10 +134,10 @@ define void @sum_of_array2(i32 %x, i32 %y, ptr nocapture %output) {
134134
ret void
135135
}
136136
; PTX-LABEL: sum_of_array2(
137-
; PTX-DAG: ld.shared.b32 {{%f[0-9]+}}, [[[BASE_REG:%(rd|r)[0-9]+]]]
138-
; PTX-DAG: ld.shared.b32 {{%f[0-9]+}}, [[[BASE_REG]]+4]
139-
; PTX-DAG: ld.shared.b32 {{%f[0-9]+}}, [[[BASE_REG]]+128]
140-
; PTX-DAG: ld.shared.b32 {{%f[0-9]+}}, [[[BASE_REG]]+132]
137+
; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG:%(rd|r)[0-9]+]]]
138+
; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG]]+4]
139+
; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG]]+128]
140+
; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG]]+132]
141141

142142

143143

@@ -203,10 +203,10 @@ define void @sum_of_array3(i32 %x, i32 %y, ptr nocapture %output) {
203203
ret void
204204
}
205205
; PTX-LABEL: sum_of_array3(
206-
; PTX-DAG: ld.shared.b32 {{%f[0-9]+}}, [[[BASE_REG:%(rd|r)[0-9]+]]]
207-
; PTX-DAG: ld.shared.b32 {{%f[0-9]+}}, [[[BASE_REG]]+4]
208-
; PTX-DAG: ld.shared.b32 {{%f[0-9]+}}, [[[BASE_REG]]+128]
209-
; PTX-DAG: ld.shared.b32 {{%f[0-9]+}}, [[[BASE_REG]]+132]
206+
; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG:%(rd|r)[0-9]+]]]
207+
; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG]]+4]
208+
; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG]]+128]
209+
; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG]]+132]
210210

211211

212212

@@ -268,10 +268,10 @@ define void @sum_of_array4(i32 %x, i32 %y, ptr nocapture %output) {
268268
ret void
269269
}
270270
; PTX-LABEL: sum_of_array4(
271-
; PTX-DAG: ld.shared.b32 {{%f[0-9]+}}, [[[BASE_REG:%(rd|r)[0-9]+]]]
272-
; PTX-DAG: ld.shared.b32 {{%f[0-9]+}}, [[[BASE_REG]]+4]
273-
; PTX-DAG: ld.shared.b32 {{%f[0-9]+}}, [[[BASE_REG]]+128]
274-
; PTX-DAG: ld.shared.b32 {{%f[0-9]+}}, [[[BASE_REG]]+132]
271+
; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG:%(rd|r)[0-9]+]]]
272+
; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG]]+4]
273+
; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG]]+128]
274+
; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG]]+132]
275275

276276

277277

@@ -307,15 +307,15 @@ entry:
307307
%0 = sext i32 %xy to i64
308308
%p0 = getelementptr inbounds float, ptr %input, i64 %0
309309
%v0 = load float, ptr %p0, align 4
310-
; PTX: ld.b32 %f{{[0-9]+}}, [[[p0:%rd[0-9]+]]]
310+
; PTX: ld.b32 %r{{[0-9]+}}, [[[p0:%rd[0-9]+]]]
311311
call void @use(float %v0)
312312

313313
%y5 = add nsw i32 %y, 5
314314
%xy5 = add nsw i32 %x, %y5
315315
%1 = sext i32 %xy5 to i64
316316
%p1 = getelementptr inbounds float, ptr %input, i64 %1
317317
%v1 = load float, ptr %p1, align 4
318-
; PTX: ld.b32 %f{{[0-9]+}}, [[[p0]]+20]
318+
; PTX: ld.b32 %r{{[0-9]+}}, [[[p0]]+20]
319319
call void @use(float %v1)
320320

321321
ret void

‎llvm/test/Transforms/StraightLineStrengthReduce/NVPTX/reassociate-geps-and-slsr.ll

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -62,28 +62,28 @@ define void @slsr_after_reassociate_geps(ptr %arr, i32 %i) {
6262
; PTX: mul.wide.s32 [[i4:%rd[0-9]+]], [[i]], 4;
6363
; PTX: add.s64 [[base1:%rd[0-9]+]], [[arr]], [[i4]];
6464
%v1 = load float, ptr %p1, align 4
65-
; PTX: ld.b32 {{%f[0-9]+}}, [[[base1]]+20];
65+
; PTX: ld.b32 {{%r[0-9]+}}, [[[base1]]+20];
6666
call void @foo(float %v1)
6767

6868
%j2 = add nsw i32 %i2, 5
6969
%p2 = getelementptr inbounds float, ptr %arr, i32 %j2
7070
; PTX: add.s64 [[base2:%rd[0-9]+]], [[base1]], [[i4]];
7171
%v2 = load float, ptr %p2, align 4
72-
; PTX: ld.b32 {{%f[0-9]+}}, [[[base2]]+20];
72+
; PTX: ld.b32 {{%r[0-9]+}}, [[[base2]]+20];
7373
call void @foo(float %v2)
7474

7575
%j3 = add nsw i32 %i3, 5
7676
%p3 = getelementptr inbounds float, ptr %arr, i32 %j3
7777
; PTX: add.s64 [[base3:%rd[0-9]+]], [[base2]], [[i4]];
7878
%v3 = load float, ptr %p3, align 4
79-
; PTX: ld.b32 {{%f[0-9]+}}, [[[base3]]+20];
79+
; PTX: ld.b32 {{%r[0-9]+}}, [[[base3]]+20];
8080
call void @foo(float %v3)
8181

8282
%j4 = add nsw i32 %i4, 5
8383
%p4 = getelementptr inbounds float, ptr %arr, i32 %j4
8484
; PTX: add.s64 [[base4:%rd[0-9]+]], [[base3]], [[i4]];
8585
%v4 = load float, ptr %p4, align 4
86-
; PTX: ld.b32 {{%f[0-9]+}}, [[[base4]]+20];
86+
; PTX: ld.b32 {{%r[0-9]+}}, [[[base4]]+20];
8787
call void @foo(float %v4)
8888

8989
ret void

0 commit comments

Comments
 (0)