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