62 lines · plain
1; RUN: opt -S -mtriple=amdgcn-unknown-unknown -mcpu=tahiti -passes=amdgpu-promote-alloca -disable-promote-alloca-to-vector < %s | FileCheck %s2 3; Both of these kernels have the same value for4; amdgpu-flat-work-group-size, except one explicitly sets it. This is5; a program visible property which should always take precedence over6; the amdgpu-waves-per-eu optimization hint.7;8; The range is incompatible with the amdgpu-waves-per-eu value, so the9; flat work group size should take precedence implying a requirement10; to support 1024 size workgroups (which exceeds the available LDS11; amount).12 13; CHECK-NOT: @no_flat_workgroup_size.stack14; CHECK-NOT: @explicit_default_workgroup_size.stack15 16; CHECK-LABEL: @no_flat_workgroup_size(17; CHECK: alloca [5 x i32]18; CHECK: store i32 4, ptr addrspace(5) %arrayidx1, align 419define amdgpu_kernel void @no_flat_workgroup_size(ptr addrspace(1) nocapture %out, ptr addrspace(1) nocapture %in) #0 {20entry:21 %stack = alloca [5 x i32], align 4, addrspace(5)22 %0 = load i32, ptr addrspace(1) %in, align 423 %arrayidx1 = getelementptr inbounds [5 x i32], ptr addrspace(5) %stack, i32 0, i32 %024 store i32 4, ptr addrspace(5) %arrayidx1, align 425 %arrayidx2 = getelementptr inbounds i32, ptr addrspace(1) %in, i32 126 %1 = load i32, ptr addrspace(1) %arrayidx2, align 427 %arrayidx3 = getelementptr inbounds [5 x i32], ptr addrspace(5) %stack, i32 0, i32 %128 store i32 5, ptr addrspace(5) %arrayidx3, align 429 %2 = load i32, ptr addrspace(5) %stack, align 430 store i32 %2, ptr addrspace(1) %out, align 431 %arrayidx12 = getelementptr inbounds [5 x i32], ptr addrspace(5) %stack, i32 0, i32 132 %3 = load i32, ptr addrspace(5) %arrayidx1233 %arrayidx13 = getelementptr inbounds i32, ptr addrspace(1) %out, i32 134 store i32 %3, ptr addrspace(1) %arrayidx1335 ret void36}37 38; CHECK-LABEL: @explicit_default_workgroup_size(39; CHECK: alloca [5 x i32]40; CHECK: store i32 4, ptr addrspace(5) %arrayidx1, align 441define amdgpu_kernel void @explicit_default_workgroup_size(ptr addrspace(1) nocapture %out, ptr addrspace(1) nocapture %in) #1 {42entry:43 %stack = alloca [5 x i32], align 4, addrspace(5)44 %0 = load i32, ptr addrspace(1) %in, align 445 %arrayidx1 = getelementptr inbounds [5 x i32], ptr addrspace(5) %stack, i32 0, i32 %046 store i32 4, ptr addrspace(5) %arrayidx1, align 447 %arrayidx2 = getelementptr inbounds i32, ptr addrspace(1) %in, i32 148 %1 = load i32, ptr addrspace(1) %arrayidx2, align 449 %arrayidx3 = getelementptr inbounds [5 x i32], ptr addrspace(5) %stack, i32 0, i32 %150 store i32 5, ptr addrspace(5) %arrayidx3, align 451 %2 = load i32, ptr addrspace(5) %stack, align 452 store i32 %2, ptr addrspace(1) %out, align 453 %arrayidx12 = getelementptr inbounds [5 x i32], ptr addrspace(5) %stack, i32 0, i32 154 %3 = load i32, ptr addrspace(5) %arrayidx1255 %arrayidx13 = getelementptr inbounds i32, ptr addrspace(1) %out, i32 156 store i32 %3, ptr addrspace(1) %arrayidx1357 ret void58}59 60attributes #0 = { "amdgpu-waves-per-eu"="1,1" }61attributes #1 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="1,1024" }62