Reland "[spirv] Switch to use common target description" (#17699)
This relands https://github.com/iree-org/iree/pull/17623.
This commit switches SPIR-V side to use the common `#iree_gpu.target` to
describe the GPU characteristics. With it we can now remove the ad-hoc
Vulkan attributes and dialects and unify how GPU are described across
various GPU compiler backends in IREE.
SPIR-V has some additional requirements that we need to account for:
We have many vendors and APIs to handle there so this commit adds
various AMD/ARM/NVIDIA/Qualcomm targets for
development purposes so that we can specify them with a shorthand.
In order to be extensible, leverage the `feature` field in
`#iree_gpu.target` to specify additional capabilities with `cap:` prefix
and extensions with `ext:` prefix. We also use the `feature` field to
specify what SPIR-V version to target with the `spirv:v1.x` format.
Right now the `SPIRVConvertGPUTarget` pass is
invoked immediately before configuration. This is to stage the changes.
As a next step we need to move
it immediately before `ConvertToSPIRV` pass.
`--iree-vulkan-target-env` is dropped given now we removed the whole
Vulkan dialect and cannot control with a `#vk.target_env` attribute
anymore.
The default `--iree-vulkan-target-triple` now becomes
`vp_android_baseline_2022`, which is a a good lowest common denominator
to guarantee the generated SPIR-V is widely accepted. We are not
considering SwiftShader now anymore like previously due to testing
purposes.
The `--iree-vulkan-target-triple` should be renamed given it's not a
triple anymore--that will happen later together with other GPU backends
(i.e., cuda/hip) to be consistent.
In order to support cooperative matrix conversion, we added
`WMMA_F16_16x16x16_F16`. For NVIDIA GPUs
we are abusing it right now without considering the concrete explicit
layout--that is fine given in Vulkan they are opaque anyway. But this
need to be fixed if we are targeting WMMA in CUDA.
We now contruct a `#iree_gpu.target` to specify
the target to drive SPIR-V CodeGen.
Progress towards https://github.com/iree-org/iree/issues/16341
ci-extra: test_nvidia_gpu,test_nvidia_a100,test_amd_mi250,
build_test_all_macos_arm64,build_and_test_android,test_on_moto-edge-x30
---------
Signed-off-by: Lei Zhang <antiagainst@gmail.com>
diff --git a/samples/transform_dialect/example_module.mlir b/samples/transform_dialect/example_module.mlir
index 13128e1..585bb25 100644
--- a/samples/transform_dialect/example_module.mlir
+++ b/samples/transform_dialect/example_module.mlir
@@ -25,19 +25,21 @@
// }
// }
-#target_env = #spirv.target_env<#spirv.vce<v1.3, [Shader, GroupNonUniform], [SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers]>, api=Vulkan, #spirv.resource_limits<max_compute_workgroup_size = [128, 128, 64], subgroup_size = 64, cooperative_matrix_properties_khr = []>>
+#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 = [64, 64],
+ max_workgroup_sizes = [128, 128, 64], max_thread_count_per_workgroup = 128, max_workgroup_memory_bytes = 16384>>
module attributes {
hal.device.targets = [
#hal.device.target<"vulkan", [
#hal.executable.target<"vulkan-spirv", "vulkan-spirv-fb", {
- spirv.target_env = #target_env
+ iree.gpu.target = #target
}>
]>
]
} {
hal.executable private @example_module_dispatch_0 {
- hal.executable.variant public @vulkan_spirv_fb target(<"vulkan-spirv", "vulkan-spirv-fb", {spirv.target_env = #target_env}>) {
+ hal.executable.variant public @vulkan_spirv_fb target(<"vulkan-spirv", "vulkan-spirv-fb", {iree.gpu.target = #target}>) {
hal.executable.export public @example_module_dispatch_0_generic_80_f32 ordinal(0) layout(
#hal.pipeline.layout<push_constants = 0, sets = [<0, bindings = [<0, storage_buffer, ReadOnly>, <1, storage_buffer>]>]>) {
^bb0(%arg0: !hal.device):
@@ -63,7 +65,7 @@
}
}
hal.executable private @example_module_dispatch_1 {
- hal.executable.variant public @vulkan_spirv_fb target(<"vulkan-spirv", "vulkan-spirv-fb", {spirv.target_env = #target_env}>) {
+ hal.executable.variant public @vulkan_spirv_fb target(<"vulkan-spirv", "vulkan-spirv-fb", {iree.gpu.target = #target}>) {
hal.executable.export public @example_module_dispatch_1_matmul_16x16x5_f32 ordinal(0) layout(
#hal.pipeline.layout<push_constants = 0, sets = [<0, bindings = [<0, storage_buffer, ReadOnly>, <1, storage_buffer, ReadOnly>, <2, storage_buffer>]>]>) {
^bb0(%arg0: !hal.device):
@@ -87,7 +89,7 @@
}
}
hal.executable private @example_module_dispatch_2 {
- hal.executable.variant public @vulkan_spirv_fb target(<"vulkan-spirv", "vulkan-spirv-fb", {spirv.target_env = #target_env}>) {
+ hal.executable.variant public @vulkan_spirv_fb target(<"vulkan-spirv", "vulkan-spirv-fb", {iree.gpu.target = #target}>) {
hal.executable.export public @example_module_dispatch_2_generic_16x16_f32 ordinal(0) layout(
#hal.pipeline.layout<push_constants = 0, sets = [<0, bindings = [<0, storage_buffer, ReadOnly>, <1, storage_buffer>]>]>) {
^bb0(%arg0: !hal.device):
