180 lines · plain
1; RUN: llc < %s -mtriple=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix PTX2; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix PTX3; RUN: opt -mtriple=nvptx-- < %s -S -passes=infer-address-spaces | FileCheck %s --check-prefix IR4; RUN: opt -mtriple=nvptx64-- < %s -S -passes=infer-address-spaces | FileCheck %s --check-prefix IR5; RUN: %if ptxas-ptr32 %{ llc < %s -mtriple=nvptx -mcpu=sm_20 | %ptxas-verify %}6; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}7 8@array = internal addrspace(3) global [10 x float] zeroinitializer, align 49@scalar = internal addrspace(3) global float 0.000000e+00, align 410 11; Verifies nvptx-favor-non-generic correctly optimizes generic address space12; usage to non-generic address space usage for the patterns we claim to handle:13; 1. load cast14; 2. store cast15; 3. load gep cast16; 4. store gep cast17; gep and cast can be an instruction or a constant expression. This function18; tries all possible combinations.19define void @ld_st_shared_f32(i32 %i, float %v) {20; IR-LABEL: @ld_st_shared_f3221; IR-NOT: addrspacecast22; PTX-LABEL: ld_st_shared_f32(23 ; load cast24 %1 = load float, ptr addrspacecast (ptr addrspace(3) @scalar to ptr), align 425 call void @use(float %1)26; PTX: ld.shared.b32 %r{{[0-9]+}}, [scalar];27 ; store cast28 store float %v, ptr addrspacecast (ptr addrspace(3) @scalar to ptr), align 429; PTX: st.shared.b32 [scalar], %r{{[0-9]+}};30 ; use syncthreads to disable optimizations across components31 call void @llvm.nvvm.barrier0()32; PTX: bar.sync 0;33 34 ; cast; load35 %2 = addrspacecast ptr addrspace(3) @scalar to ptr36 %3 = load float, ptr %2, align 437 call void @use(float %3)38; PTX: ld.shared.b32 %r{{[0-9]+}}, [scalar];39 ; cast; store40 store float %v, ptr %2, align 441; PTX: st.shared.b32 [scalar], %r{{[0-9]+}};42 call void @llvm.nvvm.barrier0()43; PTX: bar.sync 0;44 45 ; load gep cast46 %4 = load float, ptr getelementptr inbounds ([10 x float], ptr addrspacecast (ptr addrspace(3) @array to ptr), i32 0, i32 5), align 447 call void @use(float %4)48; PTX: ld.shared.b32 %r{{[0-9]+}}, [array+20];49 ; store gep cast50 store float %v, ptr getelementptr inbounds ([10 x float], ptr addrspacecast (ptr addrspace(3) @array to ptr), i32 0, i32 5), align 451; PTX: st.shared.b32 [array+20], %r{{[0-9]+}};52 call void @llvm.nvvm.barrier0()53; PTX: bar.sync 0;54 55 ; gep cast; load56 %5 = getelementptr inbounds [10 x float], ptr addrspacecast (ptr addrspace(3) @array to ptr), i32 0, i32 557 %6 = load float, ptr %5, align 458 call void @use(float %6)59; PTX: ld.shared.b32 %r{{[0-9]+}}, [array+20];60 ; gep cast; store61 store float %v, ptr %5, align 462; PTX: st.shared.b32 [array+20], %r{{[0-9]+}};63 call void @llvm.nvvm.barrier0()64; PTX: bar.sync 0;65 66 ; cast; gep; load67 %7 = addrspacecast ptr addrspace(3) @array to ptr68 %8 = getelementptr inbounds [10 x float], ptr %7, i32 0, i32 %i69 %9 = load float, ptr %8, align 470 call void @use(float %9)71; PTX: ld.shared.b32 %r{{[0-9]+}}, [%{{(r|rl|rd)[0-9]+}}];72 ; cast; gep; store73 store float %v, ptr %8, align 474; PTX: st.shared.b32 [%{{(r|rl|rd)[0-9]+}}], %r{{[0-9]+}};75 call void @llvm.nvvm.barrier0()76; PTX: bar.sync 0;77 78 ret void79}80 81; When hoisting an addrspacecast between different pointer types, replace the82; addrspacecast with a bitcast.83define i32 @ld_int_from_float() {84; IR-LABEL: @ld_int_from_float85; IR: load i32, ptr addrspace(3) @scalar86; PTX-LABEL: ld_int_from_float(87; PTX: ld.shared.b{{(32|64)}}88 %1 = load i32, ptr addrspacecast(ptr addrspace(3) @scalar to ptr), align 489 ret i32 %190}91 92define i32 @ld_int_from_global_float(ptr addrspace(1) %input, i32 %i, i32 %j) {93; IR-LABEL: @ld_int_from_global_float(94; PTX-LABEL: ld_int_from_global_float(95 %1 = addrspacecast ptr addrspace(1) %input to ptr96 %2 = getelementptr float, ptr %1, i32 %i97; IR-NEXT: getelementptr float, ptr addrspace(1) %input, i32 %i98 %3 = getelementptr float, ptr %2, i32 %j99; IR-NEXT: getelementptr float, ptr addrspace(1) {{%[^,]+}}, i32 %j100 %4 = load i32, ptr %3101; IR-NEXT: load i32, ptr addrspace(1) {{%.+}}102; PTX-LABEL: ld.global103 ret i32 %4104}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, ptr getelementptr ([10 x float], ptr addrspacecast (ptr addrspace(3) @array to ptr), i64 0, i64 1), align 4110; PTX: st.shared.b32 [array+4], 1;111 ret void112}113 114define void @rauw(ptr addrspace(1) %input) {115 %generic_input = addrspacecast ptr addrspace(1) %input to ptr116 %addr = getelementptr float, ptr %generic_input, i64 10117 %v = load float, ptr %addr118 store float %v, ptr %addr119 ret void120; IR-LABEL: @rauw(121; IR-NEXT: %addr = getelementptr float, ptr addrspace(1) %input, i64 10122; IR-NEXT: %v = load float, ptr addrspace(1) %addr123; IR-NEXT: store float %v, ptr addrspace(1) %addr124; IR-NEXT: ret void125}126 127define void @loop() {128; IR-LABEL: @loop(129entry:130 %p = addrspacecast ptr addrspace(3) @array to ptr131 %end = getelementptr float, ptr %p, i64 10132 br label %loop133 134loop:135 %i = phi ptr [ %p, %entry ], [ %i2, %loop ]136; IR: phi ptr addrspace(3) [ @array, %entry ], [ %i2, %loop ]137 %v = load float, ptr %i138; IR: %v = load float, ptr addrspace(3) %i139 call void @use(float %v)140 %i2 = getelementptr float, ptr %i, i64 1141; IR: %i2 = getelementptr float, ptr addrspace(3) %i, i64 1142 %exit_cond = icmp eq ptr %i2, %end143 br i1 %exit_cond, label %exit, label %loop144 145exit:146 ret void147}148 149@generic_end = external global ptr150 151define void @loop_with_generic_bound() {152; IR-LABEL: @loop_with_generic_bound(153entry:154 %p = addrspacecast ptr addrspace(3) @array to ptr155 %end = load ptr, ptr @generic_end156 br label %loop157 158loop:159 %i = phi ptr [ %p, %entry ], [ %i2, %loop ]160; IR: phi ptr addrspace(3) [ @array, %entry ], [ %i2, %loop ]161 %v = load float, ptr %i162; IR: %v = load float, ptr addrspace(3) %i163 call void @use(float %v)164 %i2 = getelementptr float, ptr %i, i64 1165; IR: %i2 = getelementptr float, ptr addrspace(3) %i, i64 1166 %exit_cond = icmp eq ptr %i2, %end167; IR: addrspacecast ptr addrspace(3) %i2 to ptr168; IR: icmp eq ptr %{{[0-9]+}}, %end169 br i1 %exit_cond, label %exit, label %loop170 171exit:172 ret void173}174 175declare void @llvm.nvvm.barrier0() #3176 177declare void @use(float)178 179attributes #3 = { noduplicate nounwind }180