Skip to content

Commit

Permalink
[codegen] Add max_workgroup_counts to TargetWgpAttr
Browse files Browse the repository at this point in the history 
This commit adds a max_workgroup_counts to the workgroup processor
information attribute and sets values for the known targets. Some of
these values may be underestimates as I was not able to locate
information on their values.

This field is added so that we can annotate calls to workgroup.id and
workgroup.count with upper bounds, neabling range inference and
strength reduction.

Note that in some cases (for instance, AMD) we give a
max_workgroup_counts value lower than what is actually supported
because a grid dimension greater than int32_max would be sign-extended
to a negative number to meet the 64-bit nature of `index`.

(This PR is split out of #17707)

Signed-off-by: Krzysztof Drewniak <[email protected]>
  • Loading branch information
@krzysz00
krzysz00 committed Aug 5, 2024
1 parent 09dc003 commit b92a6ce
Show file tree
Hide file tree
Showing 25 changed files with 242 additions and 114 deletions.
3 changes: 2 additions & 1 deletion compiler/plugins/target/MetalSPIRV/test/smoketest.mlir
View file
Original file line number Diff line number Diff line change
Expand Up @@ -6,7 +6,8 @@ module attributes {
#hal.executable.target<"metal-spirv", "metal-msl-fb", {
iree.gpu.target = #iree_gpu.target<arch = "", features = "spirv:v1.3,cap:Shader", wgp = <
compute = fp32|int32, storage = b32, subgroup = none, dot = none, mma = [], subgroup_size_choices = [32],
max_workgroup_sizes = [128, 128, 64], max_thread_count_per_workgroup = 128, max_workgroup_memory_bytes = 16384>>
max_workgroup_sizes = [128, 128, 64], max_thread_count_per_workgroup = 128, max_workgroup_memory_bytes = 16384,
max_workgroup_counts = [65535, 65535, 65535]>>
}>
]> : !hal.device
]
Expand Down
3 changes: 2 additions & 1 deletion compiler/plugins/target/ROCM/test/target_device_features.mlir
View file
Original file line number Diff line number Diff line change
Expand Up @@ -8,7 +8,8 @@
// GFX942-SAME: subgroup = shuffle|arithmetic, dot = dp4xi8toi32,
// GFX942-SAME: mma = [<MFMA_F32_16x16x4_F32>, <MFMA_F16_16x16x16_F32>, <MFMA_F16_32x32x8_F32>, <MFMA_F8E4M3FNUZ_16x16x32_F32>, <MFMA_I8_16x16x32_I32>, <MFMA_I8_32x32x16_I32>],
// GFX942-SAME: subgroup_size_choices = [64], max_workgroup_sizes = [1024, 1024, 1024],
// GFX942-SAME: max_thread_count_per_workgroup = 1024, max_workgroup_memory_bytes = 65536>,
// GFX942-SAME: max_thread_count_per_workgroup = 1024, max_workgroup_memory_bytes = 65536,
// GFX942-SAME: max_workgroup_counts = [2147483647, 2147483647, 2147483647]>,
// GFX942-SAME: chip = <wgp_count = 304>>

// GFX940: target = #iree_gpu.target<arch = "gfx940",
Expand Down
3 changes: 2 additions & 1 deletion compiler/plugins/target/VulkanSPIRV/test/smoketest.mlir
View file
Original file line number Diff line number Diff line change
Expand Up @@ -6,7 +6,8 @@ module attributes {
#hal.executable.target<"vulkan-spirv", "vulkan-spirv-fb", {
iree.gpu.target = #iree_gpu.target<arch = "", features = "spirv:v1.3,cap:Shader", wgp = <
compute = fp32|int32, storage = b32, subgroup = none, dot = none, mma = [], subgroup_size_choices = [32, 32],
max_workgroup_sizes = [128, 128, 64], max_thread_count_per_workgroup = 128, max_workgroup_memory_bytes = 16384>>
max_workgroup_sizes = [128, 128, 64], max_thread_count_per_workgroup = 128, max_workgroup_memory_bytes = 16384,
max_workgroup_counts = [65535, 65535, 65535]>>
}>
]> : !hal.device
]
Expand Down
3 changes: 2 additions & 1 deletion compiler/plugins/target/WebGPUSPIRV/test/smoketest.mlir
View file
Original file line number Diff line number Diff line change
Expand Up @@ -7,7 +7,8 @@ module attributes {
#hal.executable.target<"webgpu-spirv", "webgpu-wgsl-fb", {
iree.gpu.target = #iree_gpu.target<arch = "", features = "spirv:v1.0,cap:Shader,ext:SPV_KHR_storage_buffer_storage_class", wgp = <
compute = fp32|int32, storage = b32, subgroup = none, dot = none, mma = [], subgroup_size_choices = [32],
max_workgroup_sizes = [128, 128, 64], max_thread_count_per_workgroup = 128, max_workgroup_memory_bytes = 16384>>
max_workgroup_sizes = [128, 128, 64], max_thread_count_per_workgroup = 128, max_workgroup_memory_bytes = 16384,
max_workgroup_counts = [65535, 65535, 65535]>>
}>
]> : !hal.device
]
Expand Down
2 changes: 2 additions & 0 deletions compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUAttrs.td
View file
Original file line number Diff line number Diff line change
Expand Up @@ -289,6 +289,8 @@ def IREEGPU_TargetWgpAttr : AttrDef<IREEGPU_Dialect, "TargetWgp"> {
"uint32_t":$max_thread_count_per_workgroup,
// The maximal number of shared memory bytes we can allocate per workgroup.
"uint32_t":$max_workgroup_memory_bytes,
// Tthe maximum number of workgroups per X/Y/Z dimension in a dispatch.
"DenseI32ArrayAttr":$max_workgroup_counts,

// An optional extra dict
// This field allows to inject more features/limits not supported in the
Expand Down
12 changes: 8 additions & 4 deletions compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/test/target_attrs.mlir
View file
Original file line number Diff line number Diff line change
Expand Up @@ -11,15 +11,17 @@ func.func @test_target_wgp() attributes {
// CHECK-SAME: subgroup_size_choices = [32, 64],
// CHECK-SAME: max_workgroup_sizes = [1024, 1024, 1024],
// CHECK-SAME: max_thread_count_per_workgroup = 1024,
// CHECK-SAME: max_workgroup_memory_bytes = 65536>
// CHECK-SAME: max_workgroup_memory_bytes = 65536,
// CHECK-SAME: max_workgroup_counts = [2147483647, 2147483647, 2147483647]>
wgp = #iree_gpu.target_wgp<
compute = fp16|fp32|int8, storage = b16|b32,
subgroup = shuffle|arithmetic, dot = dp4xi8toi32,
mma = [<MFMA_F16_16x16x16_F32>, <MFMA_F16_32x32x8_F32>],
subgroup_size_choices = [32, 64],
max_workgroup_sizes = [1024, 1024, 1024],
max_thread_count_per_workgroup = 1024,
max_workgroup_memory_bytes = 65536
max_workgroup_memory_bytes = 65536,
max_workgroup_counts = [2147483647, 2147483647, 2147483647]
>
} { return }

Expand All @@ -37,7 +39,8 @@ func.func @test_target_wgp_none() attributes {
subgroup_size_choices = [32],
max_workgroup_sizes = [1024, 1024, 1024],
max_thread_count_per_workgroup = 1024,
max_workgroup_memory_bytes = 65536
max_workgroup_memory_bytes = 65536,
max_workgroup_counts = [2147483647, 2147483647, 2147483647]
>
} { return }

Expand Down Expand Up @@ -67,7 +70,8 @@ func.func @test_target() attributes {
subgroup_size_choices = [32, 64],
max_workgroup_sizes = [1024, 1024, 1024],
max_thread_count_per_workgroup = 1024,
max_workgroup_memory_bytes = 65536>,
max_workgroup_memory_bytes = 65536,
max_workgroup_counts = [2147483647, 2147483647, 2147483647]>,
chip = <wgp_count = 304>
>
} { return }
147 changes: 103 additions & 44 deletions compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -45,6 +45,7 @@ struct WgpDetails {
std::array<int32_t, 3> maxWorkgroupSizes;
uint32_t maxThreadSize;
uint32_t maxWorkgroupMemoryBytes;
std::array<int32_t, 3> maxWorkgroupCounts;
};

// Chip level feature/limit details
Expand Down Expand Up @@ -106,7 +107,9 @@ TargetAttr createTargetAttr(const TargetDetails &details, StringRef arch,
MMAOpsArrayAttr::get(context, mmaAttrs),
DenseI32ArrayAttr::get(context, subgroupSizes),
DenseI32ArrayAttr::get(context, wgp->maxWorkgroupSizes),
wgp->maxThreadSize, wgp->maxWorkgroupMemoryBytes, DictionaryAttr{});
wgp->maxThreadSize, wgp->maxWorkgroupMemoryBytes,
DenseI32ArrayAttr::get(context, wgp->maxWorkgroupCounts),
DictionaryAttr{});

TargetChipAttr targetChip;
if (details.chip)
Expand All @@ -118,6 +121,10 @@ TargetAttr createTargetAttr(const TargetDetails &details, StringRef arch,

//===----------------------------------------------------------------------===//
// Known AMD target details
//
// Note: the max workgroup size is given as signed int32 max because MLIR's
// `index` is signed and the workgroup ID is sign-extended, not zero-extended,
// to 64-bits.
//===----------------------------------------------------------------------===//

const WgpDetails *getCDNA3WgpDetails() {
Expand All @@ -129,11 +136,17 @@ const WgpDetails *getCDNA3WgpDetails() {
MMAIntrinsic::MFMA_I8_16x16x32_I32,
MMAIntrinsic::MFMA_I8_32x32x16_I32,
};
static const WgpDetails cdna3Wgp = {
allComputeBits, allStorageBits, allSubgroupOps,
allDotProductOps, ARRAY_SIZE(cdna3MMAOps), cdna3MMAOps,
{64, 64}, {1024, 1024, 1024}, 1024,
64 * 1024};
static const WgpDetails cdna3Wgp = {allComputeBits,
allStorageBits,
allSubgroupOps,
allDotProductOps,
ARRAY_SIZE(cdna3MMAOps),
cdna3MMAOps,
{64, 64},
{1024, 1024, 1024},
1024,
64 * 1024,
{0x7fffffff, 0x7fffffff, 0x7fffffff}};
return &cdna3Wgp;
}

Expand All @@ -142,11 +155,17 @@ const WgpDetails *getCDNA2WgpDetails() {
MMAIntrinsic::MFMA_F16_16x16x16_F32,
MMAIntrinsic::MFMA_F16_32x32x8_F32,
};
static const WgpDetails cdna2Wgp = {
allComputeBits, allStorageBits, allSubgroupOps,
allDotProductOps, ARRAY_SIZE(cdna2MMAOps), cdna2MMAOps,
{64, 64}, {1024, 1024, 1024}, 1024,
64 * 1024};
static const WgpDetails cdna2Wgp = {allComputeBits,
allStorageBits,
allSubgroupOps,
allDotProductOps,
ARRAY_SIZE(cdna2MMAOps),
cdna2MMAOps,
{64, 64},
{1024, 1024, 1024},
1024,
64 * 1024,
{0x7fffffff, 0x7fffffff, 0x7fffffff}};
return &cdna2Wgp;
}

Expand All @@ -155,11 +174,17 @@ const WgpDetails *getCDNA1WgpDetails() {
MMAIntrinsic::MFMA_F16_16x16x16_F32,
MMAIntrinsic::MFMA_F16_32x32x8_F32,
};
static const WgpDetails cdna1Wgp = {
allComputeBits, allStorageBits, allSubgroupOps,
allDotProductOps, ARRAY_SIZE(cdna1MMAOps), cdna1MMAOps,
{64, 64}, {1024, 1024, 1024}, 1024,
64 * 1024};
static const WgpDetails cdna1Wgp = {allComputeBits,
allStorageBits,
allSubgroupOps,
allDotProductOps,
ARRAY_SIZE(cdna1MMAOps),
cdna1MMAOps,
{64, 64},
{1024, 1024, 1024},
1024,
64 * 1024,
{0x7fffffff, 0x7fffffff, 0x7fffffff}};
return &cdna1Wgp;
}

Expand All @@ -168,27 +193,39 @@ const WgpDetails *getRDNA3WgpDetails() {
MMAIntrinsic::WMMA_F16_16x16x16_F32,
MMAIntrinsic::WMMA_F16_16x16x16_F16,
};
static const WgpDetails rdna3Wgp = {
allComputeBits, allStorageBits, allSubgroupOps,
allDotProductOps, ARRAY_SIZE(rdna3MMAOps), rdna3MMAOps,
{32, 64}, {1024, 1024, 1024}, 1024,
64 * 1024};
static const WgpDetails rdna3Wgp = {allComputeBits,
allStorageBits,
allSubgroupOps,
allDotProductOps,
ARRAY_SIZE(rdna3MMAOps),
rdna3MMAOps,
{32, 64},
{1024, 1024, 1024},
1024,
64 * 1024,
{0x7fffffff, 0x7fffffff, 0x7fffffff}};
return &rdna3Wgp;
}

const WgpDetails *getRDNA2WgpDetails() {
static const WgpDetails rdna2Wgp = {
allComputeBits, allStorageBits, allSubgroupOps, allDotProductOps,
/*mmaCount=*/0, /*mmaOps=*/nullptr, {32, 64}, {1024, 1024, 1024},
1024, 64 * 1024};
allComputeBits, allStorageBits,
allSubgroupOps, allDotProductOps,
/*mmaCount=*/0,
/*mmaOps=*/nullptr, {32, 64},
{1024, 1024, 1024}, 1024,
64 * 1024, {0x7fffffff, 0x7fffffff, 0x7fffffff}};
return &rdna2Wgp;
}

const WgpDetails *getRDNA1WgpDetails() {
static const WgpDetails rdna1Wgp = {
allComputeBits, allStorageBits, allSubgroupOps, DotProductOps::None,
/*mmaCount=*/0, /*mmaOps=*/nullptr, {32, 64}, {1024, 1024, 1024},
1024, 64 * 1024};
allComputeBits, allStorageBits,
allSubgroupOps, DotProductOps::None,
/*mmaCount=*/0,
/*mmaOps=*/nullptr, {32, 64},
{1024, 1024, 1024}, 1024,
64 * 1024, {0x7fffffff, 0x7fffffff, 0x7fffffff}};
return &rdna1Wgp;
}

Expand Down Expand Up @@ -281,7 +318,9 @@ std::optional<TargetDetails> getAppleTargetDetails() {
static const WgpDetails wgp = {
computeBitwdiths, allStorageBits, allSubgroupOps, allDotProductOps,
/*mmaCount=*/0, /*mmaOps=*/nullptr, {32, 32},
{1024, 1024, 1024}, 1024, 32 * 1024};
{1024, 1024, 1024}, 1024, 32 * 1024,
// Note: These values have not been checked and may be higher
{0xffff, 0xffff, 0xffff}};
// clang-format on

return TargetDetails{&wgp, nullptr};
Expand All @@ -302,7 +341,9 @@ const WgpDetails *getValhallWgpDetails() {
static const WgpDetails valhallWgp = {
computeBitwdiths, allStorageBits, allSubgroupOps, allDotProductOps,
/*mmaCount=*/0, /*mmaOps=*/nullptr, {16, 16}, {512, 512, 512},
512, 32 * 1024};
512, 32 * 1024,
// Note: These values have not been checked and may be higher
{0xffff, 0xffff, 0xffff}};
// clang-format on
return &valhallWgp;
}
Expand Down Expand Up @@ -358,11 +399,17 @@ const WgpDetails *getAmpereWgpDetails() {
MMAIntrinsic::WMMA_F16_16x16x16_F32,
MMAIntrinsic::WMMA_F16_16x16x16_F16,
};
static const WgpDetails ampereWgp = {
allComputeBits, allStorageBits, allSubgroupOps,
allDotProductOps, ARRAY_SIZE(mmaOps), mmaOps,
{32, 32}, {1024, 1024, 1024}, 1024,
163 * 1024};
static const WgpDetails ampereWgp = {allComputeBits,
allStorageBits,
allSubgroupOps,
allDotProductOps,
ARRAY_SIZE(mmaOps),
mmaOps,
{32, 32},
{1024, 1024, 1024},
1024,
163 * 1024,
{0x7fffffff, 0xffff, 0xffff}};
return &ampereWgp;
}

Expand All @@ -371,11 +418,17 @@ const WgpDetails *getTuringWgpDetails() {
MMAIntrinsic::WMMA_F16_16x16x16_F32,
MMAIntrinsic::WMMA_F16_16x16x16_F16,
};
static const WgpDetails turingWgp = {
allComputeBits, allStorageBits, allSubgroupOps,
allDotProductOps, ARRAY_SIZE(mmaOps), mmaOps,
{32, 32}, {1024, 1024, 1024}, 1024,
64 * 1024};
static const WgpDetails turingWgp = {allComputeBits,
allStorageBits,
allSubgroupOps,
allDotProductOps,
ARRAY_SIZE(mmaOps),
mmaOps,
{32, 32},
{1024, 1024, 1024},
1024,
64 * 1024,
{0x7fffffff, 0xffff, 0xffff}};
return &turingWgp;
}

Expand All @@ -388,7 +441,8 @@ const WgpDetails *getVoltaWgpDetails() {
static const WgpDetails voltaWgp = {
allComputeBits, allStorageBits, allSubgroupOps, DotProductOps::None,
ARRAY_SIZE(mmaOps), mmaOps, {32, 32}, {1024, 1024, 1024},
1024, 96 * 1024};
1024, 96 * 1024,
{0x7fffffff, 0xffff, 0xffff}};
// clang-format on
return &voltaWgp;
}
Expand All @@ -398,7 +452,8 @@ const WgpDetails *getPascalWgpDetails() {
static const WgpDetails pascalWgp = {
allComputeBits, allStorageBits, allSubgroupOps, DotProductOps::None,
0, nullptr, // Pascal does not have tensor core support.
{32, 32}, {1024, 1024, 1024}, 1024, 48 * 1024};
{32, 32}, {1024, 1024, 1024}, 1024, 48 * 1024,
{0x7fffffff, 0xffff, 0xffff}};
// clang-format on
return &pascalWgp;
}
Expand Down Expand Up @@ -479,7 +534,9 @@ const WgpDetails *getAdrenoWgpDetails() {
computeBitwdiths, storageBitwidths, allSubgroupOps,
allDotProductOps, /*mmaCount=*/0, /*mmaOps=*/nullptr,
{64, 64}, {1024, 1024, 1024}, 1024,
32 * 1024};
32 * 1024,
// Note: These values have not been checked and may be higher
{0xffff, 0xffff, 0xffff}};
// clang-format on
return &adrenoWgp;
}
Expand Down Expand Up @@ -545,7 +602,8 @@ const WgpDetails *getAndroidBaseline2022WgpDetails() {
computeBitwdiths, storageBitwidths, SubgroupOps::None,
DotProductOps::None, /*mmaCount=*/0, /*mmaOps=*/nullptr,
{64, 64}, {128, 128, 64}, 128,
16 * 1024};
16 * 1024,
{0xffff, 0xffff, 0xffff}};
// clang-format on
return &androidWgp;
}
Expand Down Expand Up @@ -645,7 +703,8 @@ TargetAttr getWebGPUTargetDetails(MLIRContext *context) {
computeBitwdiths, storageBitwidths, SubgroupOps::None,
DotProductOps::None, /*mmaCount=*/0, /*mmaOps=*/nullptr,
{32, 32}, {128, 128, 64}, 128,
16 * 1024};
16 * 1024,
{0xffff, 0xffff, 0xffff}};
// clang-format on

return createTargetAttr(
Expand Down
3 changes: 2 additions & 1 deletion compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/config_vector_distribute.mlir
View file
Original file line number Diff line number Diff line change
Expand Up @@ -93,7 +93,8 @@ func.func @conv_nhwc() {
subgroup = shuffle|arithmetic, dot = dp4xi8toi32,
mma = [],
subgroup_size_choices = [64], max_workgroup_sizes = [1024, 1024, 1024],
max_thread_count_per_workgroup = 1024, max_workgroup_memory_bytes = 65536>>
max_thread_count_per_workgroup = 1024, max_workgroup_memory_bytes = 65536,
max_workgroup_counts = [2147483647, 2147483647, 2147483647]>>
#executable_target_rocm_hsaco_fb = #hal.executable.target<"rocm", "rocm-hsaco-fb", {iree.gpu.target = #target}>
func.func @matmul_256x256x256() attributes {hal.executable.target = #executable_target_rocm_hsaco_fb} {
%cst = arith.constant 0.000000e+00 : f32
Expand Down
Loading

0 comments on commit b92a6ce

Please sign in to comment.