Index: llvm/docs/AMDGPUUsage.rst =================================================================== --- llvm/docs/AMDGPUUsage.rst +++ llvm/docs/AMDGPUUsage.rst @@ -856,6 +856,8 @@ "amdgpu-flat-work-group-size"="min,max" Specify the minimum and maximum flat work group sizes that will be specified when the kernel is dispatched. Generated by the ``amdgpu_flat_work_group_size`` CLANG attribute [CLANG-ATTR]_. + The implied default value is 1,1024. + "amdgpu-implicitarg-num-bytes"="n" Number of kernel argument bytes to add to the kernel argument block size for the implicit arguments. This varies by OS and language (for OpenCL see @@ -866,7 +868,11 @@ ``amdgpu_num_vgpr`` CLANG attribute [CLANG-ATTR]_. "amdgpu-waves-per-eu"="m,n" Specify the minimum and maximum number of waves per execution unit. Generated by the ``amdgpu_waves_per_eu`` - CLANG attribute [CLANG-ATTR]_. + CLANG attribute [CLANG-ATTR]_. This is an optimization hint, and + the backend may not be able to satisfy the request. If the specified range + is incompatible with the function's "amdgpu-flat-work-group-size" value, + the implied occupancy bounds by the workgroup size takes precedence. + "amdgpu-ieee" true/false. Specify whether the function expects the IEEE field of the mode register to be set on entry. Overrides the default for the calling convention. Index: llvm/lib/Target/AMDGPU/AMDGPUSubtarget.cpp =================================================================== --- llvm/lib/Target/AMDGPU/AMDGPUSubtarget.cpp +++ llvm/lib/Target/AMDGPU/AMDGPUSubtarget.cpp @@ -547,8 +547,6 @@ unsigned MinImpliedByFlatWorkGroupSize = getWavesPerEUForWorkGroup(FlatWorkGroupSizes.second); Default.first = MinImpliedByFlatWorkGroupSize; - bool RequestedFlatWorkGroupSize = - F.hasFnAttribute("amdgpu-flat-work-group-size"); // Requested minimum/maximum number of waves per execution unit. std::pair Requested = AMDGPU::getIntegerPairAttribute( @@ -565,8 +563,7 @@ // Make sure requested values are compatible with values implied by requested // minimum/maximum flat work group sizes. - if (RequestedFlatWorkGroupSize && - Requested.first < MinImpliedByFlatWorkGroupSize) + if (Requested.first < MinImpliedByFlatWorkGroupSize) return Default; return Requested; Index: llvm/test/CodeGen/AMDGPU/default-flat-work-group-size-overrides-waves-per-eu.ll =================================================================== --- /dev/null +++ llvm/test/CodeGen/AMDGPU/default-flat-work-group-size-overrides-waves-per-eu.ll @@ -0,0 +1,63 @@ +; RUN: opt -S -mtriple=amdgcn-unknown-unknown -mcpu=tahiti -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, i32 addrspace(5)* %arrayidx1, align 4 +define amdgpu_kernel void @no_flat_workgroup_size(i32 addrspace(1)* nocapture %out, i32 addrspace(1)* nocapture %in) #0 { +entry: + %stack = alloca [5 x i32], align 4, addrspace(5) + %0 = load i32, i32 addrspace(1)* %in, align 4 + %arrayidx1 = getelementptr inbounds [5 x i32], [5 x i32] addrspace(5)* %stack, i32 0, i32 %0 + store i32 4, i32 addrspace(5)* %arrayidx1, align 4 + %arrayidx2 = getelementptr inbounds i32, i32 addrspace(1)* %in, i32 1 + %1 = load i32, i32 addrspace(1)* %arrayidx2, align 4 + %arrayidx3 = getelementptr inbounds [5 x i32], [5 x i32] addrspace(5)* %stack, i32 0, i32 %1 + store i32 5, i32 addrspace(5)* %arrayidx3, align 4 + %arrayidx10 = getelementptr inbounds [5 x i32], [5 x i32] addrspace(5)* %stack, i32 0, i32 0 + %2 = load i32, i32 addrspace(5)* %arrayidx10, align 4 + store i32 %2, i32 addrspace(1)* %out, align 4 + %arrayidx12 = getelementptr inbounds [5 x i32], [5 x i32] addrspace(5)* %stack, i32 0, i32 1 + %3 = load i32, i32 addrspace(5)* %arrayidx12 + %arrayidx13 = getelementptr inbounds i32, i32 addrspace(1)* %out, i32 1 + store i32 %3, i32 addrspace(1)* %arrayidx13 + ret void +} + +; CHECK-LABEL: @explicit_default_workgroup_size( +; CHECK: alloca [5 x i32] +; CHECK: store i32 4, i32 addrspace(5)* %arrayidx1, align 4 +define amdgpu_kernel void @explicit_default_workgroup_size(i32 addrspace(1)* nocapture %out, i32 addrspace(1)* nocapture %in) #1 { +entry: + %stack = alloca [5 x i32], align 4, addrspace(5) + %0 = load i32, i32 addrspace(1)* %in, align 4 + %arrayidx1 = getelementptr inbounds [5 x i32], [5 x i32] addrspace(5)* %stack, i32 0, i32 %0 + store i32 4, i32 addrspace(5)* %arrayidx1, align 4 + %arrayidx2 = getelementptr inbounds i32, i32 addrspace(1)* %in, i32 1 + %1 = load i32, i32 addrspace(1)* %arrayidx2, align 4 + %arrayidx3 = getelementptr inbounds [5 x i32], [5 x i32] addrspace(5)* %stack, i32 0, i32 %1 + store i32 5, i32 addrspace(5)* %arrayidx3, align 4 + %arrayidx10 = getelementptr inbounds [5 x i32], [5 x i32] addrspace(5)* %stack, i32 0, i32 0 + %2 = load i32, i32 addrspace(5)* %arrayidx10, align 4 + store i32 %2, i32 addrspace(1)* %out, align 4 + %arrayidx12 = getelementptr inbounds [5 x i32], [5 x i32] addrspace(5)* %stack, i32 0, i32 1 + %3 = load i32, i32 addrspace(5)* %arrayidx12 + %arrayidx13 = getelementptr inbounds i32, i32 addrspace(1)* %out, i32 1 + store i32 %3, i32 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" }