Skip to content

Commit da39f3a

Browse files
AlexMacleanAlexMaclean
committed
address comments - unify FileChecks
1 parent 8a2d3eb commit da39f3a

File tree

4 files changed

+187
-381
lines changed

4 files changed

+187
-381
lines changed

‎llvm/test/CodeGen/NVPTX/surf-read-cuda.ll‎

Lines changed: 31 additions & 62 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,6 @@
11
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2-
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefix=SM20
3-
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_30 -verify-machineinstrs | FileCheck %s --check-prefix=SM30
2+
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
3+
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_30 -verify-machineinstrs | FileCheck %s
44
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
55
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_30 -verify-machineinstrs | %ptxas-verify %}
66

@@ -11,37 +11,21 @@ declare i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1))
1111

1212

1313
define void @foo(i64 %img, ptr %red, i32 %idx) {
14-
; SM20-LABEL: foo(
15-
; SM20: {
16-
; SM20-NEXT: .reg .b32 %r<3>;
17-
; SM20-NEXT: .reg .f32 %f<2>;
18-
; SM20-NEXT: .reg .b64 %rd<4>;
19-
; SM20-EMPTY:
20-
; SM20-NEXT: // %bb.0:
21-
; SM20-NEXT: ld.param.u64 %rd1, [foo_param_0];
22-
; SM20-NEXT: ld.param.u64 %rd2, [foo_param_1];
23-
; SM20-NEXT: cvta.to.global.u64 %rd3, %rd2;
24-
; SM20-NEXT: ld.param.u32 %r1, [foo_param_2];
25-
; SM20-NEXT: suld.b.1d.b32.trap {%r2}, [%rd1, {%r1}];
26-
; SM20-NEXT: cvt.rn.f32.s32 %f1, %r2;
27-
; SM20-NEXT: st.global.f32 [%rd3], %f1;
28-
; SM20-NEXT: ret;
29-
;
30-
; SM30-LABEL: foo(
31-
; SM30: {
32-
; SM30-NEXT: .reg .b32 %r<3>;
33-
; SM30-NEXT: .reg .f32 %f<2>;
34-
; SM30-NEXT: .reg .b64 %rd<4>;
35-
; SM30-EMPTY:
36-
; SM30-NEXT: // %bb.0:
37-
; SM30-NEXT: ld.param.u64 %rd1, [foo_param_0];
38-
; SM30-NEXT: ld.param.u64 %rd2, [foo_param_1];
39-
; SM30-NEXT: cvta.to.global.u64 %rd3, %rd2;
40-
; SM30-NEXT: ld.param.u32 %r1, [foo_param_2];
41-
; SM30-NEXT: suld.b.1d.b32.trap {%r2}, [%rd1, {%r1}];
42-
; SM30-NEXT: cvt.rn.f32.s32 %f1, %r2;
43-
; SM30-NEXT: st.global.f32 [%rd3], %f1;
44-
; SM30-NEXT: ret;
14+
; CHECK-LABEL: foo(
15+
; CHECK: {
16+
; CHECK-NEXT: .reg .b32 %r<3>;
17+
; CHECK-NEXT: .reg .f32 %f<2>;
18+
; CHECK-NEXT: .reg .b64 %rd<4>;
19+
; CHECK-EMPTY:
20+
; CHECK-NEXT: // %bb.0:
21+
; CHECK-NEXT: ld.param.u64 %rd1, [foo_param_0];
22+
; CHECK-NEXT: ld.param.u64 %rd2, [foo_param_1];
23+
; CHECK-NEXT: cvta.to.global.u64 %rd3, %rd2;
24+
; CHECK-NEXT: ld.param.u32 %r1, [foo_param_2];
25+
; CHECK-NEXT: suld.b.1d.b32.trap {%r2}, [%rd1, {%r1}];
26+
; CHECK-NEXT: cvt.rn.f32.s32 %f1, %r2;
27+
; CHECK-NEXT: st.global.f32 [%rd3], %f1;
28+
; CHECK-NEXT: ret;
4529
%val = tail call i32 @llvm.nvvm.suld.1d.i32.trap(i64 %img, i32 %idx)
4630
%ret = sitofp i32 %val to float
4731
store float %ret, ptr %red
@@ -51,35 +35,20 @@ define void @foo(i64 %img, ptr %red, i32 %idx) {
5135
@surf0 = internal addrspace(1) global i64 0, align 8
5236

5337
define void @bar(ptr %red, i32 %idx) {
54-
; SM20-LABEL: bar(
55-
; SM20: {
56-
; SM20-NEXT: .reg .b32 %r<3>;
57-
; SM20-NEXT: .reg .f32 %f<2>;
58-
; SM20-NEXT: .reg .b64 %rd<4>;
59-
; SM20-EMPTY:
60-
; SM20-NEXT: // %bb.0:
61-
; SM20-NEXT: ld.param.u64 %rd1, [bar_param_0];
62-
; SM20-NEXT: cvta.to.global.u64 %rd2, %rd1;
63-
; SM20-NEXT: ld.param.u32 %r1, [bar_param_1];
64-
; SM20-NEXT: suld.b.1d.b32.trap {%r2}, [surf0, {%r1}];
65-
; SM20-NEXT: cvt.rn.f32.s32 %f1, %r2;
66-
; SM20-NEXT: st.global.f32 [%rd2], %f1;
67-
; SM20-NEXT: ret;
68-
;
69-
; SM30-LABEL: bar(
70-
; SM30: {
71-
; SM30-NEXT: .reg .b32 %r<3>;
72-
; SM30-NEXT: .reg .f32 %f<2>;
73-
; SM30-NEXT: .reg .b64 %rd<4>;
74-
; SM30-EMPTY:
75-
; SM30-NEXT: // %bb.0:
76-
; SM30-NEXT: ld.param.u64 %rd1, [bar_param_0];
77-
; SM30-NEXT: cvta.to.global.u64 %rd2, %rd1;
78-
; SM30-NEXT: ld.param.u32 %r1, [bar_param_1];
79-
; SM30-NEXT: suld.b.1d.b32.trap {%r2}, [surf0, {%r1}];
80-
; SM30-NEXT: cvt.rn.f32.s32 %f1, %r2;
81-
; SM30-NEXT: st.global.f32 [%rd2], %f1;
82-
; SM30-NEXT: ret;
38+
; CHECK-LABEL: bar(
39+
; CHECK: {
40+
; CHECK-NEXT: .reg .b32 %r<3>;
41+
; CHECK-NEXT: .reg .f32 %f<2>;
42+
; CHECK-NEXT: .reg .b64 %rd<4>;
43+
; CHECK-EMPTY:
44+
; CHECK-NEXT: // %bb.0:
45+
; CHECK-NEXT: ld.param.u64 %rd1, [bar_param_0];
46+
; CHECK-NEXT: cvta.to.global.u64 %rd2, %rd1;
47+
; CHECK-NEXT: ld.param.u32 %r1, [bar_param_1];
48+
; CHECK-NEXT: suld.b.1d.b32.trap {%r2}, [surf0, {%r1}];
49+
; CHECK-NEXT: cvt.rn.f32.s32 %f1, %r2;
50+
; CHECK-NEXT: st.global.f32 [%rd2], %f1;
51+
; CHECK-NEXT: ret;
8352
%surfHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @surf0)
8453
%val = tail call i32 @llvm.nvvm.suld.1d.i32.trap(i64 %surfHandle, i32 %idx)
8554
%ret = sitofp i32 %val to float

‎llvm/test/CodeGen/NVPTX/surf-write-cuda.ll‎

Lines changed: 23 additions & 46 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,6 @@
11
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2-
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefix=SM20
3-
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_30 -verify-machineinstrs | FileCheck %s --check-prefix=SM30
2+
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
3+
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_30 -verify-machineinstrs | FileCheck %s
44
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
55
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_30 -verify-machineinstrs | %ptxas-verify %}
66

@@ -11,29 +11,17 @@ declare i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1))
1111

1212

1313
define void @foo(i64 %img, i32 %val, i32 %idx) {
14-
; SM20-LABEL: foo(
15-
; SM20: {
16-
; SM20-NEXT: .reg .b32 %r<3>;
17-
; SM20-NEXT: .reg .b64 %rd<2>;
18-
; SM20-EMPTY:
19-
; SM20-NEXT: // %bb.0:
20-
; SM20-NEXT: ld.param.u64 %rd1, [foo_param_0];
21-
; SM20-NEXT: ld.param.u32 %r1, [foo_param_1];
22-
; SM20-NEXT: ld.param.u32 %r2, [foo_param_2];
23-
; SM20-NEXT: sust.b.1d.b32.trap [%rd1, {%r2}], {%r1};
24-
; SM20-NEXT: ret;
25-
;
26-
; SM30-LABEL: foo(
27-
; SM30: {
28-
; SM30-NEXT: .reg .b32 %r<3>;
29-
; SM30-NEXT: .reg .b64 %rd<2>;
30-
; SM30-EMPTY:
31-
; SM30-NEXT: // %bb.0:
32-
; SM30-NEXT: ld.param.u64 %rd1, [foo_param_0];
33-
; SM30-NEXT: ld.param.u32 %r1, [foo_param_1];
34-
; SM30-NEXT: ld.param.u32 %r2, [foo_param_2];
35-
; SM30-NEXT: sust.b.1d.b32.trap [%rd1, {%r2}], {%r1};
36-
; SM30-NEXT: ret;
14+
; CHECK-LABEL: foo(
15+
; CHECK: {
16+
; CHECK-NEXT: .reg .b32 %r<3>;
17+
; CHECK-NEXT: .reg .b64 %rd<2>;
18+
; CHECK-EMPTY:
19+
; CHECK-NEXT: // %bb.0:
20+
; CHECK-NEXT: ld.param.u64 %rd1, [foo_param_0];
21+
; CHECK-NEXT: ld.param.u32 %r1, [foo_param_1];
22+
; CHECK-NEXT: ld.param.u32 %r2, [foo_param_2];
23+
; CHECK-NEXT: sust.b.1d.b32.trap [%rd1, {%r2}], {%r1};
24+
; CHECK-NEXT: ret;
3725
tail call void @llvm.nvvm.sust.b.1d.i32.trap(i64 %img, i32 %idx, i32 %val)
3826
ret void
3927
}
@@ -43,27 +31,16 @@ define void @foo(i64 %img, i32 %val, i32 %idx) {
4331

4432

4533
define void @bar(i32 %val, i32 %idx) {
46-
; SM20-LABEL: bar(
47-
; SM20: {
48-
; SM20-NEXT: .reg .b32 %r<3>;
49-
; SM20-NEXT: .reg .b64 %rd<2>;
50-
; SM20-EMPTY:
51-
; SM20-NEXT: // %bb.0:
52-
; SM20-NEXT: ld.param.u32 %r1, [bar_param_0];
53-
; SM20-NEXT: ld.param.u32 %r2, [bar_param_1];
54-
; SM20-NEXT: sust.b.1d.b32.trap [surf0, {%r2}], {%r1};
55-
; SM20-NEXT: ret;
56-
;
57-
; SM30-LABEL: bar(
58-
; SM30: {
59-
; SM30-NEXT: .reg .b32 %r<3>;
60-
; SM30-NEXT: .reg .b64 %rd<2>;
61-
; SM30-EMPTY:
62-
; SM30-NEXT: // %bb.0:
63-
; SM30-NEXT: ld.param.u32 %r1, [bar_param_0];
64-
; SM30-NEXT: ld.param.u32 %r2, [bar_param_1];
65-
; SM30-NEXT: sust.b.1d.b32.trap [surf0, {%r2}], {%r1};
66-
; SM30-NEXT: ret;
34+
; CHECK-LABEL: bar(
35+
; CHECK: {
36+
; CHECK-NEXT: .reg .b32 %r<3>;
37+
; CHECK-NEXT: .reg .b64 %rd<2>;
38+
; CHECK-EMPTY:
39+
; CHECK-NEXT: // %bb.0:
40+
; CHECK-NEXT: ld.param.u32 %r1, [bar_param_0];
41+
; CHECK-NEXT: ld.param.u32 %r2, [bar_param_1];
42+
; CHECK-NEXT: sust.b.1d.b32.trap [surf0, {%r2}], {%r1};
43+
; CHECK-NEXT: ret;
6744
%surfHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @surf0)
6845
tail call void @llvm.nvvm.sust.b.1d.i32.trap(i64 %surfHandle, i32 %idx, i32 %val)
6946
ret void

‎llvm/test/CodeGen/NVPTX/tex-read-cuda.ll‎

Lines changed: 55 additions & 111 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,6 @@
11
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2-
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefix=SM20
3-
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_30 -verify-machineinstrs | FileCheck %s --check-prefix=SM30
2+
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
3+
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_30 -verify-machineinstrs | FileCheck %s
44
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
55
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_30 -verify-machineinstrs | %ptxas-verify %}
66

@@ -11,35 +11,20 @@ declare { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64,
1111
declare i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1))
1212

1313
define void @foo(i64 %img, ptr %red, i32 %idx) {
14-
; SM20-LABEL: foo(
15-
; SM20: {
16-
; SM20-NEXT: .reg .b32 %r<2>;
17-
; SM20-NEXT: .reg .f32 %f<5>;
18-
; SM20-NEXT: .reg .b64 %rd<4>;
19-
; SM20-EMPTY:
20-
; SM20-NEXT: // %bb.0:
21-
; SM20-NEXT: ld.param.u64 %rd1, [foo_param_0];
22-
; SM20-NEXT: ld.param.u64 %rd2, [foo_param_1];
23-
; SM20-NEXT: cvta.to.global.u64 %rd3, %rd2;
24-
; SM20-NEXT: ld.param.u32 %r1, [foo_param_2];
25-
; SM20-NEXT: tex.1d.v4.f32.s32 {%f1, %f2, %f3, %f4}, [%rd1, {%r1}];
26-
; SM20-NEXT: st.global.f32 [%rd3], %f1;
27-
; SM20-NEXT: ret;
28-
;
29-
; SM30-LABEL: foo(
30-
; SM30: {
31-
; SM30-NEXT: .reg .b32 %r<2>;
32-
; SM30-NEXT: .reg .f32 %f<5>;
33-
; SM30-NEXT: .reg .b64 %rd<4>;
34-
; SM30-EMPTY:
35-
; SM30-NEXT: // %bb.0:
36-
; SM30-NEXT: ld.param.u64 %rd1, [foo_param_0];
37-
; SM30-NEXT: ld.param.u64 %rd2, [foo_param_1];
38-
; SM30-NEXT: cvta.to.global.u64 %rd3, %rd2;
39-
; SM30-NEXT: ld.param.u32 %r1, [foo_param_2];
40-
; SM30-NEXT: tex.1d.v4.f32.s32 {%f1, %f2, %f3, %f4}, [%rd1, {%r1}];
41-
; SM30-NEXT: st.global.f32 [%rd3], %f1;
42-
; SM30-NEXT: ret;
14+
; CHECK-LABEL: foo(
15+
; CHECK: {
16+
; CHECK-NEXT: .reg .b32 %r<2>;
17+
; CHECK-NEXT: .reg .f32 %f<5>;
18+
; CHECK-NEXT: .reg .b64 %rd<4>;
19+
; CHECK-EMPTY:
20+
; CHECK-NEXT: // %bb.0:
21+
; CHECK-NEXT: ld.param.u64 %rd1, [foo_param_0];
22+
; CHECK-NEXT: ld.param.u64 %rd2, [foo_param_1];
23+
; CHECK-NEXT: cvta.to.global.u64 %rd3, %rd2;
24+
; CHECK-NEXT: ld.param.u32 %r1, [foo_param_2];
25+
; CHECK-NEXT: tex.1d.v4.f32.s32 {%f1, %f2, %f3, %f4}, [%rd1, {%r1}];
26+
; CHECK-NEXT: st.global.f32 [%rd3], %f1;
27+
; CHECK-NEXT: ret;
4328
%val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %img, i32 %idx)
4429
%ret = extractvalue { float, float, float, float } %val, 0
4530
store float %ret, ptr %red
@@ -50,33 +35,19 @@ define void @foo(i64 %img, ptr %red, i32 %idx) {
5035
@tex0 = internal addrspace(1) global i64 0, align 8
5136

5237
define void @bar(ptr %red, i32 %idx) {
53-
; SM20-LABEL: bar(
54-
; SM20: {
55-
; SM20-NEXT: .reg .b32 %r<2>;
56-
; SM20-NEXT: .reg .f32 %f<5>;
57-
; SM20-NEXT: .reg .b64 %rd<4>;
58-
; SM20-EMPTY:
59-
; SM20-NEXT: // %bb.0:
60-
; SM20-NEXT: ld.param.u64 %rd1, [bar_param_0];
61-
; SM20-NEXT: cvta.to.global.u64 %rd2, %rd1;
62-
; SM20-NEXT: ld.param.u32 %r1, [bar_param_1];
63-
; SM20-NEXT: tex.1d.v4.f32.s32 {%f1, %f2, %f3, %f4}, [tex0, {%r1}];
64-
; SM20-NEXT: st.global.f32 [%rd2], %f1;
65-
; SM20-NEXT: ret;
66-
;
67-
; SM30-LABEL: bar(
68-
; SM30: {
69-
; SM30-NEXT: .reg .b32 %r<2>;
70-
; SM30-NEXT: .reg .f32 %f<5>;
71-
; SM30-NEXT: .reg .b64 %rd<4>;
72-
; SM30-EMPTY:
73-
; SM30-NEXT: // %bb.0:
74-
; SM30-NEXT: ld.param.u64 %rd1, [bar_param_0];
75-
; SM30-NEXT: cvta.to.global.u64 %rd2, %rd1;
76-
; SM30-NEXT: ld.param.u32 %r1, [bar_param_1];
77-
; SM30-NEXT: tex.1d.v4.f32.s32 {%f1, %f2, %f3, %f4}, [tex0, {%r1}];
78-
; SM30-NEXT: st.global.f32 [%rd2], %f1;
79-
; SM30-NEXT: ret;
38+
; CHECK-LABEL: bar(
39+
; CHECK: {
40+
; CHECK-NEXT: .reg .b32 %r<2>;
41+
; CHECK-NEXT: .reg .f32 %f<5>;
42+
; CHECK-NEXT: .reg .b64 %rd<4>;
43+
; CHECK-EMPTY:
44+
; CHECK-NEXT: // %bb.0:
45+
; CHECK-NEXT: ld.param.u64 %rd1, [bar_param_0];
46+
; CHECK-NEXT: cvta.to.global.u64 %rd2, %rd1;
47+
; CHECK-NEXT: ld.param.u32 %r1, [bar_param_1];
48+
; CHECK-NEXT: tex.1d.v4.f32.s32 {%f1, %f2, %f3, %f4}, [tex0, {%r1}];
49+
; CHECK-NEXT: st.global.f32 [%rd2], %f1;
50+
; CHECK-NEXT: ret;
8051
%texHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @tex0)
8152
%val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %texHandle, i32 %idx)
8253
%ret = extractvalue { float, float, float, float } %val, 0
@@ -87,59 +58,32 @@ define void @bar(ptr %red, i32 %idx) {
8758
declare float @texfunc(i64)
8859

8960
define void @baz(ptr %red, i32 %idx) {
90-
; SM20-LABEL: baz(
91-
; SM20: {
92-
; SM20-NEXT: .reg .b32 %r<2>;
93-
; SM20-NEXT: .reg .f32 %f<8>;
94-
; SM20-NEXT: .reg .b64 %rd<4>;
95-
; SM20-EMPTY:
96-
; SM20-NEXT: // %bb.0:
97-
; SM20-NEXT: ld.param.u64 %rd1, [baz_param_0];
98-
; SM20-NEXT: cvta.to.global.u64 %rd2, %rd1;
99-
; SM20-NEXT: ld.param.u32 %r1, [baz_param_1];
100-
; SM20-NEXT: mov.u64 %rd3, tex0;
101-
; SM20-NEXT: tex.1d.v4.f32.s32 {%f1, %f2, %f3, %f4}, [tex0, {%r1}];
102-
; SM20-NEXT: { // callseq 0, 0
103-
; SM20-NEXT: .param .b64 param0;
104-
; SM20-NEXT: st.param.b64 [param0], %rd3;
105-
; SM20-NEXT: .param .b32 retval0;
106-
; SM20-NEXT: call.uni (retval0),
107-
; SM20-NEXT: texfunc,
108-
; SM20-NEXT: (
109-
; SM20-NEXT: param0
110-
; SM20-NEXT: );
111-
; SM20-NEXT: ld.param.f32 %f5, [retval0];
112-
; SM20-NEXT: } // callseq 0
113-
; SM20-NEXT: add.rn.f32 %f7, %f1, %f5;
114-
; SM20-NEXT: st.global.f32 [%rd2], %f7;
115-
; SM20-NEXT: ret;
116-
;
117-
; SM30-LABEL: baz(
118-
; SM30: {
119-
; SM30-NEXT: .reg .b32 %r<2>;
120-
; SM30-NEXT: .reg .f32 %f<8>;
121-
; SM30-NEXT: .reg .b64 %rd<4>;
122-
; SM30-EMPTY:
123-
; SM30-NEXT: // %bb.0:
124-
; SM30-NEXT: ld.param.u64 %rd1, [baz_param_0];
125-
; SM30-NEXT: cvta.to.global.u64 %rd2, %rd1;
126-
; SM30-NEXT: ld.param.u32 %r1, [baz_param_1];
127-
; SM30-NEXT: mov.u64 %rd3, tex0;
128-
; SM30-NEXT: tex.1d.v4.f32.s32 {%f1, %f2, %f3, %f4}, [tex0, {%r1}];
129-
; SM30-NEXT: { // callseq 0, 0
130-
; SM30-NEXT: .param .b64 param0;
131-
; SM30-NEXT: st.param.b64 [param0], %rd3;
132-
; SM30-NEXT: .param .b32 retval0;
133-
; SM30-NEXT: call.uni (retval0),
134-
; SM30-NEXT: texfunc,
135-
; SM30-NEXT: (
136-
; SM30-NEXT: param0
137-
; SM30-NEXT: );
138-
; SM30-NEXT: ld.param.f32 %f5, [retval0];
139-
; SM30-NEXT: } // callseq 0
140-
; SM30-NEXT: add.rn.f32 %f7, %f1, %f5;
141-
; SM30-NEXT: st.global.f32 [%rd2], %f7;
142-
; SM30-NEXT: ret;
61+
; CHECK-LABEL: baz(
62+
; CHECK: {
63+
; CHECK-NEXT: .reg .b32 %r<2>;
64+
; CHECK-NEXT: .reg .f32 %f<8>;
65+
; CHECK-NEXT: .reg .b64 %rd<4>;
66+
; CHECK-EMPTY:
67+
; CHECK-NEXT: // %bb.0:
68+
; CHECK-NEXT: ld.param.u64 %rd1, [baz_param_0];
69+
; CHECK-NEXT: cvta.to.global.u64 %rd2, %rd1;
70+
; CHECK-NEXT: ld.param.u32 %r1, [baz_param_1];
71+
; CHECK-NEXT: mov.u64 %rd3, tex0;
72+
; CHECK-NEXT: tex.1d.v4.f32.s32 {%f1, %f2, %f3, %f4}, [tex0, {%r1}];
73+
; CHECK-NEXT: { // callseq 0, 0
74+
; CHECK-NEXT: .param .b64 param0;
75+
; CHECK-NEXT: st.param.b64 [param0], %rd3;
76+
; CHECK-NEXT: .param .b32 retval0;
77+
; CHECK-NEXT: call.uni (retval0),
78+
; CHECK-NEXT: texfunc,
79+
; CHECK-NEXT: (
80+
; CHECK-NEXT: param0
81+
; CHECK-NEXT: );
82+
; CHECK-NEXT: ld.param.f32 %f5, [retval0];
83+
; CHECK-NEXT: } // callseq 0
84+
; CHECK-NEXT: add.rn.f32 %f7, %f1, %f5;
85+
; CHECK-NEXT: st.global.f32 [%rd2], %f7;
86+
; CHECK-NEXT: ret;
14387
%texHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @tex0)
14488
%val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %texHandle, i32 %idx)
14589
%ret = extractvalue { float, float, float, float } %val, 0

0 commit comments

Comments
 (0)