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 5 6 @array = internal addrspace(3) global [10 x float] zeroinitializer, align 4 7 @scalar = internal addrspace(3) global float 0.000000e+00, align 4 8 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; 31 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; 42 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; 52 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; 63 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; 75 76 ret void 77 } 78 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 88 } 89 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 104 } 105 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 113 } 114 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 126 } 127 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 134 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 145 146 exit: 147 ret void 148 } 149 150 @generic_end = external global float* 151 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 158 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 171 172 exit: 173 ret void 174 } 175 176 declare void @llvm.nvvm.barrier0() #3 177 178 declare void @use(float) 179 180 attributes #3 = { noduplicate nounwind } 181
