[MIParser] Set RegClassOrRegBank during instruction parsing
[llvm-complete.git] / test / CodeGen / NVPTX / access-non-generic.ll
blob62520be2cf5a07e8ef0ebb36f2bf1b40719564ea
1 ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix PTX
2 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix PTX
3 ; RUN: opt -mtriple=nvptx-- < %s -S -infer-address-spaces | FileCheck %s --check-prefix IR
4 ; RUN: opt -mtriple=nvptx64-- < %s -S -infer-address-spaces | FileCheck %s --check-prefix IR
6 @array = internal addrspace(3) global [10 x float] zeroinitializer, align 4
7 @scalar = internal addrspace(3) global float 0.000000e+00, align 4
9 ; Verifies nvptx-favor-non-generic correctly optimizes generic address space
10 ; usage to non-generic address space usage for the patterns we claim to handle:
11 ; 1. load cast
12 ; 2. store cast
13 ; 3. load gep cast
14 ; 4. store gep cast
15 ; gep and cast can be an instruction or a constant expression. This function
16 ; tries all possible combinations.
17 define void @ld_st_shared_f32(i32 %i, float %v) {
18 ; IR-LABEL: @ld_st_shared_f32
19 ; IR-NOT: addrspacecast
20 ; PTX-LABEL: ld_st_shared_f32(
21   ; load cast
22   %1 = load float, float* addrspacecast (float addrspace(3)* @scalar to float*), align 4
23   call void @use(float %1)
24 ; PTX: ld.shared.f32 %f{{[0-9]+}}, [scalar];
25   ; store cast
26   store float %v, float* addrspacecast (float addrspace(3)* @scalar to float*), align 4
27 ; PTX: st.shared.f32 [scalar], %f{{[0-9]+}};
28   ; use syncthreads to disable optimizations across components
29   call void @llvm.nvvm.barrier0()
30 ; PTX: bar.sync 0;
32   ; cast; load
33   %2 = addrspacecast float addrspace(3)* @scalar to float*
34   %3 = load float, float* %2, align 4
35   call void @use(float %3)
36 ; PTX: ld.shared.f32 %f{{[0-9]+}}, [scalar];
37   ; cast; store
38   store float %v, float* %2, align 4
39 ; PTX: st.shared.f32 [scalar], %f{{[0-9]+}};
40   call void @llvm.nvvm.barrier0()
41 ; PTX: bar.sync 0;
43   ; load gep cast
44   %4 = load float, float* getelementptr inbounds ([10 x float], [10 x float]* addrspacecast ([10 x float] addrspace(3)* @array to [10 x float]*), i32 0, i32 5), align 4
45   call void @use(float %4)
46 ; PTX: ld.shared.f32 %f{{[0-9]+}}, [array+20];
47   ; store gep cast
48   store float %v, float* getelementptr inbounds ([10 x float], [10 x float]* addrspacecast ([10 x float] addrspace(3)* @array to [10 x float]*), i32 0, i32 5), align 4
49 ; PTX: st.shared.f32 [array+20], %f{{[0-9]+}};
50   call void @llvm.nvvm.barrier0()
51 ; PTX: bar.sync 0;
53   ; gep cast; load
54   %5 = getelementptr inbounds [10 x float], [10 x float]* addrspacecast ([10 x float] addrspace(3)* @array to [10 x float]*), i32 0, i32 5
55   %6 = load float, float* %5, align 4
56   call void @use(float %6)
57 ; PTX: ld.shared.f32 %f{{[0-9]+}}, [array+20];
58   ; gep cast; store
59   store float %v, float* %5, align 4
60 ; PTX: st.shared.f32 [array+20], %f{{[0-9]+}};
61   call void @llvm.nvvm.barrier0()
62 ; PTX: bar.sync 0;
64   ; cast; gep; load
65   %7 = addrspacecast [10 x float] addrspace(3)* @array to [10 x float]*
66   %8 = getelementptr inbounds [10 x float], [10 x float]* %7, i32 0, i32 %i
67   %9 = load float, float* %8, align 4
68   call void @use(float %9)
69 ; PTX: ld.shared.f32 %f{{[0-9]+}}, [%{{(r|rl|rd)[0-9]+}}];
70   ; cast; gep; store
71   store float %v, float* %8, align 4
72 ; PTX: st.shared.f32 [%{{(r|rl|rd)[0-9]+}}], %f{{[0-9]+}};
73   call void @llvm.nvvm.barrier0()
74 ; PTX: bar.sync 0;
76   ret void
79 ; When hoisting an addrspacecast between different pointer types, replace the
80 ; addrspacecast with a bitcast.
81 define i32 @ld_int_from_float() {
82 ; IR-LABEL: @ld_int_from_float
83 ; IR: load i32, i32 addrspace(3)* bitcast (float addrspace(3)* @scalar to i32 addrspace(3)*)
84 ; PTX-LABEL: ld_int_from_float(
85 ; PTX: ld.shared.u{{(32|64)}}
86   %1 = load i32, i32* addrspacecast(float addrspace(3)* @scalar to i32*), align 4
87   ret i32 %1
90 define i32 @ld_int_from_global_float(float addrspace(1)* %input, i32 %i, i32 %j) {
91 ; IR-LABEL: @ld_int_from_global_float(
92 ; PTX-LABEL: ld_int_from_global_float(
93   %1 = addrspacecast float addrspace(1)* %input to float*
94   %2 = getelementptr float, float* %1, i32 %i
95 ; IR-NEXT: getelementptr float, float addrspace(1)* %input, i32 %i
96   %3 = getelementptr float, float* %2, i32 %j
97 ; IR-NEXT: getelementptr float, float addrspace(1)* {{%[^,]+}}, i32 %j
98   %4 = bitcast float* %3 to i32*
99 ; IR-NEXT: bitcast float addrspace(1)* {{%[^ ]+}} to i32 addrspace(1)*
100   %5 = load i32, i32* %4
101 ; IR-NEXT: load i32, i32 addrspace(1)* {{%.+}}
102 ; PTX-LABEL: ld.global
103   ret i32 %5
106 define void @nested_const_expr() {
107 ; PTX-LABEL: nested_const_expr(
108   ; store 1 to bitcast(gep(addrspacecast(array), 0, 1))
109   store i32 1, i32* bitcast (float* getelementptr ([10 x float], [10 x float]* addrspacecast ([10 x float] addrspace(3)* @array to [10 x float]*), i64 0, i64 1) to i32*), align 4
110 ; PTX: mov.u32 %r1, 1;
111 ; PTX-NEXT: st.shared.u32 [array+4], %r1;
112   ret void
115 define void @rauw(float addrspace(1)* %input) {
116   %generic_input = addrspacecast float addrspace(1)* %input to float*
117   %addr = getelementptr float, float* %generic_input, i64 10
118   %v = load float, float* %addr
119   store float %v, float* %addr
120   ret void
121 ; IR-LABEL: @rauw(
122 ; IR-NEXT: %addr = getelementptr float, float addrspace(1)* %input, i64 10
123 ; IR-NEXT: %v = load float, float addrspace(1)* %addr
124 ; IR-NEXT: store float %v, float addrspace(1)* %addr
125 ; IR-NEXT: ret void
128 define void @loop() {
129 ; IR-LABEL: @loop(
130 entry:
131   %p = addrspacecast [10 x float] addrspace(3)* @array to float*
132   %end = getelementptr float, float* %p, i64 10
133   br label %loop
135 loop:
136   %i = phi float* [ %p, %entry ], [ %i2, %loop ]
137 ; IR: phi float addrspace(3)* [ %p, %entry ], [ %i2, %loop ]
138   %v = load float, float* %i
139 ; IR: %v = load float, float addrspace(3)* %i
140   call void @use(float %v)
141   %i2 = getelementptr float, float* %i, i64 1
142 ; IR: %i2 = getelementptr float, float addrspace(3)* %i, i64 1
143   %exit_cond = icmp eq float* %i2, %end
144   br i1 %exit_cond, label %exit, label %loop
146 exit:
147   ret void
150 @generic_end = external global float*
152 define void @loop_with_generic_bound() {
153 ; IR-LABEL: @loop_with_generic_bound(
154 entry:
155   %p = addrspacecast [10 x float] addrspace(3)* @array to float*
156   %end = load float*, float** @generic_end
157   br label %loop
159 loop:
160   %i = phi float* [ %p, %entry ], [ %i2, %loop ]
161 ; IR: phi float addrspace(3)* [ %p, %entry ], [ %i2, %loop ]
162   %v = load float, float* %i
163 ; IR: %v = load float, float addrspace(3)* %i
164   call void @use(float %v)
165   %i2 = getelementptr float, float* %i, i64 1
166 ; IR: %i2 = getelementptr float, float addrspace(3)* %i, i64 1
167   %exit_cond = icmp eq float* %i2, %end
168 ; IR: addrspacecast float addrspace(3)* %i2 to float*
169 ; IR: icmp eq float* %{{[0-9]+}}, %end
170   br i1 %exit_cond, label %exit, label %loop
172 exit:
173   ret void
176 declare void @llvm.nvvm.barrier0() #3
178 declare void @use(float)
180 attributes #3 = { noduplicate nounwind }