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-favor-non-generic -dce | 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 8 ; Verifies nvptx-favor-non-generic correctly optimizes generic address space 9 ; usage to non-generic address space usage for the patterns we claim to handle: 10 ; 1. load cast 11 ; 2. store cast 12 ; 3. load gep cast 13 ; 4. store gep cast 14 ; gep and cast can be an instruction or a constant expression. This function 15 ; tries all possible combinations. 16 define float @ld_st_shared_f32(i32 %i, float %v) { 17 ; IR-LABEL: @ld_st_shared_f32 18 ; IR-NOT: addrspacecast 19 ; PTX-LABEL: ld_st_shared_f32( 20 ; load cast 21 %1 = load float, float* addrspacecast (float addrspace(3)* @scalar to float*), align 4 22 ; PTX: ld.shared.f32 %f{{[0-9]+}}, [scalar]; 23 ; store cast 24 store float %v, float* addrspacecast (float addrspace(3)* @scalar to float*), align 4 25 ; PTX: st.shared.f32 [scalar], %f{{[0-9]+}}; 26 ; use syncthreads to disable optimizations across components 27 call void @llvm.cuda.syncthreads() 28 ; PTX: bar.sync 0; 29 30 ; cast; load 31 %2 = addrspacecast float addrspace(3)* @scalar to float* 32 %3 = load float, float* %2, align 4 33 ; PTX: ld.shared.f32 %f{{[0-9]+}}, [scalar]; 34 ; cast; store 35 store float %v, float* %2, align 4 36 ; PTX: st.shared.f32 [scalar], %f{{[0-9]+}}; 37 call void @llvm.cuda.syncthreads() 38 ; PTX: bar.sync 0; 39 40 ; load gep cast 41 %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 42 ; PTX: ld.shared.f32 %f{{[0-9]+}}, [array+20]; 43 ; store gep cast 44 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 45 ; PTX: st.shared.f32 [array+20], %f{{[0-9]+}}; 46 call void @llvm.cuda.syncthreads() 47 ; PTX: bar.sync 0; 48 49 ; gep cast; load 50 %5 = getelementptr inbounds [10 x float], [10 x float]* addrspacecast ([10 x float] addrspace(3)* @array to [10 x float]*), i32 0, i32 5 51 %6 = load float, float* %5, align 4 52 ; PTX: ld.shared.f32 %f{{[0-9]+}}, [array+20]; 53 ; gep cast; store 54 store float %v, float* %5, align 4 55 ; PTX: st.shared.f32 [array+20], %f{{[0-9]+}}; 56 call void @llvm.cuda.syncthreads() 57 ; PTX: bar.sync 0; 58 59 ; cast; gep; load 60 %7 = addrspacecast [10 x float] addrspace(3)* @array to [10 x float]* 61 %8 = getelementptr inbounds [10 x float], [10 x float]* %7, i32 0, i32 %i 62 %9 = load float, float* %8, align 4 63 ; PTX: ld.shared.f32 %f{{[0-9]+}}, [%{{(r|rl|rd)[0-9]+}}]; 64 ; cast; gep; store 65 store float %v, float* %8, align 4 66 ; PTX: st.shared.f32 [%{{(r|rl|rd)[0-9]+}}], %f{{[0-9]+}}; 67 call void @llvm.cuda.syncthreads() 68 ; PTX: bar.sync 0; 69 70 %sum2 = fadd float %1, %3 71 %sum3 = fadd float %sum2, %4 72 %sum4 = fadd float %sum3, %6 73 %sum5 = fadd float %sum4, %9 74 ret float %sum5 75 } 76 77 ; When hoisting an addrspacecast between different pointer types, replace the 78 ; addrspacecast with a bitcast. 79 define i32 @ld_int_from_float() { 80 ; IR-LABEL: @ld_int_from_float 81 ; IR: load i32, i32 addrspace(3)* bitcast (float addrspace(3)* @scalar to i32 addrspace(3)*) 82 ; PTX-LABEL: ld_int_from_float( 83 ; PTX: ld.shared.u{{(32|64)}} 84 %1 = load i32, i32* addrspacecast(float addrspace(3)* @scalar to i32*), align 4 85 ret i32 %1 86 } 87 88 define i32 @ld_int_from_global_float(float addrspace(1)* %input, i32 %i, i32 %j) { 89 ; IR-LABEL: @ld_int_from_global_float( 90 ; PTX-LABEL: ld_int_from_global_float( 91 %1 = addrspacecast float addrspace(1)* %input to float* 92 %2 = getelementptr float, float* %1, i32 %i 93 ; IR-NEXT: getelementptr float, float addrspace(1)* %input, i32 %i 94 %3 = getelementptr float, float* %2, i32 %j 95 ; IR-NEXT: getelementptr float, float addrspace(1)* {{%[^,]+}}, i32 %j 96 %4 = bitcast float* %3 to i32* 97 ; IR-NEXT: bitcast float addrspace(1)* {{%[^ ]+}} to i32 addrspace(1)* 98 %5 = load i32, i32* %4 99 ; IR-NEXT: load i32, i32 addrspace(1)* {{%.+}} 100 ; PTX-LABEL: ld.global 101 ret i32 %5 102 } 103 104 define void @nested_const_expr() { 105 ; PTX-LABEL: nested_const_expr( 106 ; store 1 to bitcast(gep(addrspacecast(array), 0, 1)) 107 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 108 ; PTX: mov.u32 %r1, 1; 109 ; PTX-NEXT: st.shared.u32 [array+4], %r1; 110 ret void 111 } 112 113 define void @rauw(float addrspace(1)* %input) { 114 %generic_input = addrspacecast float addrspace(1)* %input to float* 115 %addr = getelementptr float, float* %generic_input, i64 10 116 %v = load float, float* %addr 117 store float %v, float* %addr 118 ret void 119 ; IR-LABEL: @rauw( 120 ; IR-NEXT: %1 = getelementptr float, float addrspace(1)* %input, i64 10 121 ; IR-NEXT: %v = load float, float addrspace(1)* %1 122 ; IR-NEXT: store float %v, float addrspace(1)* %1 123 ; IR-NEXT: ret void 124 } 125 126 declare void @llvm.cuda.syncthreads() #3 127 128 attributes #3 = { noduplicate nounwind } 129 130
