; RUN: opt -S -mtriple=amdgcn-unknown-unknown -mcpu=tahiti -passes=amdgpu-promote-alloca -disable-promote-alloca-to-vector < %s | FileCheck %s

; Both of these kernels have the same value for
; amdgpu-flat-work-group-size, except one explicitly sets it. This is
; a program visible property which should always take precedence over
; the amdgpu-waves-per-eu optimization hint.
;
; The range is incompatible with the amdgpu-waves-per-eu value, so the
; flat work group size should take precedence implying a requirement
; to support 1024 size workgroups (which exceeds the available LDS
; amount).

; CHECK-NOT: @no_flat_workgroup_size.stack
; CHECK-NOT: @explicit_default_workgroup_size.stack

; CHECK-LABEL: @no_flat_workgroup_size(
; CHECK: alloca [5 x i32]
; CHECK: store i32 4, ptr addrspace(5) %arrayidx1, align 4
define amdgpu_kernel void @no_flat_workgroup_size(ptr addrspace(1) nocapture %out, ptr addrspace(1) nocapture %in) #0 {
entry:
  %stack = alloca [5 x i32], align 4, addrspace(5)
  %0 = load i32, ptr addrspace(1) %in, align 4
  %arrayidx1 = getelementptr inbounds [5 x i32], ptr addrspace(5) %stack, i32 0, i32 %0
  store i32 4, ptr addrspace(5) %arrayidx1, align 4
  %arrayidx2 = getelementptr inbounds i32, ptr addrspace(1) %in, i32 1
  %1 = load i32, ptr addrspace(1) %arrayidx2, align 4
  %arrayidx3 = getelementptr inbounds [5 x i32], ptr addrspace(5) %stack, i32 0, i32 %1
  store i32 5, ptr addrspace(5) %arrayidx3, align 4
  %2 = load i32, ptr addrspace(5) %stack, align 4
  store i32 %2, ptr addrspace(1) %out, align 4
  %arrayidx12 = getelementptr inbounds [5 x i32], ptr addrspace(5) %stack, i32 0, i32 1
  %3 = load i32, ptr addrspace(5) %arrayidx12
  %arrayidx13 = getelementptr inbounds i32, ptr addrspace(1) %out, i32 1
  store i32 %3, ptr addrspace(1) %arrayidx13
  ret void
}

; CHECK-LABEL: @explicit_default_workgroup_size(
; CHECK: alloca [5 x i32]
; CHECK: store i32 4, ptr addrspace(5) %arrayidx1, align 4
define amdgpu_kernel void @explicit_default_workgroup_size(ptr addrspace(1) nocapture %out, ptr addrspace(1) nocapture %in) #1 {
entry:
  %stack = alloca [5 x i32], align 4, addrspace(5)
  %0 = load i32, ptr addrspace(1) %in, align 4
  %arrayidx1 = getelementptr inbounds [5 x i32], ptr addrspace(5) %stack, i32 0, i32 %0
  store i32 4, ptr addrspace(5) %arrayidx1, align 4
  %arrayidx2 = getelementptr inbounds i32, ptr addrspace(1) %in, i32 1
  %1 = load i32, ptr addrspace(1) %arrayidx2, align 4
  %arrayidx3 = getelementptr inbounds [5 x i32], ptr addrspace(5) %stack, i32 0, i32 %1
  store i32 5, ptr addrspace(5) %arrayidx3, align 4
  %2 = load i32, ptr addrspace(5) %stack, align 4
  store i32 %2, ptr addrspace(1) %out, align 4
  %arrayidx12 = getelementptr inbounds [5 x i32], ptr addrspace(5) %stack, i32 0, i32 1
  %3 = load i32, ptr addrspace(5) %arrayidx12
  %arrayidx13 = getelementptr inbounds i32, ptr addrspace(1) %out, i32 1
  store i32 %3, ptr addrspace(1) %arrayidx13
  ret void
}

attributes #0 = { "amdgpu-waves-per-eu"="1,1" }
attributes #1 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="1,1024" }