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