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. 17define 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. 81define 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 90define 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 106define 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 115define 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 128define void @loop() { 129; IR-LABEL: @loop( 130entry: 131 %p = addrspacecast [10 x float] addrspace(3)* @array to float* 132 %end = getelementptr float, float* %p, i64 10 133 br label %loop 134 135loop: 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 146exit: 147 ret void 148} 149 150@generic_end = external global float* 151 152define void @loop_with_generic_bound() { 153; IR-LABEL: @loop_with_generic_bound( 154entry: 155 %p = addrspacecast [10 x float] addrspace(3)* @array to float* 156 %end = load float*, float** @generic_end 157 br label %loop 158 159loop: 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 172exit: 173 ret void 174} 175 176declare void @llvm.nvvm.barrier0() #3 177 178declare void @use(float) 179 180attributes #3 = { noduplicate nounwind } 181