From 7b9fb121ea1127a6d199904a0f2b0409573bc52d Mon Sep 17 00:00:00 2001 From: Lei Zhang Date: Tue, 18 Jun 2024 15:16:20 -0700 Subject: [PATCH] [spirv] Switch to use common target description (#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 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. #### Vulkan `--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. #### Metal / WebGPU 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 --- .../plugins/target/MetalSPIRV/BUILD.bazel | 1 + .../plugins/target/MetalSPIRV/CMakeLists.txt | 1 + .../target/MetalSPIRV/MetalSPIRVTarget.cpp | 67 +-- .../target/MetalSPIRV/test/smoketest.mlir | 4 +- .../ROCM/test/target_device_features.mlir | 2 +- .../plugins/target/VulkanSPIRV/BUILD.bazel | 3 +- .../plugins/target/VulkanSPIRV/CMakeLists.txt | 3 +- .../target/VulkanSPIRV/VulkanSPIRVTarget.cpp | 73 +-- .../target/VulkanSPIRV/test/smoketest.mlir | 4 +- .../plugins/target/WebGPUSPIRV/CMakeLists.txt | 1 + .../target/WebGPUSPIRV/WebGPUSPIRVTarget.cpp | 26 +- .../target/WebGPUSPIRV/test/smoketest.mlir | 4 +- .../API/Internal/IREEReduceToolEntryPoint.cpp | 15 +- .../Codegen/Dialect/GPU/IR/IREEGPUAttrs.cpp | 42 +- .../Codegen/Dialect/GPU/IR/IREEGPUEnums.td | 8 +- .../Dialect/GPU/TargetUtils/KnownTargets.cpp | 308 +++++++++- .../Dialect/GPU/TargetUtils/KnownTargets.h | 31 +- .../compiler/Codegen/LLVMGPU/KernelConfig.cpp | 22 +- .../iree/compiler/Codegen/SPIRV/BUILD.bazel | 1 + .../compiler/Codegen/SPIRV/CMakeLists.txt | 1 + .../iree/compiler/Codegen/SPIRV/Passes.cpp | 2 + .../src/iree/compiler/Codegen/SPIRV/Passes.h | 4 + .../src/iree/compiler/Codegen/SPIRV/Passes.td | 7 + .../Codegen/SPIRV/SPIRVConvertGPUTarget.cpp | 288 ++++++++++ .../compiler/Codegen/SPIRV/test/BUILD.bazel | 1 + .../Codegen/SPIRV/test/CMakeLists.txt | 1 + .../SPIRV/test/convert_gpu_target.mlir | 36 ++ .../iree/compiler/Dialect/Vulkan/BUILD.bazel | 11 - .../compiler/Dialect/Vulkan/CMakeLists.txt | 13 - .../compiler/Dialect/Vulkan/IR/BUILD.bazel | 87 --- .../compiler/Dialect/Vulkan/IR/CMakeLists.txt | 59 -- .../Dialect/Vulkan/IR/VulkanAttributes.cpp | 359 ------------ .../Dialect/Vulkan/IR/VulkanAttributes.h | 89 --- .../Dialect/Vulkan/IR/VulkanAttributes.td | 134 ----- .../compiler/Dialect/Vulkan/IR/VulkanBase.td | 199 ------- .../Dialect/Vulkan/IR/VulkanDialect.cpp | 18 - .../Dialect/Vulkan/IR/VulkanDialect.h | 37 -- .../Dialect/Vulkan/IR/VulkanTypes.cpp | 13 - .../compiler/Dialect/Vulkan/IR/VulkanTypes.h | 20 - .../Dialect/Vulkan/IR/test/BUILD.bazel | 26 - .../Dialect/Vulkan/IR/test/CMakeLists.txt | 23 - .../Dialect/Vulkan/IR/test/target_env.mlir | 150 ----- .../compiler/Dialect/Vulkan/Utils/BUILD.bazel | 32 -- .../Dialect/Vulkan/Utils/CMakeLists.txt | 31 - .../Vulkan/Utils/TargetEnvironment.cpp | 222 -------- .../Dialect/Vulkan/Utils/TargetEnvironment.h | 36 -- .../Dialect/Vulkan/Utils/TargetTriple.cpp | 539 ------------------ .../Dialect/Vulkan/Utils/TargetTriple.h | 67 --- .../Dialect/Vulkan/Utils/test/BUILD.bazel | 37 -- .../Dialect/Vulkan/Utils/test/CMakeLists.txt | 27 - .../Utils/test/target_env_conversion.mlir | 86 --- compiler/src/iree/compiler/Tools/BUILD.bazel | 1 - .../src/iree/compiler/Tools/CMakeLists.txt | 1 - .../iree/compiler/Tools/init_iree_dialects.h | 6 +- .../vulkan/shaders/example.mlir | 9 +- .../vulkan/shaders/example_inline.mlir | 9 +- .../vulkan/shaders/example_transform.mlir | 10 +- .../shaders/example_transform_spec.mlir | 10 +- samples/transform_dialect/example_module.mlir | 12 +- 59 files changed, 784 insertions(+), 2545 deletions(-) create mode 100644 compiler/src/iree/compiler/Codegen/SPIRV/SPIRVConvertGPUTarget.cpp create mode 100644 compiler/src/iree/compiler/Codegen/SPIRV/test/convert_gpu_target.mlir delete mode 100644 compiler/src/iree/compiler/Dialect/Vulkan/BUILD.bazel delete mode 100644 compiler/src/iree/compiler/Dialect/Vulkan/CMakeLists.txt delete mode 100644 compiler/src/iree/compiler/Dialect/Vulkan/IR/BUILD.bazel delete mode 100644 compiler/src/iree/compiler/Dialect/Vulkan/IR/CMakeLists.txt delete mode 100644 compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanAttributes.cpp delete mode 100644 compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanAttributes.h delete mode 100644 compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanAttributes.td delete mode 100644 compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanBase.td delete mode 100644 compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanDialect.cpp delete mode 100644 compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanDialect.h delete mode 100644 compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanTypes.cpp delete mode 100644 compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanTypes.h delete mode 100644 compiler/src/iree/compiler/Dialect/Vulkan/IR/test/BUILD.bazel delete mode 100644 compiler/src/iree/compiler/Dialect/Vulkan/IR/test/CMakeLists.txt delete mode 100644 compiler/src/iree/compiler/Dialect/Vulkan/IR/test/target_env.mlir delete mode 100644 compiler/src/iree/compiler/Dialect/Vulkan/Utils/BUILD.bazel delete mode 100644 compiler/src/iree/compiler/Dialect/Vulkan/Utils/CMakeLists.txt delete mode 100644 compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetEnvironment.cpp delete mode 100644 compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetEnvironment.h delete mode 100644 compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetTriple.cpp delete mode 100644 compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetTriple.h delete mode 100644 compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/BUILD.bazel delete mode 100644 compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/CMakeLists.txt delete mode 100644 compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/target_env_conversion.mlir diff --git a/compiler/plugins/target/MetalSPIRV/BUILD.bazel b/compiler/plugins/target/MetalSPIRV/BUILD.bazel index 9773eff4786d..ede556649d07 100644 --- a/compiler/plugins/target/MetalSPIRV/BUILD.bazel +++ b/compiler/plugins/target/MetalSPIRV/BUILD.bazel @@ -26,6 +26,7 @@ iree_compiler_cc_library( ":SPIRVToMSL", "//compiler/src/iree/compiler/Codegen/Common", "//compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR:IREECodegenDialect", + "//compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils:KnownTargets", "//compiler/src/iree/compiler/Codegen/SPIRV", "//compiler/src/iree/compiler/Codegen/Utils", "//compiler/src/iree/compiler/Dialect/Flow/IR", diff --git a/compiler/plugins/target/MetalSPIRV/CMakeLists.txt b/compiler/plugins/target/MetalSPIRV/CMakeLists.txt index 4dd1b0614e3a..678a37a1a4d7 100644 --- a/compiler/plugins/target/MetalSPIRV/CMakeLists.txt +++ b/compiler/plugins/target/MetalSPIRV/CMakeLists.txt @@ -36,6 +36,7 @@ iree_cc_library( MLIRVectorDialect iree::compiler::Codegen::Common iree::compiler::Codegen::Dialect::Codegen::IR::IREECodegenDialect + iree::compiler::Codegen::Dialect::GPU::TargetUtils::KnownTargets iree::compiler::Codegen::SPIRV iree::compiler::Codegen::Utils iree::compiler::Dialect::Flow::IR diff --git a/compiler/plugins/target/MetalSPIRV/MetalSPIRVTarget.cpp b/compiler/plugins/target/MetalSPIRV/MetalSPIRVTarget.cpp index 6ea1cbf7719f..25e8e51843c9 100644 --- a/compiler/plugins/target/MetalSPIRV/MetalSPIRVTarget.cpp +++ b/compiler/plugins/target/MetalSPIRV/MetalSPIRVTarget.cpp @@ -8,6 +8,7 @@ #include "compiler/plugins/target/MetalSPIRV/MetalTargetPlatform.h" #include "compiler/plugins/target/MetalSPIRV/SPIRVToMSL.h" #include "iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenDialect.h" +#include "iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.h" #include "iree/compiler/Codegen/SPIRV/Passes.h" #include "iree/compiler/Dialect/Flow/IR/FlowDialect.h" #include "iree/compiler/Dialect/HAL/Target/TargetRegistry.h" @@ -19,9 +20,7 @@ #include "llvm/TargetParser/Triple.h" #include "mlir/Dialect/GPU/IR/GPUDialect.h" #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" -#include "mlir/Dialect/SPIRV/IR/SPIRVEnums.h" #include "mlir/Dialect/SPIRV/IR/SPIRVOps.h" -#include "mlir/Dialect/SPIRV/IR/TargetAndABI.h" #include "mlir/Dialect/Vector/IR/VectorOps.h" #include "mlir/Target/SPIRV/Serialization.h" @@ -52,60 +51,6 @@ struct MetalSPIRVOptions { }; } // namespace -static spirv::TargetEnvAttr getMetalTargetEnv(MLIRContext *context) { - using spirv::Capability; - using spirv::Extension; - - // Capabilities and limits according to Metal 3 devices. - const std::array extensions = { - Extension::SPV_KHR_16bit_storage, - Extension::SPV_KHR_8bit_storage, - Extension::SPV_KHR_storage_buffer_storage_class, - Extension::SPV_KHR_variable_pointers, - }; - const std::array capabilities = { - Capability::Shader, - Capability::Int8, - Capability::Int16, - Capability::Int64, - Capability::Float16, - Capability::UniformAndStorageBuffer8BitAccess, - Capability::StorageBuffer8BitAccess, - Capability::StoragePushConstant8, - Capability::StorageUniform16, - Capability::StorageBuffer16BitAccess, - Capability::StoragePushConstant16, - Capability::GroupNonUniform, - Capability::GroupNonUniformVote, - Capability::GroupNonUniformArithmetic, - Capability::GroupNonUniformBallot, - Capability::GroupNonUniformShuffle, - Capability::GroupNonUniformShuffleRelative, - Capability::GroupNonUniformQuad, - Capability::StoragePushConstant16, - Capability::VariablePointers, - Capability::VariablePointersStorageBuffer, - }; - auto limits = spirv::ResourceLimitsAttr::get( - context, - /*max_compute_shared_memory_size=*/32768, - /*max_compute_workgroup_invocations=*/1024, - /*max_compute_workgroup_size=*/ - Builder(context).getI32ArrayAttr({1024, 1024, 1024}), - /*subgroup_size=*/32, - /*min_subgroup_size=*/std::nullopt, - /*max_subgroup_size=*/std::nullopt, - /*cooperative_matrix_properties_khr=*/ArrayAttr{}, - /*cooperative_matrix_properties_nv=*/ArrayAttr{}); - - auto triple = spirv::VerCapExtAttr::get(spirv::Version::V_1_3, capabilities, - extensions, context); - // Further assuming Apple GPUs. - return spirv::TargetEnvAttr::get( - triple, limits, spirv::ClientAPI::Metal, spirv::Vendor::Apple, - spirv::DeviceType::IntegratedGPU, spirv::TargetEnvAttr::kUnknownDeviceID); -} - // TODO: MetalOptions for choosing the Metal version. class MetalTargetDevice : public TargetDevice { public: @@ -145,20 +90,20 @@ class MetalSPIRVTargetBackend : public TargetBackend { MLIRContext *context, StringRef deviceID, DictionaryAttr deviceConfigAttr, SmallVectorImpl &executableTargetAttrs) const override { - executableTargetAttrs.push_back( - getExecutableTarget(context, getMetalTargetEnv(context))); + executableTargetAttrs.push_back(getExecutableTarget(context)); } IREE::HAL::ExecutableTargetAttr - getExecutableTarget(MLIRContext *context, - spirv::TargetEnvAttr targetEnv) const { + getExecutableTarget(MLIRContext *context) const { Builder b(context); SmallVector configItems; auto addConfig = [&](StringRef name, Attribute value) { configItems.emplace_back(b.getStringAttr(name), value); }; - addConfig(spirv::getTargetEnvAttrName(), targetEnv); + if (auto target = GPU::getMetalTargetDetails(context)) { + addConfig("iree.gpu.target", target); + } return b.getAttr( b.getStringAttr("metal-spirv"), b.getStringAttr("metal-msl-fb"), diff --git a/compiler/plugins/target/MetalSPIRV/test/smoketest.mlir b/compiler/plugins/target/MetalSPIRV/test/smoketest.mlir index 84dc61e90945..720e00b2f835 100644 --- a/compiler/plugins/target/MetalSPIRV/test/smoketest.mlir +++ b/compiler/plugins/target/MetalSPIRV/test/smoketest.mlir @@ -4,7 +4,9 @@ module attributes { hal.device.targets = [ #hal.device.target<"metal", [ #hal.executable.target<"metal-spirv", "metal-msl-fb", { - spirv.target_env = #spirv.target_env<#spirv.vce, #spirv.resource_limits<>> + iree.gpu.target = #iree_gpu.target> }> ]> ] diff --git a/compiler/plugins/target/ROCM/test/target_device_features.mlir b/compiler/plugins/target/ROCM/test/target_device_features.mlir index 9f0124651c65..0f427c52fd00 100644 --- a/compiler/plugins/target/ROCM/test/target_device_features.mlir +++ b/compiler/plugins/target/ROCM/test/target_device_features.mlir @@ -15,7 +15,7 @@ // GFX940-SAME: mma = [, ] // GFX1100: target = #iree_gpu.target] +// GFX1100-SAME: mma = [, ] // GFX1100-SAME: subgroup_size_choices = [32, 64] // GFX941: target = #iree_gpu.target( + // TODO: Rename this as target given it's not a triple anymore. "iree-vulkan-target-triple", targetTriple, llvm::cl::desc( "Vulkan target triple controlling the SPIR-V environment.")); - binder.opt( - "iree-vulkan-target-env", targetEnv, - llvm::cl::desc( - "Vulkan target environment as #vk.target_env attribute assembly.")); binder.opt( "iree-vulkan-experimental-indirect-bindings", indirectBindings, llvm::cl::desc( @@ -56,31 +53,6 @@ struct VulkanSPIRVTargetOptions { }; } // namespace -// Returns the Vulkan target environment for conversion. -static spirv::TargetEnvAttr -getSPIRVTargetEnv(const std::string &vulkanTargetTripleOrEnv, - MLIRContext *context) { - if (!vulkanTargetTripleOrEnv.empty()) { - if (vulkanTargetTripleOrEnv[0] != '#') { - // Parse target triple. - return convertTargetEnv( - Vulkan::getTargetEnvForTriple(context, vulkanTargetTripleOrEnv)); - } - - // Parse `#vk.target_env<...` attribute assembly. - if (auto attr = parseAttribute(vulkanTargetTripleOrEnv, context)) { - if (auto vkTargetEnv = llvm::dyn_cast(attr)) { - return convertTargetEnv(vkTargetEnv); - } - } - emitError(Builder(context).getUnknownLoc()) - << "cannot parse vulkan target environment as #vk.target_env " - "attribute: '" - << vulkanTargetTripleOrEnv << "'"; - } - return {}; -} - // TODO: VulkanOptions for choosing the Vulkan version and extensions/features. class VulkanTargetDevice : public TargetDevice { public: @@ -119,35 +91,32 @@ class VulkanSPIRVTargetBackend : public TargetBackend { MLIRContext *context, StringRef deviceID, DictionaryAttr deviceConfigAttr, SmallVectorImpl &executableTargetAttrs) const override { - std::string targetTripleOrEnv; - if (!options_.targetEnv.empty()) { - // TODO(scotttodd): assert if triple is set too? (mutually exclusive) - targetTripleOrEnv = options_.targetEnv; - } else if (!options_.targetTriple.empty()) { - targetTripleOrEnv = options_.targetTriple; - } else { - targetTripleOrEnv = "unknown-unknown-unknown"; - } - - executableTargetAttrs.push_back(getExecutableTarget( - context, getSPIRVTargetEnv(targetTripleOrEnv, context), - options_.indirectBindings)); + executableTargetAttrs.push_back( + getExecutableTarget(context, options_.indirectBindings)); } IREE::HAL::ExecutableTargetAttr - getExecutableTarget(MLIRContext *context, spirv::TargetEnvAttr targetEnv, - bool indirectBindings) const { + getExecutableTarget(MLIRContext *context, bool indirectBindings) const { Builder b(context); SmallVector configItems; auto addConfig = [&](StringRef name, Attribute value) { configItems.emplace_back(b.getStringAttr(name), value); }; - addConfig(spirv::getTargetEnvAttrName(), targetEnv); if (indirectBindings) { addConfig("hal.bindings.indirect", b.getUnitAttr()); } + // We only care about the architecture right now. + StringRef arch = StringRef(options_.targetTriple).split("-").first; + if (auto target = GPU::getVulkanTargetDetails(arch, context)) { + addConfig("iree.gpu.target", target); + } else { + emitError(b.getUnknownLoc(), "Unknown Vulkan target '") + << options_.targetTriple << "'"; + return nullptr; + } + return IREE::HAL::ExecutableTargetAttr::get( context, b.getStringAttr("vulkan-spirv"), indirectBindings ? b.getStringAttr("vulkan-spirv-fb-ptr") @@ -156,8 +125,8 @@ class VulkanSPIRVTargetBackend : public TargetBackend { } void getDependentDialects(DialectRegistry ®istry) const override { - registry.insert(); + registry.insert(); } void diff --git a/compiler/plugins/target/VulkanSPIRV/test/smoketest.mlir b/compiler/plugins/target/VulkanSPIRV/test/smoketest.mlir index 68d654297a6a..f8d81592b778 100644 --- a/compiler/plugins/target/VulkanSPIRV/test/smoketest.mlir +++ b/compiler/plugins/target/VulkanSPIRV/test/smoketest.mlir @@ -4,7 +4,9 @@ module attributes { hal.device.targets = [ #hal.device.target<"vulkan", [ #hal.executable.target<"vulkan-spirv", "vulkan-spirv-fb", { - spirv.target_env = #spirv.target_env<#spirv.vce, #spirv.resource_limits<>> + iree.gpu.target = #iree_gpu.target> }> ]> ] diff --git a/compiler/plugins/target/WebGPUSPIRV/CMakeLists.txt b/compiler/plugins/target/WebGPUSPIRV/CMakeLists.txt index d98dcf261979..caf4460d1d51 100644 --- a/compiler/plugins/target/WebGPUSPIRV/CMakeLists.txt +++ b/compiler/plugins/target/WebGPUSPIRV/CMakeLists.txt @@ -48,6 +48,7 @@ iree_cc_library( MLIRSPIRVTransforms SPIRV-Tools iree::compiler::Codegen::Dialect::Codegen::IR::IREECodegenDialect + iree::compiler::Codegen::Dialect::GPU::TargetUtils::KnownTargets iree::compiler::Codegen::SPIRV iree::compiler::Dialect::Flow::IR iree::compiler::Dialect::HAL::Target diff --git a/compiler/plugins/target/WebGPUSPIRV/WebGPUSPIRVTarget.cpp b/compiler/plugins/target/WebGPUSPIRV/WebGPUSPIRVTarget.cpp index 8397eb10e4f1..0a376912652a 100644 --- a/compiler/plugins/target/WebGPUSPIRV/WebGPUSPIRVTarget.cpp +++ b/compiler/plugins/target/WebGPUSPIRV/WebGPUSPIRVTarget.cpp @@ -6,18 +6,16 @@ #include "compiler/plugins/target/WebGPUSPIRV/SPIRVToWGSL.h" #include "iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenDialect.h" +#include "iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.h" #include "iree/compiler/Codegen/SPIRV/Passes.h" #include "iree/compiler/Codegen/WGSL/Passes.h" #include "iree/compiler/Dialect/Flow/IR/FlowDialect.h" #include "iree/compiler/Dialect/HAL/Target/TargetRegistry.h" -#include "iree/compiler/Dialect/HAL/Transforms/Passes.h" #include "iree/compiler/PluginAPI/Client.h" #include "iree/compiler/Utils/FlatbufferUtils.h" #include "iree/schemas/wgsl_executable_def_builder.h" #include "llvm/Support/CommandLine.h" -#include "llvm/Support/FileSystem.h" #include "llvm/Support/FormatVariadic.h" -#include "llvm/Support/ToolOutputFile.h" #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/GPU/IR/GPUDialect.h" #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" @@ -43,18 +41,6 @@ struct WebGPUSPIRVOptions { } }; -// TODO(scotttodd): provide a proper target environment for WebGPU. -static spirv::TargetEnvAttr getWebGPUTargetEnv(MLIRContext *context) { - // TODO(scotttodd): find list of SPIR-V extensions supported by WebGPU/WGSL - auto triple = spirv::VerCapExtAttr::get( - spirv::Version::V_1_0, {spirv::Capability::Shader}, - {spirv::Extension::SPV_KHR_storage_buffer_storage_class}, context); - return spirv::TargetEnvAttr::get( - triple, spirv::getDefaultResourceLimits(context), - spirv::ClientAPI::WebGPU, spirv::Vendor::Unknown, - spirv::DeviceType::Unknown, spirv::TargetEnvAttr::kUnknownDeviceID); -} - // TODO: WebGPUOptions for choosing the version/extensions/etc. class WebGPUTargetDevice : public TargetDevice { public: @@ -94,20 +80,20 @@ class WebGPUSPIRVTargetBackend : public TargetBackend { MLIRContext *context, StringRef deviceID, DictionaryAttr deviceConfigAttr, SmallVectorImpl &executableTargetAttrs) const override { - executableTargetAttrs.push_back( - getExecutableTarget(context, getWebGPUTargetEnv(context))); + executableTargetAttrs.push_back(getExecutableTarget(context)); } IREE::HAL::ExecutableTargetAttr - getExecutableTarget(MLIRContext *context, - spirv::TargetEnvAttr targetEnv) const { + getExecutableTarget(MLIRContext *context) const { Builder b(context); SmallVector configItems; auto addConfig = [&](StringRef name, Attribute value) { configItems.emplace_back(b.getStringAttr(name), value); }; - addConfig(spirv::getTargetEnvAttrName(), targetEnv); + if (auto target = GPU::getWebGPUTargetDetails(context)) { + addConfig("iree.gpu.target", target); + } return b.getAttr( b.getStringAttr("webgpu-spirv"), b.getStringAttr("webgpu-wgsl-fb"), diff --git a/compiler/plugins/target/WebGPUSPIRV/test/smoketest.mlir b/compiler/plugins/target/WebGPUSPIRV/test/smoketest.mlir index 1a17240ac6bc..31f361b1ab5f 100644 --- a/compiler/plugins/target/WebGPUSPIRV/test/smoketest.mlir +++ b/compiler/plugins/target/WebGPUSPIRV/test/smoketest.mlir @@ -5,7 +5,9 @@ module attributes { hal.device.targets = [ #hal.device.target<"webgpu", [ #hal.executable.target<"webgpu-spirv", "webgpu-wgsl-fb", { - spirv.target_env = #spirv.target_env<#spirv.vce, #spirv.resource_limits<>> + iree.gpu.target = #iree_gpu.target> }> ]> ] diff --git a/compiler/src/iree/compiler/API/Internal/IREEReduceToolEntryPoint.cpp b/compiler/src/iree/compiler/API/Internal/IREEReduceToolEntryPoint.cpp index 7e0d2014b268..4cb61ec3c421 100644 --- a/compiler/src/iree/compiler/API/Internal/IREEReduceToolEntryPoint.cpp +++ b/compiler/src/iree/compiler/API/Internal/IREEReduceToolEntryPoint.cpp @@ -8,13 +8,13 @@ #include "iree/compiler/tool_entry_points_api.h" #include "iree/compiler/Tools/init_dialects.h" +#include "llvm/Support/CommandLine.h" #include "llvm/Support/InitLLVM.h" #include "llvm/Support/Process.h" #include "llvm/Support/SourceMgr.h" #include "llvm/Support/ToolOutputFile.h" #include "mlir/Bytecode/BytecodeWriter.h" #include "mlir/IR/AsmState.h" -#include "mlir/IR/Dialect.h" #include "mlir/IR/MLIRContext.h" #include "mlir/Parser/Parser.h" #include "mlir/Pass/PassManager.h" @@ -50,11 +50,11 @@ static OwningOpRef loadModule(MLIRContext &context, static LogicalResult ireeReduceMainFromCL(int argc, char **argv, MLIRContext ®istry) { - llvm::cl::OptionCategory ireeReduceCategory("iree-reduce options"); + cl::OptionCategory ireeReduceCategory("iree-reduce options"); - llvm::cl::opt testScript(cl::Positional, cl::Required, - cl::desc(""), - cl::cat(ireeReduceCategory)); + cl::opt testScript(cl::Positional, cl::Required, + cl::desc(""), + cl::cat(ireeReduceCategory)); cl::opt inputFilename(cl::Positional, cl::desc(""), cl::init("-"), @@ -74,12 +74,11 @@ static LogicalResult ireeReduceMainFromCL(int argc, char **argv, "output-bytecode", cl::desc("Output the final output as bytecode."), cl::init(false), llvm::cl::cat(ireeReduceCategory)); - llvm::cl::HideUnrelatedOptions(ireeReduceCategory); + cl::HideUnrelatedOptions(ireeReduceCategory); InitLLVM y(argc, argv); - llvm::cl::ParseCommandLineOptions(argc, argv, - "IREE test case reduction tool.\n"); + cl::ParseCommandLineOptions(argc, argv, "IREE test case reduction tool.\n"); // When reading from stdin and the input is a tty, it is often a user mistake // and the process "appears to be stuck". Print a message to let the user know diff --git a/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUAttrs.cpp b/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUAttrs.cpp index 5df8a3632844..d88dc84389cd 100644 --- a/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUAttrs.cpp +++ b/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUAttrs.cpp @@ -8,7 +8,6 @@ #include #include "iree-dialects/Dialect/VectorExt/IR/VectorExtDialect.h" -#include "iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenInterfaces.h" #include "iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUDialect.h" #include "iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUEnums.h" #include "iree/compiler/Codegen/Dialect/GPU/TargetUtils/ConfigUtils.h" @@ -17,7 +16,6 @@ #include "llvm/ADT/STLForwardCompat.h" #include "llvm/ADT/SmallVector.h" #include "llvm/ADT/TypeSwitch.h" -#include "llvm/ADT/iterator_range.h" #include "llvm/Support/Debug.h" #include "llvm/Support/ErrorHandling.h" #include "mlir/Dialect/AMDGPU/IR/AMDGPUDialect.h" @@ -216,6 +214,9 @@ static OpaqueMmaLayout getOpaqueMFMALayout(MLIRContext *context, case MMAIntrinsic::WMMA_F16_16x16x16_F32: { return OpaqueMmaLayout{16, 16, 16, f16, f16, f32}; } + case MMAIntrinsic::WMMA_F16_16x16x16_F16: { + return OpaqueMmaLayout{16, 16, 16, f16, f16, f16}; + } } llvm_unreachable("unhandled mfma layout type"); return OpaqueMmaLayout{}; @@ -278,7 +279,8 @@ static ConcreteMmaLayout getConcreteMFMALayout(MLIRContext *context, return ConcreteMmaLayout{opaqueLayout, aMLayout, aKLayout, bKLayout, bNLayout, cMLayout, cNLayout}; } - case MMAIntrinsic::WMMA_F16_16x16x16_F32: { + case MMAIntrinsic::WMMA_F16_16x16x16_F32: + case MMAIntrinsic::WMMA_F16_16x16x16_F16: { // #outer = #iree_vector_ext.per_dim_layout<[LANEX], [16]> // #inner = #iree_vector_ext.per_dim_layout<[LANEY, VECTORX], [4, 4]> // #layout_a = #iree_vector_ext.layout<#outer, #inner> @@ -369,7 +371,8 @@ MMAAttr::getABCVectorTypes() const { auto cType = VectorType::get({16}, getCType()); return std::make_tuple(aType, bType, cType); } - case MMAIntrinsic::WMMA_F16_16x16x16_F32: { + case MMAIntrinsic::WMMA_F16_16x16x16_F32: + case MMAIntrinsic::WMMA_F16_16x16x16_F16: { auto aType = VectorType::get({16}, getAType()); auto bType = VectorType::get({16}, getBType()); auto cType = VectorType::get({8}, getCType()); @@ -392,6 +395,7 @@ int64_t MMAAttr::getBlockSize() const { switch (getIntrinsic().getValue()) { case MMAIntrinsic::MFMA_F16_16x16x16_F32: case MMAIntrinsic::MFMA_F16_32x32x8_F32: + case MMAIntrinsic::WMMA_F16_16x16x16_F16: case MMAIntrinsic::WMMA_F16_16x16x16_F32: { return 1; } @@ -406,7 +410,8 @@ int64_t MMAAttr::getSubgroupSize() const { case MMAIntrinsic::MFMA_F16_32x32x8_F32: { return 64; } - case MMAIntrinsic::WMMA_F16_16x16x16_F32: { + case MMAIntrinsic::WMMA_F16_16x16x16_F32: + case MMAIntrinsic::WMMA_F16_16x16x16_F16: { return 32; } } @@ -420,7 +425,8 @@ SmallVector MMAAttr::getADataDuplicate() const { case MMAIntrinsic::MFMA_F16_32x32x8_F32: { break; } - case MMAIntrinsic::WMMA_F16_16x16x16_F32: { + case MMAIntrinsic::WMMA_F16_16x16x16_F32: + case MMAIntrinsic::WMMA_F16_16x16x16_F16: { return {2, 1}; } } @@ -434,7 +440,8 @@ SmallVector MMAAttr::getBDataDuplicate() const { case MMAIntrinsic::MFMA_F16_32x32x8_F32: { break; } - case MMAIntrinsic::WMMA_F16_16x16x16_F32: { + case MMAIntrinsic::WMMA_F16_16x16x16_F32: + case MMAIntrinsic::WMMA_F16_16x16x16_F16: { return {1, 2}; } } @@ -455,7 +462,8 @@ MMAAttr::SingleSubgroupLayout MMAAttr::getASingleSubgroupLayoutCount() const { case MMAIntrinsic::MFMA_F16_32x32x8_F32: { return {/*outer=*/{1, 1}, /*thread=*/{32, 2}, /*element=*/{1, 4}}; } - case MMAIntrinsic::WMMA_F16_16x16x16_F32: { + case MMAIntrinsic::WMMA_F16_16x16x16_F32: + case MMAIntrinsic::WMMA_F16_16x16x16_F16: { return {/*outer=*/{1, 1}, /*thread=*/{16, 1}, /*element=*/{1, 16}}; } } @@ -470,7 +478,8 @@ MMAAttr::SingleSubgroupLayout MMAAttr::getBSingleSubgroupLayoutCount() const { case MMAIntrinsic::MFMA_F16_32x32x8_F32: { return {/*outer=*/{1, 1}, /*thread=*/{2, 32}, /*element=*/{4, 1}}; } - case MMAIntrinsic::WMMA_F16_16x16x16_F32: { + case MMAIntrinsic::WMMA_F16_16x16x16_F32: + case MMAIntrinsic::WMMA_F16_16x16x16_F16: { return {/*outer=*/{1, 1}, /*thread=*/{1, 16}, /*element=*/{16, 1}}; } } @@ -485,7 +494,8 @@ MMAAttr::SingleSubgroupLayout MMAAttr::getCSingleSubgroupLayoutCount() const { case MMAIntrinsic::MFMA_F16_32x32x8_F32: { return {/*outer=*/{4, 1}, /*thread=*/{2, 32}, /*element=*/{4, 1}}; } - case MMAIntrinsic::WMMA_F16_16x16x16_F32: { + case MMAIntrinsic::WMMA_F16_16x16x16_F32: + case MMAIntrinsic::WMMA_F16_16x16x16_F16: { return {/*outer=*/{8, 1}, /*thread=*/{2, 16}, /*element=*/{1, 1}}; } } @@ -496,7 +506,8 @@ MMAAttr::SingleSubgroupLayout MMAAttr::getASingleSubgroupLayoutOrder() const { switch (getIntrinsic().getValue()) { case MMAIntrinsic::MFMA_F16_16x16x16_F32: case MMAIntrinsic::MFMA_F16_32x32x8_F32: - case MMAIntrinsic::WMMA_F16_16x16x16_F32: { + case MMAIntrinsic::WMMA_F16_16x16x16_F32: + case MMAIntrinsic::WMMA_F16_16x16x16_F16: { return {/*outer=*/{0, 1}, /*thread=*/{1, 0}, /*element=*/{0, 1}}; } } @@ -507,7 +518,8 @@ MMAAttr::SingleSubgroupLayout MMAAttr::getBSingleSubgroupLayoutOrder() const { switch (getIntrinsic().getValue()) { case MMAIntrinsic::MFMA_F16_16x16x16_F32: case MMAIntrinsic::MFMA_F16_32x32x8_F32: - case MMAIntrinsic::WMMA_F16_16x16x16_F32: { + case MMAIntrinsic::WMMA_F16_16x16x16_F32: + case MMAIntrinsic::WMMA_F16_16x16x16_F16: { return {/*outer=*/{0, 1}, /*thread=*/{0, 1}, /*element=*/{1, 0}}; } } @@ -518,7 +530,8 @@ MMAAttr::SingleSubgroupLayout MMAAttr::getCSingleSubgroupLayoutOrder() const { switch (getIntrinsic().getValue()) { case MMAIntrinsic::MFMA_F16_16x16x16_F32: case MMAIntrinsic::MFMA_F16_32x32x8_F32: - case MMAIntrinsic::WMMA_F16_16x16x16_F32: { + case MMAIntrinsic::WMMA_F16_16x16x16_F32: + case MMAIntrinsic::WMMA_F16_16x16x16_F16: { return {/*outer=*/{0, 1}, /*thread=*/{0, 1}, /*element=*/{1, 0}}; } } @@ -549,7 +562,8 @@ FailureOr MMAAttr::buildMmaOperation(OpBuilder &builder, Location loc, rhs, acc) .getResult(); } - case MMAIntrinsic::WMMA_F16_16x16x16_F32: { + case MMAIntrinsic::WMMA_F16_16x16x16_F32: + case MMAIntrinsic::WMMA_F16_16x16x16_F16: { return builder.create(loc, resultType, lhs, rhs, acc) .getResult(); } diff --git a/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUEnums.td b/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUEnums.td index 5c1bead3a3a3..a7abbb65ceb6 100644 --- a/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUEnums.td +++ b/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUEnums.td @@ -98,15 +98,19 @@ class IREEGPU_I32MmaEnumAttr let genSpecializedAttr = 0; } +// Format: __xx_ def MFMA_F16_16x16x16_F32 : I32EnumAttrCase<"MFMA_F16_16x16x16_F32", 0>; -def MFMA_F16_32x32x8_F32 : I32EnumAttrCase<"MFMA_F16_32x32x8_F32", 1>; +def MFMA_F16_32x32x8_F32 : I32EnumAttrCase<"MFMA_F16_32x32x8_F32", 1>; +// TODO: Create separate WMMA ops for AMD and NVIDIA GPUs def WMMA_F16_16x16x16_F32 : I32EnumAttrCase<"WMMA_F16_16x16x16_F32", 2>; +def WMMA_F16_16x16x16_F16 : I32EnumAttrCase<"WMMA_F16_16x16x16_F16", 3>; def IREEGPU_MMAIntrinsic : IREEGPU_I32MmaEnumAttr<"MMAIntrinsic", "Descriptor for different MMA intrinsics", [ MFMA_F16_16x16x16_F32, MFMA_F16_32x32x8_F32, - WMMA_F16_16x16x16_F32 + WMMA_F16_16x16x16_F32, + WMMA_F16_16x16x16_F16 ]>; def MMA_LHS : I32EnumAttrCase<"Lhs", 0>; diff --git a/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.cpp b/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.cpp index e02e759abbde..bdfcace37f3b 100644 --- a/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.cpp +++ b/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.cpp @@ -8,6 +8,7 @@ #include #include "iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUAttrs.h" +#include "iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUEnums.h" #include "llvm/ADT/StringSwitch.h" #include "mlir/IR/BuiltinAttributes.h" @@ -156,6 +157,7 @@ const WgpDetails *getCDNA1WgpDetails() { const WgpDetails *getRDNA3WgpDetails() { static const MMAIntrinsic rdna3MMAOps[] = { MMAIntrinsic::WMMA_F16_16x16x16_F32, + MMAIntrinsic::WMMA_F16_16x16x16_F16, }; static const WgpDetails rdna3Wgp = { allComputeBits, allStorageBits, allSubgroupOps, @@ -165,11 +167,29 @@ const WgpDetails *getRDNA3WgpDetails() { 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}; + 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}; + return &rdna1Wgp; +} + std::optional getAMDGPUTargetDetails(StringRef target) { const WgpDetails *cdna3Wgp = getCDNA3WgpDetails(); const WgpDetails *cdna2Wgp = getCDNA2WgpDetails(); const WgpDetails *cdna1Wgp = getCDNA1WgpDetails(); const WgpDetails *rdna3Wgp = getRDNA3WgpDetails(); + const WgpDetails *rdna2Wgp = getRDNA2WgpDetails(); + const WgpDetails *rdna1Wgp = getRDNA1WgpDetails(); // "AMD Instinct MI300 Series Product Offerings" in Page 23 of // https://www.amd.com/content/dam/amd/en/documents/instinct-tech-docs/white-papers/amd-cdna-3-white-paper.pdf @@ -215,6 +235,10 @@ std::optional getAMDGPUTargetDetails(StringRef target) { .Case("rx7700xt", TargetDetails{rdna3Wgp, &rx7700xtChip}) .Cases("rdna3", "gfx1100", "gfx1101", "gfx1102", "gfx1103", "gfx1150", "gfx1151", TargetDetails{rdna3Wgp, nullptr}) + .Cases("rdna2", "gfx1030", "gfx1031", "gfx1032", "gfx1033", "gfx1034", + "gfx1035", "gfx1036", TargetDetails{rdna2Wgp, nullptr}) + .Cases("rdna1", "gfx1010", "gfx1011", "gfx1012", "gfx1013", + TargetDetails{rdna1Wgp, nullptr}) .Default(std::nullopt); } @@ -222,41 +246,136 @@ StringRef normalizeAMDGPUTarget(StringRef target) { if (target.starts_with("gfx")) return target; + // We cannot accept rdnaN as a target for LLVM AMDGPU backend; so the + // following is only meant for Vulkan but not HIP. + if (target.starts_with("rdna")) + return target; + return llvm::StringSwitch(target.lower()) .Case("mi300x", "gfx942") .Case("mi300a", "gfx940") .Cases("mi250x", "mi250", "mi210", "cdna2", "gfx90a") + .Case("cdna1", "gfx908") .Cases("rx7900xtx", "rx7900xt", "gfx1100") .Cases("rx7800xt", "rx7700xt", "gfx1101") .Default(StringRef()); } +//===----------------------------------------------------------------------===// +// Known Apple target details +//===----------------------------------------------------------------------===// + +std::optional getAppleTargetDetails() { + ComputeBitwidths computeBitwdiths = + allIntComputeBits | ComputeBitwidths::FP32 | ComputeBitwidths::FP16; + // clang-format off + static const WgpDetails wgp = { + computeBitwdiths, allStorageBits, allSubgroupOps, allDotProductOps, + /*mmaCount=*/0, /*mmaOps=*/nullptr, {32, 32}, + {1024, 1024, 1024}, 1024, 32 * 1024}; + // clang-format on + + return TargetDetails{&wgp, nullptr}; +} + +//===----------------------------------------------------------------------===// +// Known ARM target details +//===----------------------------------------------------------------------===// + +const WgpDetails *getValhallWgpDetails() { + ComputeBitwidths computeBitwdiths = + allIntComputeBits | ComputeBitwidths::FP32 | ComputeBitwidths::FP16; + // clang-format off + static const WgpDetails valhallWgp = { + computeBitwdiths, allStorageBits, allSubgroupOps, allDotProductOps, + /*mmaCount=*/0, /*mmaOps=*/nullptr, {16}, {512, 512, 512}, + 512, 32 * 1024}; + // clang-format on + return &valhallWgp; +} + +std::optional getARMGPUTargetDetails(StringRef target) { + const WgpDetails *valhallWgp = getValhallWgpDetails(); + + // Note that the underlying GPU may have certain capabilities but the Android + // version and driver stack may not expose them. So the following is just and + // will always be approximate. + + return llvm::StringSwitch>(target.lower()) + // Mali-G715: https://vulkan.gpuinfo.org/displayreport.php?id=29754 + .Cases("mali-g715", "mali-g615", "valhall4", + TargetDetails{valhallWgp, nullptr}) + // Mali-G710: https://vulkan.gpuinfo.org/displayreport.php?id=30471 + .Cases("mali-g710", "mali-g510", "mali-g310", "valhall3", + TargetDetails{valhallWgp, nullptr}) + // Mali-G78: https://vulkan.gpuinfo.org/displayreport.php?id=29994 + .Cases("mali-g78", "valhall2", TargetDetails{valhallWgp, nullptr}) + // Mali-G57: https://vulkan.gpuinfo.org/displayreport.php?id=24636 + .Cases("mali-g77", "mali-g57", "valhall1", "valhall", + TargetDetails{valhallWgp, nullptr}) + .Default(std::nullopt); +} + +StringRef normalizeARMGPUTarget(StringRef target) { + if (target == "valhall") + return "valhall1"; + if (target.starts_with("valhall")) + return target; + + return llvm::StringSwitch(target.lower()) + .Cases("mali-g715", "mali-g615", "valhall4") + .Cases("mali-g710", "mali-g510", "mali-g310", "valhall3") + .Case("mali-78", "valhall2") + .Cases("mali-g77", "mali-g57", "valhall1") + .Default(""); +} + //===----------------------------------------------------------------------===// // Known NVIDIA target details //===----------------------------------------------------------------------===// +// FIXME: In the following query functions, we are using AMD WMMA intrinsics +// that have different layout from NVIDIA WMMA intrinsics. This is fine given +// right now we only use this to indicate target features for Vulkan, where all +// cooperative matrix layouts are opaque. We need to create NVIDIA specific WMMA +// intrinsics if we need to have explicit layout analysis and register mapping. + const WgpDetails *getAmpereWgpDetails() { + static const MMAIntrinsic mmaOps[] = { + MMAIntrinsic::WMMA_F16_16x16x16_F32, + MMAIntrinsic::WMMA_F16_16x16x16_F16, + }; static const WgpDetails ampereWgp = { - allComputeBits, allStorageBits, allSubgroupOps, allDotProductOps, 0, - nullptr, // TODO: Add tensor core operations - {32, 32}, {1024, 1024, 1024}, 1024, 163 * 1024}; + allComputeBits, allStorageBits, allSubgroupOps, + allDotProductOps, ARRAY_SIZE(mmaOps), mmaOps, + {32, 32}, {1024, 1024, 1024}, 1024, + 163 * 1024}; return &ereWgp; } const WgpDetails *getTuringWgpDetails() { + static const MMAIntrinsic mmaOps[] = { + MMAIntrinsic::WMMA_F16_16x16x16_F32, + MMAIntrinsic::WMMA_F16_16x16x16_F16, + }; static const WgpDetails turingWgp = { - allComputeBits, allStorageBits, allSubgroupOps, allDotProductOps, 0, - nullptr, // TODO: Add tensor core operations - {32, 32}, {1024, 1024, 1024}, 1024, 64 * 1024}; + allComputeBits, allStorageBits, allSubgroupOps, + allDotProductOps, ARRAY_SIZE(mmaOps), mmaOps, + {32, 32}, {1024, 1024, 1024}, 1024, + 64 * 1024}; return &turingWgp; } const WgpDetails *getVoltaWgpDetails() { + static const MMAIntrinsic mmaOps[] = { + MMAIntrinsic::WMMA_F16_16x16x16_F32, + MMAIntrinsic::WMMA_F16_16x16x16_F16, + }; // clang-format off static const WgpDetails voltaWgp = { - allComputeBits, allStorageBits, allSubgroupOps, DotProductOps::None, - 0, nullptr, // TODO: Add tensor core operations - {32, 32}, {1024, 1024, 1024}, 1024, 96 * 1024}; + allComputeBits, allStorageBits, allSubgroupOps, DotProductOps::None, + ARRAY_SIZE(mmaOps), mmaOps, {32, 32}, {1024, 1024, 1024}, + 1024, 96 * 1024}; // clang-format on return &voltaWgp; } @@ -332,15 +451,126 @@ StringRef normalizeNVIDIAGPUTarget(StringRef target) { .Default(StringRef()); } +//===----------------------------------------------------------------------===// +// Known Qualcomm target details +//===----------------------------------------------------------------------===// + +const WgpDetails *getAdrenoWgpDetails() { + auto computeBitwdiths = ComputeBitwidths::Int32 | ComputeBitwidths::Int16 | + ComputeBitwidths::Int8 | ComputeBitwidths::FP32 | + ComputeBitwidths::FP16; + auto storageBitwidths = + StorageBitwidths::B64 | StorageBitwidths::B32 | StorageBitwidths::B16; + // clang-format off + static const WgpDetails adrenoWgp = { + computeBitwdiths, storageBitwidths, allSubgroupOps, + allDotProductOps, /*mmaCount=*/0, /*mmaOps=*/nullptr, + {64}, {1024, 1024, 1024}, 1024, + 32 * 1024}; + // clang-format on + return &adrenoWgp; +} + +bool verifyQualcommGPUTarget(StringRef target) { + if (target == "adreno") + return true; + + StringRef t = target; + if (!t.consume_front("adreno-")) + return false; + + // The can exist an optional L at the end. + if (t.ends_with("l")) + t = t.drop_back(); + + // Check whether we have a product number + unsigned number = 0; + // StringRef::consumeInteger() returns true to signify errors. + if (t.size() != 3 || t.consumeInteger(10, number)) + return false; + + return true; +} + +std::optional getQualcommGPUTargetDetails(StringRef target) { + const WgpDetails *adrenoWgp = getAdrenoWgpDetails(); + + // Note that the underlying GPU may have certain capabilities but the Android + // version and driver stack may not expose them. So the following is just and + // will always be approximate. + + // Adreno GPUs are quite opaque regarding their generational information. + // So right now we only have one target description for all cases. + // + // Though some example Adreno GPUs: + // Adreno-750: https://vulkan.gpuinfo.org/displayreport.php?id=27414 + // Adreno-740: https://vulkan.gpuinfo.org/displayreport.php?id=19218 + // Adreno-730: https://vulkan.gpuinfo.org/displayreport.php?id=19382 + if (verifyQualcommGPUTarget(target)) + return TargetDetails{adrenoWgp, nullptr}; + + return std::nullopt; +} + +//===----------------------------------------------------------------------===// +// Vulkan profile details +//===----------------------------------------------------------------------===// + +const WgpDetails *getAndroidBaseline2022WgpDetails() { + // The following details are from + // https://github.com/KhronosGroup/Vulkan-Profiles/blob/main/profiles/VP_ANDROID_baseline_2022.json + + auto computeBitwdiths = ComputeBitwidths::Int32 | ComputeBitwidths::FP32; + auto storageBitwidths = StorageBitwidths::B32; + // FIXME: We cannot have a fixed subgroup size to target a profile; need to + // have different targets for different subgroup sizes, or change CodeGen to + // use symbolic subgroup size values, which can be hard for reduction. + // It's kinda fine now given we don't allow any subgroup ops anyway here.. + + // clang-format off + static const WgpDetails androidWgp = { + computeBitwdiths, storageBitwidths, SubgroupOps::None, + DotProductOps::None, /*mmaCount=*/0, /*mmaOps=*/nullptr, + {64, 64}, {128, 128, 64}, 128, + 16 * 1024}; + // clang-format on + return &androidWgp; +} + +std::optional getAndroidProfileDetails(StringRef target) { + const WgpDetails *baseline2022Wgp = getAndroidBaseline2022WgpDetails(); + + return llvm::StringSwitch>(target.lower()) + .Case("vp_android_baseline_2022", TargetDetails{baseline2022Wgp, nullptr}) + .Default(std::nullopt); +} + } // namespace //===----------------------------------------------------------------------===// // Query functions //===----------------------------------------------------------------------===// +TargetAttr getMetalTargetDetails(MLIRContext *context) { + return createTargetAttr(*getAppleTargetDetails(), /*arch=*/"", + /*features=*/"spirv:v1.3,cap:Shader", context); +} + +TargetAttr getCUDATargetDetails(StringRef target, StringRef features, + MLIRContext *context) { + if (std::optional details = getNVIDIAGPUTargetDetails(target)) + return createTargetAttr(*details, normalizeNVIDIAGPUTarget(target), + features, context); + return nullptr; +} + +StringRef normalizeCUDATarget(StringRef target) { + return normalizeNVIDIAGPUTarget(target); +} + TargetAttr getHIPTargetDetails(StringRef target, StringRef features, MLIRContext *context) { - if (auto details = getAMDGPUTargetDetails(target)) { + if (std::optional details = getAMDGPUTargetDetails(target)) { return createTargetAttr(*details, normalizeAMDGPUTarget(target), features, context); } @@ -351,16 +581,62 @@ StringRef normalizeHIPTarget(StringRef target) { return normalizeAMDGPUTarget(target); } -TargetAttr getCUDATargetDetails(StringRef target, StringRef features, - MLIRContext *context) { - if (auto details = getNVIDIAGPUTargetDetails(target)) +TargetAttr getVulkanTargetDetails(llvm::StringRef target, + MLIRContext *context) { + // Go through each vendor's target details. This assumes we won't have + // duplicated product or microarchitecture names among vendors, which should + // be the case. + + // For mobile GPUs we target Vulkan 1.1, which accepts SPIR-V 1.3 as the + // maximum. But the VK_KHR_spirv_1_4 extension is commonly available so we use + // SPIR-V 1.4. For non-mobile GPUs we target Vulkan 1.3, which accepts + // SPIR-V 1.6 as the maximum. + + if (std::optional details = getAMDGPUTargetDetails(target)) { + return createTargetAttr(*details, normalizeAMDGPUTarget(target), + /*features=*/"spirv:v1.6,cap:Shader", context); + } + if (std::optional details = getARMGPUTargetDetails(target)) { + return createTargetAttr(*details, normalizeARMGPUTarget(target), + /*features=*/"spirv:v1.4,cap:Shader", context); + } + if (std::optional details = + getNVIDIAGPUTargetDetails(target)) { return createTargetAttr(*details, normalizeNVIDIAGPUTarget(target), - features, context); + /*features=*/"spirv:v1.6,cap:Shader", context); + } + if (std::optional details = + getQualcommGPUTargetDetails(target)) { + return createTargetAttr(*details, target, + /*features=*/"spirv:v1.4,cap:Shader", context); + } + + // Go through common profiles if not hit in the above. + + if (std::optional details = getAndroidProfileDetails(target)) { + return createTargetAttr(*details, target, + /*features=*/"spirv:v1.3,cap:Shader", context); + } return nullptr; } -StringRef normalizeCUDATarget(StringRef target) { - return normalizeNVIDIAGPUTarget(target); +TargetAttr getWebGPUTargetDetails(MLIRContext *context) { + // TODO(scotttodd): find list of SPIR-V capabilities and extensions supported + // by WebGPU/WGSL. + auto computeBitwdiths = ComputeBitwidths::Int32 | ComputeBitwidths::FP32; + auto storageBitwidths = StorageBitwidths::B32; + // clang-format off + static const WgpDetails wgp = { + computeBitwdiths, storageBitwidths, SubgroupOps::None, + DotProductOps::None, /*mmaCount=*/0, /*mmaOps=*/nullptr, + {32, 32}, {128, 128, 64}, 128, + 16 * 1024}; + // clang-format on + + return createTargetAttr( + {&wgp, nullptr}, /*arch=*/"", + "spirv:v1.0,cap:Shader,ext:SPV_KHR_storage_buffer_storage_class", + context); } TargetAttr getFullTarget(StringRef targetAPI, StringRef aliasTarget, diff --git a/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.h b/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.h index ffe9a15b4663..d9698cc912f0 100644 --- a/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.h +++ b/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.h @@ -12,6 +12,22 @@ namespace mlir::iree_compiler::IREE::GPU { +// Returns a TargetAttr to target Metal via SPIR-V CodeGen. +TargetAttr getMetalTargetDetails(MLIRContext *context); + +// Returns a TargetAttr to describe the details of the given |target|, which can +// be a product name like "rtx3090", an microarchitecture name like "ampere", or +// a compute capability like "sm_80", with a list of comma-separated target +// |features|. Returns a null TargetAttr if the given |target| is not +// recognized. +TargetAttr getCUDATargetDetails(llvm::StringRef target, + llvm::StringRef features, MLIRContext *context); + +// Normalizes the given CUDA |target| to the gfx target commonly used for +// compiling towards CUDA. For example, "sm_80" for "a100", "sm_89" for "ada". +// if the given |target| is not recognized. +StringRef normalizeCUDATarget(StringRef target); + // Returns a TargetAttr to describe the details of the given |target|, which can // be a product name like "rx7900xtx", an microarchitecture name like "rdna3", // or a compiler target like "gfx1100", with a list of comma-separated @@ -26,16 +42,13 @@ TargetAttr getHIPTargetDetails(llvm::StringRef target, llvm::StringRef features, StringRef normalizeHIPTarget(StringRef target); // Returns a TargetAttr to describe the details of the given |target|, which can -// be a product name like "rtx3090", an microarchitecture name like "ampere", or -// a compute capability like "sm_80", with a list of comma-separated target -// |features|. TargetAttr if the given |target| is not recognized. -TargetAttr getCUDATargetDetails(llvm::StringRef target, - llvm::StringRef features, MLIRContext *context); +// be a product name like "rtx3090"/"mali-g710"/"adreno" or an microarchitecture +// name like "ampere"/"valhall". Returns a null TargetAttr if the given |target| +// is not recognized. +TargetAttr getVulkanTargetDetails(llvm::StringRef target, MLIRContext *context); -// Normalizes the given CUDA |target| to the gfx target commonly used for -// compiling towards CUDA. For example, "sm_80" for "a100", "sm_89" for "ada". -// if the given |target| is not recognized. -StringRef normalizeCUDATarget(StringRef target); +// Returns a TargetAttr to target WebGPU via SPIR-V CodeGen. +TargetAttr getWebGPUTargetDetails(MLIRContext *context); // Returns the full target of the given |aliasTarget| with a list of // comma-separated target |features|. Returns null target if unknown. diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp b/compiler/src/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp index d73740f54c79..39678b9a570f 100644 --- a/compiler/src/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp +++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp @@ -74,9 +74,6 @@ namespace { using CodeGenPipeline = IREE::Codegen::DispatchLoweringPassPipeline; -constexpr StringLiteral kCudaTarget = "cuda"; -constexpr StringLiteral kRocmTarget = "rocm"; - // Threshold used to determine whether a matmul dimension is 'very skinny'. constexpr int64_t kVerySkinnyDimThreshold = 4; @@ -92,6 +89,10 @@ constexpr unsigned softwarePipelineDepthSimt = 0; } // namespace +bool isROCmBackend(IREE::GPU::TargetAttr target) { + return target.getArch().starts_with("gfx"); +} + //====---------------------------------------------------------------------===// // Matmul Configuration Helpers //====---------------------------------------------------------------------===// @@ -576,6 +577,8 @@ static LogicalResult setVectorDistributionConfig(IREE::GPU::TargetAttr target, mlir::FunctionOpInterface entryPoint, Operation *computeOp) { + if (!isROCmBackend(target)) + return failure(); if (!clGPUEnableVectorDistribution) { LDBG("Vector Distribution not enabled, skipping..."); @@ -1188,15 +1191,6 @@ static bool isMatvecLike(linalg::LinalgOp linalgOp) { // Warp Reduction Pipeline Configuration //====---------------------------------------------------------------------===// -bool isROCmBackend(mlir::FunctionOpInterface entryPoint) { - if (auto targetAttr = IREE::HAL::ExecutableTargetAttr::lookup(entryPoint)) { - if (auto backend = targetAttr.getBackend()) { - return backend.getValue() == "rocm"; - } - } - return false; -} - /// Set the configuration for reductions that can be mapped to warp reductions. static LogicalResult setWarpReductionConfig(IREE::GPU::TargetAttr target, @@ -1367,8 +1361,8 @@ setWarpReductionConfig(IREE::GPU::TargetAttr target, // // TODO: This is enabled for matvec on ROCm for now. We should // validate this strategy and extend to more linalg generics and to CUDA. - if (isROCmBackend(entryPoint) && - llvm::none_of(bounds, ShapedType::isDynamic) && isMatvecLike(op)) { + if (isROCmBackend(target) && llvm::none_of(bounds, ShapedType::isDynamic) && + isMatvecLike(op)) { int64_t lastParallelBound = bounds[parallelDims.back()]; int64_t numParallelReductions = 1; const int64_t maxParallelFactor = groupSize / 4; diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/BUILD.bazel b/compiler/src/iree/compiler/Codegen/SPIRV/BUILD.bazel index 5c804451cbb3..e72fdc56f0d8 100644 --- a/compiler/src/iree/compiler/Codegen/SPIRV/BUILD.bazel +++ b/compiler/src/iree/compiler/Codegen/SPIRV/BUILD.bazel @@ -58,6 +58,7 @@ iree_compiler_cc_library( "Passes.cpp", "SPIRVAnnotateWinogradLoops.cpp", "SPIRVBreakDownLargeVector.cpp", + "SPIRVConvertGPUTarget.cpp", "SPIRVEmulateI64.cpp", "SPIRVEraseStorageBufferStaticShape.cpp", "SPIRVFinalVectorLowering.cpp", diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/CMakeLists.txt b/compiler/src/iree/compiler/Codegen/SPIRV/CMakeLists.txt index 13632df10b5c..1378bbc4483c 100644 --- a/compiler/src/iree/compiler/Codegen/SPIRV/CMakeLists.txt +++ b/compiler/src/iree/compiler/Codegen/SPIRV/CMakeLists.txt @@ -57,6 +57,7 @@ iree_cc_library( "Passes.cpp" "SPIRVAnnotateWinogradLoops.cpp" "SPIRVBreakDownLargeVector.cpp" + "SPIRVConvertGPUTarget.cpp" "SPIRVEmulateI64.cpp" "SPIRVEraseStorageBufferStaticShape.cpp" "SPIRVFinalVectorLowering.cpp" diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/Passes.cpp b/compiler/src/iree/compiler/Codegen/SPIRV/Passes.cpp index 538e3ba6340f..75eb17d9aa72 100644 --- a/compiler/src/iree/compiler/Codegen/SPIRV/Passes.cpp +++ b/compiler/src/iree/compiler/Codegen/SPIRV/Passes.cpp @@ -631,6 +631,8 @@ static void buildSPIRVCodegenConfigurationPassPipelineImpl( void buildSPIRVCodegenConfigurationPassPipeline( OpPassManager &variantPassManager) { + // TODO: move the following pass to be immediately before ConvertToSPIRVPass. + variantPassManager.addPass(createSPIRVConvertGPUTargetPass()); OpPassManager &modulePassManager = variantPassManager.nest(); buildSPIRVCodegenConfigurationPassPipelineImpl(modulePassManager); } diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/Passes.h b/compiler/src/iree/compiler/Codegen/SPIRV/Passes.h index e1dc8306d9e6..a0b0d1693a9a 100644 --- a/compiler/src/iree/compiler/Codegen/SPIRV/Passes.h +++ b/compiler/src/iree/compiler/Codegen/SPIRV/Passes.h @@ -86,6 +86,10 @@ createSPIRVAnnotateWinogradLoopsPass(); std::unique_ptr> createSPIRVBreakDownLargeVectorPass(); +// Converts #iree_gpu.target into #spirv.target_env. +std::unique_ptr> +createSPIRVConvertGPUTargetPass(); + /// Emulates bfloat 16 ops with 32-bit float ops. std::unique_ptr> createSPIRVEmulateBf16Pass(); diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/Passes.td b/compiler/src/iree/compiler/Codegen/SPIRV/Passes.td index 29b396de3927..dc94eb2b4d04 100644 --- a/compiler/src/iree/compiler/Codegen/SPIRV/Passes.td +++ b/compiler/src/iree/compiler/Codegen/SPIRV/Passes.td @@ -33,6 +33,13 @@ def SPIRVBreakDownLargeVector : InterfacePass<"iree-spirv-breakdown-large-vector let constructor = "mlir::iree_compiler::createSPIRVBreakDownLargeVectorPass()"; } +def SPIRVConvertGPUTarget : + Pass<"iree-spirv-convert-gpu-target", + "mlir::iree_compiler::IREE::HAL::ExecutableVariantOp"> { + let summary = "Convert #iree_gpu.target into #spirv.target_env"; + let constructor = "mlir::iree_compiler::createSPIRVConvertGPUTargetPass()"; +} + def SPIRVEmulateI64 : InterfacePass<"iree-spirv-emulate-i64", "mlir::FunctionOpInterface"> { let summary = "Emulate 64-bit integer ops with 32-bit integer ops"; diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVConvertGPUTarget.cpp b/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVConvertGPUTarget.cpp new file mode 100644 index 000000000000..3b8c3ba83021 --- /dev/null +++ b/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVConvertGPUTarget.cpp @@ -0,0 +1,288 @@ +// Copyright 2024 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#include "iree/compiler/Codegen/SPIRV/PassDetail.h" +#include "iree/compiler/Codegen/SPIRV/Passes.h" +#include "iree/compiler/Codegen/Utils/GPUUtils.h" +#include "iree/compiler/Dialect/HAL/IR/HALOps.h" +#include "llvm/ADT/STLExtras.h" +#include "llvm/ADT/StringExtras.h" +#include "llvm/ADT/StringSwitch.h" +#include "mlir/Dialect/SPIRV/IR/SPIRVAttributes.h" +#include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" +#include "mlir/Dialect/SPIRV/IR/SPIRVEnums.h" +#include "mlir/Dialect/SPIRV/IR/TargetAndABI.h" +#include "mlir/IR/Builders.h" +#include "mlir/IR/BuiltinAttributes.h" +#include "mlir/Pass/Pass.h" + +namespace mlir::iree_compiler { + +namespace { + +using IREE::GPU::ComputeBitwidths; +using IREE::GPU::DotProductOps; +using IREE::GPU::StorageBitwidths; +using IREE::GPU::SubgroupOps; + +using spirv::Capability; +using spirv::ClientAPI; +using spirv::Extension; +using spirv::Vendor; +using spirv::Version; + +//===----------------------------------------------------------------------===// +// Freeform features +//===----------------------------------------------------------------------===// + +// Scans the given |features| list and pushes SPIR-V version specification of +// 'spirv:v1.x' format into |caps|. +std::optional deduceVersion(ArrayRef features) { + for (StringRef feature : features) { + if (feature.consume_front("spirv:v1.")) { + return llvm::StringSwitch>(feature) + .Case("6", Version::V_1_6) + .Case("5", Version::V_1_5) + .Case("4", Version::V_1_4) + .Case("3", Version::V_1_3) + .Case("2", Version::V_1_2) + .Case("1", Version::V_1_1) + .Case("0", Version::V_1_0) + .Default(std::nullopt); + } + } + return std::nullopt; +} + +// Scans the given |features| list and pushes capability specification with +// 'cap:' prefix into |caps|. +std::optional processCapabilities(ArrayRef features, + SetVector &caps) { + for (StringRef feature : features) { + if (feature.consume_front("cap:")) { + if (std::optional cap = spirv::symbolizeCapability(feature)) + caps.insert(*cap); + } + } + return std::nullopt; +} + +// Scans the given |features| list and pushes extension specification with +// 'ext:' prefix into |exts|. +std::optional processExtensions(ArrayRef features, + SetVector &exts) { + for (StringRef feature : features) { + if (feature.consume_front("ext:")) { + if (std::optional ext = spirv::symbolizeExtension(feature)) + exts.insert(*ext); + } + } + return std::nullopt; +} + +//===----------------------------------------------------------------------===// +// Client API and vendor +//===----------------------------------------------------------------------===// + +ClientAPI deduceClientAPI(StringRef backend) { + return llvm::StringSwitch(backend) + .Case("vulkan", ClientAPI::Vulkan) + .Case("metal", ClientAPI::Metal) + .Case("webgpu", ClientAPI::WebGPU) + .Case("opencl", ClientAPI::OpenCL) + .Default(ClientAPI::Unknown); +} + +Vendor deduceVendor(StringRef arch) { + if (arch.starts_with("gfx") || arch.starts_with("rdna")) + return Vendor::AMD; + if (arch.starts_with("mali")) + return Vendor::ARM; + if (arch.starts_with("sm_")) + return Vendor::NVIDIA; + if (arch.starts_with("adreno")) + return Vendor::Qualcomm; + return Vendor::Unknown; +} + +//===----------------------------------------------------------------------===// +// Workgroup-processor features and limits +//===----------------------------------------------------------------------===// + +void addComputeFeatures(ComputeBitwidths compute, SetVector &caps, + SetVector &exts) { + if (bitEnumContainsAny(compute, ComputeBitwidths::FP64)) + caps.insert(Capability::Float64); + // FP32 does not need special capabilities or extensions. + if (bitEnumContainsAny(compute, ComputeBitwidths::FP16)) + caps.insert(Capability::Float16); + + if (bitEnumContainsAny(compute, ComputeBitwidths::Int64)) + caps.insert(Capability::Int64); + // Int32 does not need special capabilities or extensions. + if (bitEnumContainsAny(compute, ComputeBitwidths::Int16)) + caps.insert(Capability::Int16); + if (bitEnumContainsAny(compute, ComputeBitwidths::Int8)) + caps.insert(Capability::Int8); +} + +void addStorageFeatures(StorageBitwidths storage, SetVector &caps, + SetVector &exts) { + // 64bit does not need special capabilities or extensions. + // 32bit does not need special capabilities or extensions. + if (bitEnumContainsAny(storage, StorageBitwidths::B16)) { + caps.insert(Capability::StorageBuffer16BitAccess); + caps.insert(Capability::StorageUniform16); + caps.insert(Capability::StoragePushConstant16); + exts.insert(Extension::SPV_KHR_16bit_storage); + } + if (bitEnumContainsAny(storage, StorageBitwidths::B8)) { + caps.insert(Capability::StorageBuffer8BitAccess); + caps.insert(Capability::UniformAndStorageBuffer8BitAccess); + caps.insert(Capability::StoragePushConstant8); + exts.insert(Extension::SPV_KHR_8bit_storage); + } +} + +void addSubgroupFeatures(SubgroupOps subgroup, SetVector &caps, + SetVector &exts) { + if (bitEnumContainsAny(subgroup, SubgroupOps::Shuffle)) { + caps.insert(Capability::GroupNonUniformShuffle); + caps.insert(Capability::GroupNonUniformShuffleRelative); + } + if (bitEnumContainsAny(subgroup, SubgroupOps::Arithmetic)) { + caps.insert(Capability::GroupNonUniformArithmetic); + } +} + +void addDotProductFeatures(ComputeBitwidths compute, DotProductOps dotProduct, + SetVector &caps, + SetVector &exts) { + if (bitEnumContainsAny(dotProduct, DotProductOps::DP4xI8ToI32)) { + caps.insert(Capability::DotProduct); + caps.insert(Capability::DotProductInput4x8BitPacked); // Use i32 input + caps.insert(Capability::DotProductInputAll); // Use vector<*> input + if (bitEnumContainsAny(compute, ComputeBitwidths::Int8)) { + caps.insert(Capability::DotProductInput4x8Bit); // Use vector<4xi8> input + } + exts.insert(Extension::SPV_KHR_integer_dot_product); + } +} + +void addMatrixFeatures(IREE::GPU::MMAOpsArrayAttr mmaOps, + SetVector &caps, SetVector &exts, + SetVector &coopMatAttrs) { + if (!mmaOps.empty()) { + caps.insert(Capability::CooperativeMatrixKHR); + exts.insert(Extension::SPV_KHR_cooperative_matrix); + } +} + +spirv::ResourceLimitsAttr convertLimits(StringRef arch, + IREE::GPU::TargetWgpAttr wgp) { + MLIRContext *context = wgp.getContext(); + Builder b(context); + + SmallVector coopMatAttrs; + for (IREE::GPU::MMAAttr mmaOp : wgp.getMma()) { + auto [mSize, nSize, kSize] = mmaOp.getMNKShape(); + auto [aType, bType, cType] = mmaOp.getABCElementTypes(); + coopMatAttrs.push_back(spirv::CooperativeMatrixPropertiesKHRAttr::get( + context, mSize, nSize, kSize, aType, bType, cType, cType, + false /*saturatingAccumulation*/, + spirv::ScopeAttr::get(context, spirv::Scope::Subgroup))); + } + + ArrayRef subgroupSizes = wgp.getSubgroupSizeChoices().asArrayRef(); + const int minSubgroupSize = *llvm::min_element(subgroupSizes); + const int maxSubgroupSize = *llvm::max_element(subgroupSizes); + // This is mostly to match RDNA behavior on Vulkan--RDNA supports either 32 or + // 64 as subgroup sizes; the default subgroup size is 64. + const int preferredSubgroupSize = maxSubgroupSize; + + return spirv::ResourceLimitsAttr::get( + context, wgp.getMaxWorkgroupMemoryBytes(), + wgp.getMaxThreadCountPerWorkgroup(), + b.getI32ArrayAttr(wgp.getMaxWorkgroupSizes().asArrayRef()), + preferredSubgroupSize, minSubgroupSize, maxSubgroupSize, + ArrayAttr::get(context, coopMatAttrs), ArrayAttr{}); +} + +//===----------------------------------------------------------------------===// +// Target specification conversion +//===----------------------------------------------------------------------===// + +FailureOr +convertGPUTarget(IREE::HAL::ExecutableVariantOp variant) { + IREE::HAL::ExecutableTargetAttr target = variant.getTarget(); + IREE::GPU::TargetAttr gpuTarget = getGPUTargetAttr(target); + + SmallVector features; + llvm::SplitString(gpuTarget.getFeatures(), features, ","); + + SetVector caps; + SetVector exts; + SetVector coopMatAttrs; + + std::optional version = deduceVersion(features); + if (!version) { + return variant.emitError("cannot deduce spirv version from target " + "features; need to specify 'spirv1.x'"); + } + processCapabilities(features, caps); + processExtensions(features, exts); + + IREE::GPU::TargetWgpAttr wgp = gpuTarget.getWgp(); + ComputeBitwidths compute = wgp.getCompute().getValue(); + addComputeFeatures(compute, caps, exts); + addStorageFeatures(wgp.getStorage().getValue(), caps, exts); + addSubgroupFeatures(wgp.getSubgroup().getValue(), caps, exts); + addDotProductFeatures(compute, wgp.getDot().getValue(), caps, exts); + addMatrixFeatures(wgp.getMma(), caps, exts, coopMatAttrs); + + auto triple = spirv::VerCapExtAttr::get( + *version, caps.getArrayRef(), exts.getArrayRef(), variant.getContext()); + return spirv::TargetEnvAttr::get( + triple, convertLimits(gpuTarget.getArch(), wgp), + deduceClientAPI(target.getBackend()), deduceVendor(gpuTarget.getArch()), + spirv::DeviceType::Unknown, spirv::TargetEnvAttr::kUnknownDeviceID); +} + +struct SPIRVConvertGPUTargetPass final + : SPIRVConvertGPUTargetBase { + void getDependentDialects(DialectRegistry ®istry) const override { + registry.insert(); + } + + void runOnOperation() override { + IREE::HAL::ExecutableVariantOp variant = getOperation(); + IREE::HAL::ExecutableTargetAttr target = variant.getTarget(); + + FailureOr spirvTarget = convertGPUTarget(variant); + if (failed(spirvTarget)) + return signalPassFailure(); + + Builder b(&getContext()); + auto attrs = llvm::to_vector(target.getConfiguration().getValue()); + attrs.emplace_back(b.getStringAttr(spirv::getTargetEnvAttrName()), + *spirvTarget); + auto configAttr = b.getDictionaryAttr(attrs); + + auto halTarget = IREE::HAL::ExecutableTargetAttr::get( + target.getContext(), target.getBackend(), target.getFormat(), + configAttr); + variant.setTargetAttr(halTarget); + } +}; + +} // namespace + +std::unique_ptr> +createSPIRVConvertGPUTargetPass() { + return std::make_unique(); +} + +} // namespace mlir::iree_compiler diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/test/BUILD.bazel b/compiler/src/iree/compiler/Codegen/SPIRV/test/BUILD.bazel index ed6a8be75185..eb5947497d0b 100644 --- a/compiler/src/iree/compiler/Codegen/SPIRV/test/BUILD.bazel +++ b/compiler/src/iree/compiler/Codegen/SPIRV/test/BUILD.bazel @@ -39,6 +39,7 @@ iree_lit_test_suite( "config_nvidia_matmul_cooperative_ops.mlir", "config_user.mlir", "convert_to_spirv.mlir", + "convert_gpu_target.mlir", "emulate_i64.mlir", "erase_storage_buffer_static_shape.mlir", "illegal_configuration.mlir", diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/test/CMakeLists.txt b/compiler/src/iree/compiler/Codegen/SPIRV/test/CMakeLists.txt index 3dc827725e1f..273e581158e1 100644 --- a/compiler/src/iree/compiler/Codegen/SPIRV/test/CMakeLists.txt +++ b/compiler/src/iree/compiler/Codegen/SPIRV/test/CMakeLists.txt @@ -34,6 +34,7 @@ iree_lit_test_suite( "config_nvidia_matmul.mlir" "config_nvidia_matmul_cooperative_ops.mlir" "config_user.mlir" + "convert_gpu_target.mlir" "convert_to_spirv.mlir" "emulate_i64.mlir" "erase_storage_buffer_static_shape.mlir" diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/test/convert_gpu_target.mlir b/compiler/src/iree/compiler/Codegen/SPIRV/test/convert_gpu_target.mlir new file mode 100644 index 000000000000..b1f809293db2 --- /dev/null +++ b/compiler/src/iree/compiler/Codegen/SPIRV/test/convert_gpu_target.mlir @@ -0,0 +1,36 @@ +// RUN: iree-opt --split-input-file --pass-pipeline='builtin.module(hal.executable(hal.executable.variant(iree-spirv-convert-gpu-target)))' %s | FileCheck %s + +hal.executable @dispatch { +hal.executable.variant public @vulkan_spirv_fb target(<"vulkan-spirv", "vulkan-spirv-fb", { + iree.gpu.target = #iree_gpu.target, ], + subgroup_size_choices = [32, 64], max_workgroup_sizes = [1024, 1024, 1024], max_thread_count_per_workgroup = 1024, max_workgroup_memory_bytes = 65536>>}>) { + hal.executable.export public @dispatch ordinal(0) layout(#hal.pipeline.layout]>]>) { + ^bb0(%arg0: !hal.device): + %x, %y, %z = flow.dispatch.workgroup_count_from_slice + hal.return %x, %y, %z : index, index, index + } + builtin.module { + func.func @dispatch() { + return + } + } +} +} + +// CHECK: spirv.target_env = #spirv.target_env<#spirv.vce, +// CHECK-SAME: AMD, +// CHECK-SAME: #spirv.resource_limits>, +// CHECK-SAME: #spirv.coop_matrix_props_khr> +// CHECK-SAME: ]>> diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/BUILD.bazel b/compiler/src/iree/compiler/Dialect/Vulkan/BUILD.bazel deleted file mode 100644 index 236a47446725..000000000000 --- a/compiler/src/iree/compiler/Dialect/Vulkan/BUILD.bazel +++ /dev/null @@ -1,11 +0,0 @@ -# Copyright 2020 The IREE Authors -# -# Licensed under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception - -package( - default_visibility = ["//visibility:public"], - features = ["layering_check"], - licenses = ["notice"], # Apache 2.0 -) diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/CMakeLists.txt b/compiler/src/iree/compiler/Dialect/Vulkan/CMakeLists.txt deleted file mode 100644 index 487e4f10fcf4..000000000000 --- a/compiler/src/iree/compiler/Dialect/Vulkan/CMakeLists.txt +++ /dev/null @@ -1,13 +0,0 @@ -################################################################################ -# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from # -# compiler/src/iree/compiler/Dialect/Vulkan/BUILD.bazel # -# # -# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary # -# CMake-only content. # -# # -# To disable autogeneration for this file entirely, delete this header. # -################################################################################ - -iree_add_all_subdirs() - -### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ### diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/IR/BUILD.bazel b/compiler/src/iree/compiler/Dialect/Vulkan/IR/BUILD.bazel deleted file mode 100644 index da4b65ecd178..000000000000 --- a/compiler/src/iree/compiler/Dialect/Vulkan/IR/BUILD.bazel +++ /dev/null @@ -1,87 +0,0 @@ -# Copyright 2020 The IREE Authors -# -# Licensed under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception - -load("//build_tools/bazel:build_defs.oss.bzl", "iree_compiler_cc_library", "iree_gentbl_cc_library", "iree_td_library") -load("//build_tools/bazel:enforce_glob.bzl", "enforce_glob") - -package( - default_visibility = ["//visibility:public"], - features = ["layering_check"], - licenses = ["notice"], # Apache 2.0 -) - -iree_td_library( - name = "td_files", - srcs = enforce_glob( - [ - "VulkanAttributes.td", - "VulkanBase.td", - ], - include = ["*.td"], - ), - deps = ["@llvm-project//mlir:OpBaseTdFiles"], -) - -iree_compiler_cc_library( - name = "IR", - srcs = [ - "VulkanAttributes.cpp", - "VulkanAttributes.cpp.inc", - "VulkanDialect.cpp", - "VulkanEnums.cpp.inc", - "VulkanTypes.cpp", - ], - hdrs = [ - "VulkanAttributes.h", - "VulkanAttributes.h.inc", - "VulkanDialect.h", - "VulkanEnums.h.inc", - "VulkanTypes.h", - ], - deps = [ - ":VulkanAttrsGen", - ":VulkanEnumsGen", - "//compiler/src/iree/compiler/Dialect/Util/IR", - "@llvm-project//llvm:Support", - "@llvm-project//mlir:IR", - "@llvm-project//mlir:SPIRVDialect", - "@llvm-project//mlir:Support", - ], -) - -iree_gentbl_cc_library( - name = "VulkanAttrsGen", - tbl_outs = [ - ( - ["--gen-attrdef-decls"], - "VulkanAttributes.h.inc", - ), - ( - ["--gen-attrdef-defs"], - "VulkanAttributes.cpp.inc", - ), - ], - tblgen = "@llvm-project//mlir:mlir-tblgen", - td_file = "VulkanAttributes.td", - deps = [":td_files"], -) - -iree_gentbl_cc_library( - name = "VulkanEnumsGen", - tbl_outs = [ - ( - ["--gen-enum-decls"], - "VulkanEnums.h.inc", - ), - ( - ["--gen-enum-defs"], - "VulkanEnums.cpp.inc", - ), - ], - tblgen = "@llvm-project//mlir:mlir-tblgen", - td_file = "VulkanBase.td", - deps = [":td_files"], -) diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/IR/CMakeLists.txt b/compiler/src/iree/compiler/Dialect/Vulkan/IR/CMakeLists.txt deleted file mode 100644 index 3b03c56e8582..000000000000 --- a/compiler/src/iree/compiler/Dialect/Vulkan/IR/CMakeLists.txt +++ /dev/null @@ -1,59 +0,0 @@ -################################################################################ -# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from # -# compiler/src/iree/compiler/Dialect/Vulkan/IR/BUILD.bazel # -# # -# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary # -# CMake-only content. # -# # -# To disable autogeneration for this file entirely, delete this header. # -################################################################################ - -iree_add_all_subdirs() - -iree_cc_library( - NAME - IR - HDRS - "VulkanAttributes.h" - "VulkanAttributes.h.inc" - "VulkanDialect.h" - "VulkanEnums.h.inc" - "VulkanTypes.h" - SRCS - "VulkanAttributes.cpp" - "VulkanAttributes.cpp.inc" - "VulkanDialect.cpp" - "VulkanEnums.cpp.inc" - "VulkanTypes.cpp" - DEPS - ::VulkanAttrsGen - ::VulkanEnumsGen - LLVMSupport - MLIRIR - MLIRSPIRVDialect - MLIRSupport - iree::compiler::Dialect::Util::IR - PUBLIC -) - -iree_tablegen_library( - NAME - VulkanAttrsGen - TD_FILE - "VulkanAttributes.td" - OUTS - --gen-attrdef-decls VulkanAttributes.h.inc - --gen-attrdef-defs VulkanAttributes.cpp.inc -) - -iree_tablegen_library( - NAME - VulkanEnumsGen - TD_FILE - "VulkanBase.td" - OUTS - --gen-enum-decls VulkanEnums.h.inc - --gen-enum-defs VulkanEnums.cpp.inc -) - -### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ### diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanAttributes.cpp b/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanAttributes.cpp deleted file mode 100644 index dc33c2b2cb9f..000000000000 --- a/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanAttributes.cpp +++ /dev/null @@ -1,359 +0,0 @@ -// Copyright 2020 The IREE Authors -// -// Licensed under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception - -#include "iree/compiler/Dialect/Vulkan/IR/VulkanAttributes.h" - -#include "iree/compiler/Dialect/Vulkan/IR/VulkanDialect.h" -#include "iree/compiler/Dialect/Vulkan/IR/VulkanTypes.h" -#include "llvm/ADT/TypeSwitch.h" -#include "llvm/Support/SMLoc.h" -#include "mlir/Dialect/SPIRV/IR/SPIRVAttributes.h" -#include "mlir/IR/AttributeSupport.h" -#include "mlir/IR/Builders.h" -#include "mlir/IR/BuiltinAttributes.h" -#include "mlir/IR/BuiltinTypes.h" -#include "mlir/IR/Diagnostics.h" -#include "mlir/IR/DialectImplementation.h" -#include "mlir/IR/Location.h" - -#define GET_ATTRDEF_CLASSES -#include "iree/compiler/Dialect/Vulkan/IR/VulkanAttributes.cpp.inc" // IWYU pragma: keep - -namespace mlir::iree_compiler::IREE::Vulkan { - -//===----------------------------------------------------------------------===// -// TargetEnv -//===----------------------------------------------------------------------===// - -namespace detail { -struct TargetEnvAttributeStorage : public AttributeStorage { - using KeyTy = std::tuple; - - TargetEnvAttributeStorage(Attribute version, Attribute revision, - Attribute extensions, spirv::Vendor vendorID, - spirv::DeviceType deviceType, uint32_t deviceID, - Attribute capabilities) - : version(version), revision(revision), extensions(extensions), - capabilities(capabilities), vendorID(vendorID), deviceType(deviceType), - deviceID(deviceID) {} - - bool operator==(const KeyTy &key) const { - return key == std::make_tuple(version, revision, extensions, vendorID, - deviceType, deviceID, capabilities); - } - - static TargetEnvAttributeStorage * - construct(AttributeStorageAllocator &allocator, const KeyTy &key) { - return new (allocator.allocate()) - TargetEnvAttributeStorage(std::get<0>(key), std::get<1>(key), - std::get<2>(key), std::get<3>(key), - std::get<4>(key), std::get<5>(key), - std::get<6>(key)); - } - - Attribute version; - Attribute revision; - Attribute extensions; - Attribute capabilities; - spirv::Vendor vendorID; - spirv::DeviceType deviceType; - uint32_t deviceID; -}; -} // namespace detail - -TargetEnvAttr TargetEnvAttr::get(Vulkan::Version version, uint32_t revision, - ArrayRef extensions, - spirv::Vendor vendorID, - spirv::DeviceType deviceType, - uint32_t deviceID, - CapabilitiesAttr capabilities) { - mlir::Builder builder(capabilities.getContext()); - llvm::SmallVector extAttrs; - extAttrs.reserve(extensions.size()); - for (auto ext : extensions) { - extAttrs.push_back(ExtensionAttr::get(builder.getContext(), ext)); - } - return get(builder.getI32IntegerAttr(static_cast(version)), - builder.getI32IntegerAttr(revision), - builder.getArrayAttr(extAttrs), vendorID, deviceType, deviceID, - capabilities); -} - -TargetEnvAttr TargetEnvAttr::get(IntegerAttr version, IntegerAttr revision, - ArrayAttr extensions, spirv::Vendor vendorID, - spirv::DeviceType deviceType, - uint32_t deviceID, - CapabilitiesAttr capabilities) { - assert(version && revision && extensions && capabilities); - MLIRContext *context = version.getContext(); - return Base::get(context, version, revision, extensions, vendorID, deviceType, - deviceID, capabilities); -} - -StringRef TargetEnvAttr::getKindName() { return "target_env"; } - -Version TargetEnvAttr::getVersion() { - return static_cast( - llvm::cast(getImpl()->version).getValue().getZExtValue()); -} - -unsigned TargetEnvAttr::getRevision() { - return llvm::cast(getImpl()->revision).getValue().getZExtValue(); -} - -TargetEnvAttr::ext_iterator::ext_iterator(ArrayAttr::iterator it) - : llvm::mapped_iterator( - it, [](Attribute attr) { - return llvm::cast(attr).getValue(); - }) {} - -TargetEnvAttr::ext_range TargetEnvAttr::getExtensions() { - auto range = getExtensionsAttr().getValue(); - return {ext_iterator(range.begin()), ext_iterator(range.end())}; -} - -ArrayAttr TargetEnvAttr::getExtensionsAttr() { - return llvm::cast(getImpl()->extensions); -} - -spirv::Vendor TargetEnvAttr::getVendorID() { return getImpl()->vendorID; } - -spirv::DeviceType TargetEnvAttr::getDeviceType() { - return getImpl()->deviceType; -} - -uint32_t TargetEnvAttr::getDeviceID() { return getImpl()->deviceID; } - -CapabilitiesAttr TargetEnvAttr::getCapabilitiesAttr() { - return llvm::cast(getImpl()->capabilities); -} - -LogicalResult -TargetEnvAttr::verify(function_ref emitError, - IntegerAttr version, IntegerAttr revision, - ArrayAttr extensions, spirv::Vendor /*vendorID*/, - spirv::DeviceType /*deviceType*/, uint32_t /*deviceID*/, - CapabilitiesAttr capabilities) { - if (!version.getType().isInteger(32)) - return emitError() << "expected 32-bit integer for version"; - - if (!revision.getType().isInteger(32)) - return emitError() << "expected 32-bit integer for revision"; - - return success(); -} - -//===----------------------------------------------------------------------===// -// Attribute Parsing -//===----------------------------------------------------------------------===// - -namespace { - -/// Parses a comma-separated list of keywords, invokes `processKeyword` on each -/// of the parsed keyword, and returns failure if any error occurs. -ParseResult parseKeywordList( - DialectAsmParser &parser, - function_ref processKeyword) { - if (parser.parseLSquare()) - return failure(); - - // Special case for empty list. - if (succeeded(parser.parseOptionalRSquare())) - return success(); - - // Keep parsing the keyword and an optional comma following it. If the comma - // is successfully parsed, then we have more keywords to parse. - do { - auto loc = parser.getCurrentLocation(); - StringRef keyword; - if (parser.parseKeyword(&keyword) || failed(processKeyword(loc, keyword))) - return failure(); - } while (succeeded(parser.parseOptionalComma())); - - if (parser.parseRSquare()) - return failure(); - - return success(); -} - -/// Parses a TargetEnvAttr. -Attribute parseTargetAttr(DialectAsmParser &parser) { - if (parser.parseLess()) - return {}; - - Builder &builder = parser.getBuilder(); - - IntegerAttr versionAttr; - { - auto loc = parser.getCurrentLocation(); - StringRef version; - if (parser.parseKeyword(&version) || parser.parseComma()) - return {}; - - if (auto versionSymbol = symbolizeVersion(version)) { - versionAttr = - builder.getI32IntegerAttr(static_cast(*versionSymbol)); - } else { - parser.emitError(loc, "unknown Vulkan version: ") << version; - return {}; - } - } - - IntegerAttr revisionAttr; - { - unsigned revision = 0; - // TODO(antiagainst): it would be nice to parse rN instad of r(N). - if (parser.parseKeyword("r") || parser.parseLParen() || - parser.parseInteger(revision) || parser.parseRParen() || - parser.parseComma()) - return {}; - revisionAttr = builder.getI32IntegerAttr(revision); - } - - ArrayAttr extensionsAttr; - { - SmallVector extensions; - llvm::SMLoc errorloc; - StringRef errorKeyword; - - MLIRContext *context = parser.getContext(); - auto processExtension = [&](llvm::SMLoc loc, StringRef extension) { - if (std::optional symbol = symbolizeExtension(extension)) { - extensions.push_back(ExtensionAttr::get(context, *symbol)); - return success(); - } - return errorloc = loc, errorKeyword = extension, failure(); - }; - if (parseKeywordList(parser, processExtension) || parser.parseComma()) { - if (!errorKeyword.empty()) - parser.emitError(errorloc, "unknown Vulkan extension: ") - << errorKeyword; - return {}; - } - - extensionsAttr = builder.getArrayAttr(extensions); - } - - // Parse vendor:device-type[:device-id] - spirv::Vendor vendorID = spirv::Vendor::Unknown; - spirv::DeviceType deviceType = spirv::DeviceType::Unknown; - uint32_t deviceID = spirv::TargetEnvAttr::kUnknownDeviceID; - { - auto loc = parser.getCurrentLocation(); - StringRef vendorStr; - if (parser.parseKeyword(&vendorStr)) - return {}; - if (auto vendorSymbol = spirv::symbolizeVendor(vendorStr)) { - vendorID = *vendorSymbol; - } else { - parser.emitError(loc, "unknown vendor: ") << vendorStr; - } - - loc = parser.getCurrentLocation(); - StringRef deviceTypeStr; - if (parser.parseColon() || parser.parseKeyword(&deviceTypeStr)) - return {}; - if (auto deviceTypeSymbol = spirv::symbolizeDeviceType(deviceTypeStr)) { - deviceType = *deviceTypeSymbol; - } else { - parser.emitError(loc, "unknown device type: ") << deviceTypeStr; - } - - loc = parser.getCurrentLocation(); - if (succeeded(parser.parseOptionalColon())) { - if (parser.parseInteger(deviceID)) - return {}; - } - - if (parser.parseComma()) - return {}; - } - - CapabilitiesAttr capabilities; - if (parser.parseAttribute(capabilities)) - return {}; - - if (parser.parseGreater()) - return {}; - - return TargetEnvAttr::get(versionAttr, revisionAttr, extensionsAttr, vendorID, - deviceType, deviceID, capabilities); -} -} // namespace - -Attribute VulkanDialect::parseAttribute(DialectAsmParser &parser, - Type type) const { - // Vulkan attributes do not have type. - if (type) { - parser.emitError(parser.getNameLoc(), "unexpected type"); - return {}; - } - - // Parse the kind keyword first. - StringRef attrKind; - Attribute attr; - OptionalParseResult result = - generatedAttributeParser(parser, &attrKind, type, attr); - if (result.has_value()) { - if (failed(result.value())) - return {}; - return attr; - } - - if (attrKind == TargetEnvAttr::getKindName()) - return parseTargetAttr(parser); - - parser.emitError(parser.getNameLoc(), "unknown Vulkan attriubte kind: ") - << attrKind; - return {}; -} - -//===----------------------------------------------------------------------===// -// Attribute Printing -//===----------------------------------------------------------------------===// - -namespace { -void print(TargetEnvAttr targetEnv, DialectAsmPrinter &printer) { - auto &os = printer.getStream(); - printer << TargetEnvAttr::getKindName() << "<" - << stringifyVersion(targetEnv.getVersion()) << ", r(" - << targetEnv.getRevision() << "), ["; - interleaveComma(targetEnv.getExtensions(), os, - [&](Extension ext) { os << stringifyExtension(ext); }); - printer << "], " << spirv::stringifyVendor(targetEnv.getVendorID()); - printer << ":" << spirv::stringifyDeviceType(targetEnv.getDeviceType()); - auto deviceID = targetEnv.getDeviceID(); - if (deviceID != spirv::TargetEnvAttr::kUnknownDeviceID) { - printer << ":" << targetEnv.getDeviceID(); - } - printer << ", " << targetEnv.getCapabilitiesAttr() << ">"; -} -} // namespace - -void VulkanDialect::printAttribute(Attribute attr, - DialectAsmPrinter &printer) const { - if (succeeded(generatedAttributePrinter(attr, printer))) - return; - - if (auto targetEnv = llvm::dyn_cast(attr)) - return print(targetEnv, printer); - - assert(false && "unhandled Vulkan attribute kind"); -} - -//===----------------------------------------------------------------------===// -// Registration -//===----------------------------------------------------------------------===// - -void VulkanDialect::registerAttributes() { - addAttributes(); -} - -} // namespace mlir::iree_compiler::IREE::Vulkan diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanAttributes.h b/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanAttributes.h deleted file mode 100644 index 1175db6bad6c..000000000000 --- a/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanAttributes.h +++ /dev/null @@ -1,89 +0,0 @@ -// Copyright 2020 The IREE Authors -// -// Licensed under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception - -#ifndef IREE_COMPILER_DIALECT_VULKAN_IR_VULKANATTRIBUTES_H_ -#define IREE_COMPILER_DIALECT_VULKAN_IR_VULKANATTRIBUTES_H_ - -#include "iree/compiler/Dialect/Vulkan/IR/VulkanTypes.h" -#include "mlir/Dialect/SPIRV/IR/SPIRVTypes.h" -#include "mlir/IR/BuiltinAttributes.h" - -#define GET_ATTRDEF_CLASSES -#include "iree/compiler/Dialect/Vulkan/IR/VulkanAttributes.h.inc" // IWYU pragma: export - -namespace mlir::iree_compiler::IREE::Vulkan { - -namespace detail { -struct TargetEnvAttributeStorage; -} // namespace detail - -/// An attribute that specifies the target version, supported extensions, and -/// resource limits. These information describles a Vulkan target environment. -class TargetEnvAttr - : public Attribute::AttrBase { -public: - using Base::Base; - - static constexpr StringLiteral name = "vk.target_env"; - - /// Gets a TargetEnvAttr instance. - // TODO(antiagainst): support other physical device core properties, physical - // device core features and per-extension features. - static TargetEnvAttr get(Version version, uint32_t revision, - ArrayRef extensions, - spirv::Vendor vendorID, spirv::DeviceType deviceType, - uint32_t deviceID, CapabilitiesAttr capabilities); - static TargetEnvAttr get(IntegerAttr version, IntegerAttr revision, - ArrayAttr extensions, spirv::Vendor vendorID, - spirv::DeviceType deviceType, uint32_t deviceID, - CapabilitiesAttr capabilities); - - /// Returns the attribute kind's name (without the 'vk.' prefix). - static StringRef getKindName(); - - /// Returns the target Vulkan version; e.g., for 1.1.120, it should be V_1_1. - Version getVersion(); - - /// Returns the target Vulkan revision; e.g., for 1.1.120, it should be 120. - unsigned getRevision(); - - struct ext_iterator final - : public llvm::mapped_iterator { - explicit ext_iterator(ArrayAttr::iterator it); - }; - using ext_range = llvm::iterator_range; - - /// Returns the target Vulkan instance and device extensions. - ext_range getExtensions(); - /// Returns the target Vulkan instance and device extensions as an string - /// array attribute. - ArrayAttr getExtensionsAttr(); - - /// Returns the vendor ID. - spirv::Vendor getVendorID(); - - /// Returns the device type. - spirv::DeviceType getDeviceType(); - - /// Returns the device ID. - uint32_t getDeviceID(); - - /// Returns the dictionary attribute containing various Vulkan capabilities - /// bits. - CapabilitiesAttr getCapabilitiesAttr(); - - static LogicalResult verify(function_ref emitError, - IntegerAttr version, IntegerAttr revision, - ArrayAttr extensions, spirv::Vendor vendorID, - spirv::DeviceType deviceType, uint32_t deviceID, - CapabilitiesAttr capabilities); -}; - -} // namespace mlir::iree_compiler::IREE::Vulkan - -#endif // IREE_COMPILER_DIALECT_VULKAN_IR_VULKANATTRIBUTES_H_ diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanAttributes.td b/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanAttributes.td deleted file mode 100644 index fcd0ccf65c04..000000000000 --- a/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanAttributes.td +++ /dev/null @@ -1,134 +0,0 @@ -// Copyright 2020 The IREE Authors -// -// Licensed under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception - -#ifndef IREE_DIALECT_VULKAN_VULKANATTRIBUTES -#define IREE_DIALECT_VULKAN_VULKANATTRIBUTES - -include "iree/compiler/Dialect/Vulkan/IR/VulkanBase.td" - -class VK_Attr - : AttrDef { - let mnemonic = attrMnemonic; - let assemblyFormat = "`<` struct(params) `>`"; -} - -// Attribute that can be used to specify the configuration of the -// cooperative matrix multiply instructions supported by the target -// device. This corresponds to `VkCooperativeMatrixPropertiesKHR` structure: -// https://registry.khronos.org/vulkan/specs/1.3-extensions/man/html/VkCooperativeMatrixPropertiesKHR.html -def VK_CooperativeMatrixPropertiesKHRAttr : - VK_Attr<"CooperativeMatrixPropertiesKHR", "coop_matrix_props"> { - let parameters = (ins - "uint32_t":$mSize, - "uint32_t":$nSize, - "uint32_t":$kSize, - "::mlir::Type":$aType, - "::mlir::Type":$bType, - "::mlir::Type":$cType, - "::mlir::Type":$resultType, - "bool":$accSat, - "::mlir::iree_compiler::IREE::Vulkan::ScopeKHRAttr":$scope - ); -} - -// TODO(antiagainst): consider auto-generating this file (or part of it) from -// vk.xml: -// https://raw.githubusercontent.com/KhronosGroup/Vulkan-Docs/main/xml/vk.xml - -// Dictionary attribute containing various Vulkan capability bits. This is -// aggregated from various Vulkan properties, limits, features from the spec. -// -// Note that we are using UnitAttr for booleans to allow omitting to mean false. -// TODO(antiagainst): support DefaultValuedAttr in StrucctAttr to allow -// specifying defaults for non-boolean fields. -def VK_CapabilitiesAttr : VK_Attr<"Capabilities", "caps"> { - let parameters = (ins - // Core Vulkan 1.0 physical device properties. - // - // This corresponds to the `VkPhysicalDeviceProperties` structure: - // https://www.khronos.org/registry/vulkan/specs/1.2-extensions/man/html/VkPhysicalDeviceProperties.html - "int":$maxComputeSharedMemorySize, - "int":$maxComputeWorkGroupInvocations, - "::mlir::DenseIntElementsAttr":$maxComputeWorkGroupSize, - - // Core Vulkan 1.0 physical device features. - // - // This corresponds to the `VkPhysicalDeviceFeatures` structure: - // https://www.khronos.org/registry/vulkan/specs/1.2-extensions/man/html/VkPhysicalDeviceFeatures.html - OptionalParameter<"::mlir::UnitAttr">:$shaderFloat64, - OptionalParameter<"::mlir::UnitAttr">:$shaderInt16, - OptionalParameter<"::mlir::UnitAttr">:$shaderInt64, - - // Core Vulkan 1.1 physical device subgroup properties. - // - // This corresponds to the `VkPhysicalDeviceSubgroupProperties` structure: - // https://www.khronos.org/registry/vulkan/specs/1.2-extensions/man/html/VkPhysicalDeviceSubgroupProperties.html - - // TODO(antiagainst): StructAttr does not actually support attribute kinds - // that are derived from IntegerAttr well. So the nice parsing/printing for - // VK_SubgroupFeatureAttr does not really kick in here. We need to enhance - // upstream MLIR. - "::mlir::iree_compiler::IREE::Vulkan::SubgroupFeatureAttr":$subgroupFeatures, - "int":$subgroupSize, - - // VK_EXT_subgroup_size_control features. - // - // This corresponds to the `VkPhysicalDeviceSubgroupSizeControlProperties` structure: - // https://registry.khronos.org/vulkan/specs/1.3-extensions/man/html/VkPhysicalDeviceSubgroupSizeControlPropertiesEXT.html - OptionalParameter<"::std::optional">:$minSubgroupSize, - OptionalParameter<"::std::optional">:$maxSubgroupSize, - - // VK_KHR_16bit_storage features. - // - // This corresponds to the `VkPhysicalDevice16BitStorageFeatures` structure: - // https://www.khronos.org/registry/vulkan/specs/1.2-extensions/man/html/VkPhysicalDevice16BitStorageFeatures.html - OptionalParameter<"::mlir::UnitAttr">:$storageBuffer16BitAccess, - OptionalParameter<"::mlir::UnitAttr">:$storagePushConstant16, - OptionalParameter<"::mlir::UnitAttr">:$uniformAndStorageBuffer16BitAccess, - - // VK_KHR_8bit_storage features. - // - // This corresponds to the `VkPhysicalDevice8BitStorageFeatures` structure: - // https://www.khronos.org/registry/vulkan/specs/1.2-extensions/man/html/VkPhysicalDevice8BitStorageFeatures.html - OptionalParameter<"::mlir::UnitAttr">:$storageBuffer8BitAccess, - OptionalParameter<"::mlir::UnitAttr">:$storagePushConstant8, - OptionalParameter<"::mlir::UnitAttr">:$uniformAndStorageBuffer8BitAccess, - - // VK_KHR_device_buffer_address features. - // This corresponds to the only capability implied by the extensions: - // https://registry.khronos.org/vulkan/specs/1.3-extensions/man/html/VK_KHR_buffer_device_address.html#_new_spir_v_capabilities - OptionalParameter<"::mlir::UnitAttr">:$physicalDeviceBufferAddresses, - - // VK_KHR_shader_float16_int8 features. - // - // This corresponds to the `VkPhysicalDeviceShaderFloat16Int8Features` - // structure: - // https://www.khronos.org/registry/vulkan/specs/1.2-extensions/man/html/VkPhysicalDeviceShaderFloat16Int8Features.html - OptionalParameter<"::mlir::UnitAttr">:$shaderFloat16, - OptionalParameter<"::mlir::UnitAttr">:$shaderInt8, - - // VK_KHR_shader_integer_dot_product features. - // - // This corresponds to the `VkPhysicalDeviceShaderIntegerDotProductFeatures` - // structure: - // https://registry.khronos.org/vulkan/specs/1.3-extensions/man/html/VkPhysicalDeviceShaderIntegerDotProductFeaturesKHR.html - OptionalParameter<"::mlir::UnitAttr">:$shaderIntegerDotProduct, - - // VK_KHR_variable_pointers features. - // This corresponds to the `VkPhysicalDeviceVariablePointersFeatures` - // structure: - // https://www.khronos.org/registry/vulkan/specs/1.2-extensions/man/html/VkPhysicalDeviceVariablePointersFeatures.html - OptionalParameter<"::mlir::UnitAttr">:$variablePointersStorageBuffer, - OptionalParameter<"::mlir::UnitAttr">:$variablePointers, - - // VkCooperativeMatrixPropertiesKHR features. - // This corresponds to `VkCoooperativeMatrixPropertiesKHR` structure: - // https://registry.khronos.org/vulkan/specs/1.3-extensions/man/html/VK_KHR_cooperative_matrix.html - DefaultValuedParameter<"ArrayAttr", "nullptr">:$cooperativeMatrixPropertiesKHR - ); -} - -#endif // IREE_DIALECT_VULKAN_VULKANATTRIBUTES diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanBase.td b/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanBase.td deleted file mode 100644 index c25611181d26..000000000000 --- a/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanBase.td +++ /dev/null @@ -1,199 +0,0 @@ -// Copyright 2020 The IREE Authors -// -// Licensed under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception - -#ifndef IREE_DIALECT_VULKAN_BASE -#define IREE_DIALECT_VULKAN_BASE - -include "mlir/IR/OpBase.td" -include "mlir/IR/EnumAttr.td" - -//===----------------------------------------------------------------------===// -// Vulkan dialect definition -//===----------------------------------------------------------------------===// - -def VK_Dialect : Dialect { - let name = "vk"; - let cppNamespace = "::mlir::iree_compiler::IREE::Vulkan"; - - let summary = "The Vulkan dialect in IREE"; - let description = [{ - Vulkan is a new generation graphics and compute API that provides - high-efficiency, cross-platform access to modern GPUs used in a wide - variety of devices from PCs and consoles to mobile phones and embedded - platforms. See https://www.khronos.org/vulkan for more details regarding - Vulkan itself. - - This is not a full-fledged Vulkan dialect that models common Vulkan concepts - in intermediate representation to be amenable to compiler analysis and - transformation. IREE has the HAL dialect for that purpose. Instead, this - dialect contains useful utilities for targeting Vulkan both in CodeGen and - runtime. - }]; -} - -//===----------------------------------------------------------------------===// -// Utility definitions -//===----------------------------------------------------------------------===// - -// A predicate that checks whether `$_self` is a known enum case for the -// enum class with `name`. -class VK_IsKnownBitEnumCaseFor : - CPred<"::mlir::iree_compiler::IREE::Vulkan::symbolize" # name # "(" - "cast($_self).getValue().getZExtValue()).hasValue()">; -class VK_IsKnownIntEnumCaseFor : - CPred<"::mlir::iree_compiler::IREE::Vulkan::symbolize" # name # "(" - "cast($_self).getValue().getZExtValue()).hasValue()">; - -// Wrapper over base I32BitEnumAttr to set common fields. -class VK_BitEnumAttr cases> : - I32BitEnumAttr { - let predicate = And<[I32Attr.predicate, VK_IsKnownBitEnumCaseFor]>; - let cppNamespace = "::mlir::iree_compiler::IREE::Vulkan"; -} - -class VK_I32Enum cases> : - I32EnumAttr { - let predicate = And<[I32Attr.predicate, VK_IsKnownIntEnumCaseFor]>; - let cppNamespace = "::mlir::iree_compiler::IREE::Vulkan"; -} - -class VK_I32EnumAttr cases> : - EnumAttr, mnemonic> { - let cppNamespace = "::mlir::iree_compiler::IREE::Vulkan"; - let assemblyFormat = "`<` $value `>`"; -} - -//===----------------------------------------------------------------------===// -// Target environment -//===----------------------------------------------------------------------===// - -def VK_V_1_0 : I32EnumAttrCase<"V_1_0", 0, "v1.0">; -def VK_V_1_1 : I32EnumAttrCase<"V_1_1", 1, "v1.1">; -def VK_V_1_2 : I32EnumAttrCase<"V_1_2", 2, "v1.2">; -def VK_V_1_3 : I32EnumAttrCase<"V_1_3", 3, "v1.3">; - -def VK_VersionAttr : VK_I32Enum<"Version", "valid Vulkan version", [ - VK_V_1_0, VK_V_1_1, VK_V_1_2, VK_V_1_3 -]>; - -def VK_KHR_16bit_storage : I32EnumAttrCase<"VK_KHR_16bit_storage", 0>; -def VK_KHR_8bit_storage : I32EnumAttrCase<"VK_KHR_8bit_storage", 1>; -def VK_KHR_shader_float16_int8 : I32EnumAttrCase<"VK_KHR_shader_float16_int8", 2>; -def VK_KHR_shader_integer_dot_product : I32EnumAttrCase<"VK_KHR_shader_integer_dot_product", 3>; -def VK_KHR_spirv_1_4 : I32EnumAttrCase<"VK_KHR_spirv_1_4", 4>; -def VK_KHR_storage_buffer_storage_class : I32EnumAttrCase<"VK_KHR_storage_buffer_storage_class", 5>; -def VK_KHR_variable_pointers: I32EnumAttrCase<"VK_KHR_variable_pointers", 6>; -def VK_EXT_subgroup_size_control : I32EnumAttrCase<"VK_EXT_subgroup_size_control", 7>; -def VK_KHR_cooperative_matrix : I32EnumAttrCase<"VK_KHR_cooperative_matrix", 8>; -def VK_KHR_buffer_device_address : I32EnumAttrCase<"VK_KHR_buffer_device_address", 9>; - -def VK_ExtensionAttr : - VK_I32EnumAttr<"Extension", "supported Vulkan extension", "extension", [ - VK_KHR_16bit_storage, VK_KHR_8bit_storage, VK_KHR_shader_float16_int8, - VK_KHR_shader_integer_dot_product, VK_KHR_spirv_1_4, - VK_KHR_storage_buffer_storage_class, VK_KHR_variable_pointers, - VK_EXT_subgroup_size_control, VK_KHR_cooperative_matrix, - VK_KHR_buffer_device_address - ]>; - -//===----------------------------------------------------------------------===// -// Target triple -//===----------------------------------------------------------------------===// - -def VK_TTA_Unknown : I32EnumAttrCase<"Unknown", 0, "unknown">; -// Software emulated GPU -def VK_TTA_CPU : I32EnumAttrCase<"CPU", 1, "cpu">; -// AMD GPU -def VK_TTA_RDNAv1 : I32EnumAttrCase<"AMD_RDNAv1", 100, "rdna1">; -def VK_TTA_RDNAv2 : I32EnumAttrCase<"AMD_RDNAv2", 101, "rdna2">; -def VK_TTA_RDNAv3 : I32EnumAttrCase<"AMD_RDNAv3", 102, "rdna3">; -// Apple Silicon GPU -def VK_TTA_M1 : I32EnumAttrCase<"Apple_M1", 200, "m1">; -// ARM Mali GPU -def VK_TTA_Valhall : I32EnumAttrCase<"ARM_Valhall", 300, "valhall">; -// NVIDIA GPU -def VK_TTA_Turing : I32EnumAttrCase<"NV_Turing", 400, "turing">; -def VK_TTA_Ampere : I32EnumAttrCase<"NV_Ampere", 401, "ampere">; -def VK_TTA_Pascal : I32EnumAttrCase<"NV_Pascal", 402, "pascal">; -// Qualcomm Adreno GPU -def VK_TTA_Adreno : I32EnumAttrCase<"QC_Adreno", 500, "adreno">; -// Intel ARC GPU -def VK_TTA_Arc : I32EnumAttrCase<"Intel_Arc", 600, "arc">; - -def VK_TargetArchAttr : VK_I32Enum< - "TargetTripleArch", "recognized target architecture", [ - VK_TTA_Unknown, VK_TTA_CPU, VK_TTA_RDNAv1, VK_TTA_RDNAv2, - VK_TTA_RDNAv3, VK_TTA_M1, VK_TTA_Valhall, VK_TTA_Turing, VK_TTA_Ampere, - VK_TTA_Pascal, VK_TTA_Adreno, VK_TTA_Arc, - ]>; - -def VK_TTP_Unknown : I32EnumAttrCase<"Unknown", 0, "unknown">; -// Qualcomm Adreno GPU -def VK_TTP_Adreno640 : I32EnumAttrCase<"Adreno_640", 100, "a640">; -def VK_TTP_Adreno650 : I32EnumAttrCase<"Adreno_650", 101, "a650">; -def VK_TTP_Adreno660 : I32EnumAttrCase<"Adreno_660", 102, "a660">; -// Software emulated GPU -def VK_TTP_SwiftShader : I32EnumAttrCase<"SwiftShader", 200, "swiftshader">; -// Translation layers -def VK_TTP_MoltenVK : I32EnumAttrCase<"MoltenVK", 300, "moltenvk">; - -def VK_TargetProductAttr : VK_I32Enum< - "TargetTripleProduct", "recognized target product", [ - VK_TTP_Unknown, VK_TTP_Adreno650, VK_TTP_Adreno660, VK_TTP_SwiftShader, - VK_TTP_MoltenVK, - ]>; - -def VK_TTOS_Unknown : I32EnumAttrCase<"Unknown", 0, "unknown">; -def VK_TTOS_Linux : I32EnumAttrCase<"Linux", 1, "linux">; -def VK_TTOS_iOS : I32EnumAttrCase<"iOS", 2, "iOS">; -def VK_TTOS_macOS : I32EnumAttrCase<"macOS", 3, "macos">; -def VK_TTOS_Windows : I32EnumAttrCase<"Windows", 4, "windows">; -// API Level 30 => Android 11 -def VK_TTOS_Android30 : I32EnumAttrCase<"Android30", 5, "android30">; -// API Level 31 => Android 12 -def VK_TTOS_Android31 : I32EnumAttrCase<"Android31", 6, "android31">; - -def VK_TargetOSAttr : VK_I32Enum< - "TargetTripleOS", "recognized target operating system", [ - VK_TTOS_Unknown, VK_TTOS_Linux, VK_TTOS_iOS, VK_TTOS_macOS, - VK_TTOS_Windows, VK_TTOS_Android30, VK_TTOS_Android31, - ]>; - -//===----------------------------------------------------------------------===// -// Subgroup features -//===----------------------------------------------------------------------===// - -def VK_SF_Basic : I32BitEnumAttrCase<"Basic", 0x001>; -def VK_SF_Vote : I32BitEnumAttrCase<"Vote", 0x002>; -def VK_SF_Arithmetic : I32BitEnumAttrCase<"Arithmetic", 0x004>; -def VK_SF_Ballot : I32BitEnumAttrCase<"Ballot", 0x008>; -def VK_SF_Shuffle : I32BitEnumAttrCase<"Shuffle", 0x010>; -def VK_SF_ShuffleRelative : I32BitEnumAttrCase<"ShuffleRelative", 0x020>; -def VK_SF_Clustered : I32BitEnumAttrCase<"Clustered", 0x040>; -def VK_SF_Quad : I32BitEnumAttrCase<"Quad", 0x080>; -def VK_SF_PartitionedNV : I32BitEnumAttrCase<"PartitionedNV", 0x100>; - -def VK_SubgroupFeatureAttr : VK_BitEnumAttr< - "SubgroupFeature", "supported Vulkan subgroup feature", [ - VK_SF_Basic, VK_SF_Vote, VK_SF_Arithmetic, VK_SF_Ballot, VK_SF_Shuffle, - VK_SF_ShuffleRelative, VK_SF_Clustered, VK_SF_Quad, VK_SF_PartitionedNV - ]>; - -// Matches VkScopeKHR and VkScopeNV. -def VK_SKHR_Device : I32EnumAttrCase<"Device", 1>; -def VK_SKHR_Workgroup : I32EnumAttrCase<"Workgroup", 2>; -def VK_SKHR_Subgroup : I32EnumAttrCase<"Subgroup", 3>; -def VK_SKHR_QueueFamily : I32EnumAttrCase<"QueueFamily", 5>; - -def VK_ScopeKHR_Attr : - VK_I32EnumAttr<"ScopeKHR", "valid VkScopeKHR", "scope", [ - VK_SKHR_Device, VK_SKHR_Workgroup, VK_SKHR_Subgroup, - VK_SKHR_QueueFamily - ]>; - -#endif // IREE_DIALECT_VULKAN_BASE diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanDialect.cpp b/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanDialect.cpp deleted file mode 100644 index 2e78febb1eb7..000000000000 --- a/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanDialect.cpp +++ /dev/null @@ -1,18 +0,0 @@ -// Copyright 2020 The IREE Authors -// -// Licensed under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception - -#include "iree/compiler/Dialect/Vulkan/IR/VulkanDialect.h" - -#include "iree/compiler/Dialect/Vulkan/IR/VulkanAttributes.h" - -namespace mlir::iree_compiler::IREE::Vulkan { - -VulkanDialect::VulkanDialect(MLIRContext *context) - : Dialect(getDialectNamespace(), context, TypeID::get()) { - registerAttributes(); -} - -} // namespace mlir::iree_compiler::IREE::Vulkan diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanDialect.h b/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanDialect.h deleted file mode 100644 index 9cb3d010008f..000000000000 --- a/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanDialect.h +++ /dev/null @@ -1,37 +0,0 @@ -// Copyright 2020 The IREE Authors -// -// Licensed under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception - -#ifndef IREE_COMPILER_DIALECT_VULKAN_IR_VULKANDIALECT_H_ -#define IREE_COMPILER_DIALECT_VULKAN_IR_VULKANDIALECT_H_ - -#include "mlir/IR/Dialect.h" - -namespace mlir::iree_compiler::IREE::Vulkan { - -class VulkanDialect : public Dialect { -public: - explicit VulkanDialect(MLIRContext *context); - - static StringRef getDialectNamespace() { return "vk"; } - - //===--------------------------------------------------------------------===// - // Attribute - //===--------------------------------------------------------------------===// - - /// Parses an attribute registered to this dialect. - Attribute parseAttribute(DialectAsmParser &parser, Type type) const override; - - /// Prints an attribute registered to this dialect. - void printAttribute(Attribute, DialectAsmPrinter &printer) const override; - -private: - /// Register the attributes of this dialect. - void registerAttributes(); -}; - -} // namespace mlir::iree_compiler::IREE::Vulkan - -#endif // IREE_COMPILER_DIALECT_VULKAN_IR_VULKANDIALECT_H_ diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanTypes.cpp b/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanTypes.cpp deleted file mode 100644 index fc67767cb64a..000000000000 --- a/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanTypes.cpp +++ /dev/null @@ -1,13 +0,0 @@ -// Copyright 2020 The IREE Authors -// -// Licensed under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception - -#include "iree/compiler/Dialect/Vulkan/IR/VulkanTypes.h" - -#include "llvm/ADT/StringExtras.h" // IWYU pragma: keep - -// clang-format off: must be included after all LLVM/MLIR headers. -#include "iree/compiler/Dialect/Vulkan/IR/VulkanEnums.cpp.inc" // IWYU pragma: keep -// clang-format on diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanTypes.h b/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanTypes.h deleted file mode 100644 index 2422a85b171f..000000000000 --- a/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanTypes.h +++ /dev/null @@ -1,20 +0,0 @@ -// Copyright 2020 The IREE Authors -// -// Licensed under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception - -#ifndef IREE_COMPILER_DIALECT_VULKAN_IR_VULKANTYPES_H_ -#define IREE_COMPILER_DIALECT_VULKAN_IR_VULKANTYPES_H_ - -#include "llvm/ADT/DenseMapInfo.h" -#include "llvm/ADT/StringRef.h" -#include "llvm/ADT/StringSwitch.h" -#include "mlir/IR/BuiltinAttributes.h" -#include "mlir/IR/BuiltinTypes.h" - -// clang-format off: must be included after all LLVM/MLIR headers. -#include "iree/compiler/Dialect/Vulkan/IR/VulkanEnums.h.inc" // IWYU pragma: export -// clang-format on - -#endif // IREE_COMPILER_DIALECT_VULKAN_IR_VULKANTYPES_H_ diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/IR/test/BUILD.bazel b/compiler/src/iree/compiler/Dialect/Vulkan/IR/test/BUILD.bazel deleted file mode 100644 index bbddf7da3581..000000000000 --- a/compiler/src/iree/compiler/Dialect/Vulkan/IR/test/BUILD.bazel +++ /dev/null @@ -1,26 +0,0 @@ -# Copyright 2020 The IREE Authors -# -# Licensed under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception - -load("//build_tools/bazel:enforce_glob.bzl", "enforce_glob") -load("//build_tools/bazel:iree_lit_test.bzl", "iree_lit_test_suite") - -package( - features = ["layering_check"], - licenses = ["notice"], # Apache 2.0 -) - -iree_lit_test_suite( - name = "lit", - srcs = enforce_glob( - ["target_env.mlir"], - include = ["*.mlir"], - ), - cfg = "//compiler:lit.cfg.py", - tools = [ - "//tools:iree-opt", - "@llvm-project//llvm:FileCheck", - ], -) diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/IR/test/CMakeLists.txt b/compiler/src/iree/compiler/Dialect/Vulkan/IR/test/CMakeLists.txt deleted file mode 100644 index cebe847f2b4a..000000000000 --- a/compiler/src/iree/compiler/Dialect/Vulkan/IR/test/CMakeLists.txt +++ /dev/null @@ -1,23 +0,0 @@ -################################################################################ -# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from # -# compiler/src/iree/compiler/Dialect/Vulkan/IR/test/BUILD.bazel # -# # -# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary # -# CMake-only content. # -# # -# To disable autogeneration for this file entirely, delete this header. # -################################################################################ - -iree_add_all_subdirs() - -iree_lit_test_suite( - NAME - lit - SRCS - "target_env.mlir" - TOOLS - FileCheck - iree-opt -) - -### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ### diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/IR/test/target_env.mlir b/compiler/src/iree/compiler/Dialect/Vulkan/IR/test/target_env.mlir deleted file mode 100644 index 343f1aa749ee..000000000000 --- a/compiler/src/iree/compiler/Dialect/Vulkan/IR/test/target_env.mlir +++ /dev/null @@ -1,150 +0,0 @@ -// Test parsing and printing Vulkan target environment attribute. - -// RUN: iree-opt --allow-unregistered-dialect --split-input-file --verify-diagnostics %s | FileCheck %s - -"vk_configure_op"() { - // CHECK: #vk.target_env : vector<3xi32> - // CHECK-SAME: subgroupFeatures = 63 : i32, - // CHECK-SAME: subgroupSize = 4 - // CHECK-SAME: >> - target_env = #vk.target_env: vector<3xi32>, - subgroupFeatures = 63 : i32, - subgroupSize = 4 - >> -} : () -> () - -// ----- - -"vk_configure_op"() { - // CHECK: #vk.target_env - // CHECK-SAME: VK_KHR_cooperative_matrix - // CHECK-SAME: cooperativeMatrixPropertiesKHR = - // CHECK-SAME: #vk.coop_matrix_props> - // CHECK-SAME: #vk.coop_matrix_props> - target_env = - #vk.target_env : vector<3xi32>, - subgroupFeatures = 63: i32, subgroupSize = 32, - cooperativeMatrixPropertiesKHR = [ - #vk.coop_matrix_props< - mSize = 8, nSize = 8, kSize = 32, - aType = i8, bType = i8, cType = i32, resultType = i32, - accSat = false, scope = #vk.scope>, - #vk.coop_matrix_props< - mSize = 8, nSize = 8, kSize = 16, - aType = f16, bType = f16, cType = f16, resultType = f16, - accSat = false, scope = #vk.scope> - ] - >> -} : () -> () - -// ----- - -"vk_configure_op"() { - // CHECK: Qualcomm:IntegratedGPU:100925441 - // CHECK-SAME: shaderFloat64 - // CHECK-SAME: shaderInt16 - target_env = #vk.target_env: vector<3xi32>, - subgroupFeatures = 63: i32, - subgroupSize = 4, - shaderFloat64 = unit, shaderInt16 = unit - >> -} : () -> () - -// ----- - -"unknown_vulkan_version"() { - // expected-error @+1 {{unknown Vulkan version: v10.8}} - target_env = #vk.target_env: vector<3xi32> - >> -} : () -> () - -// ----- - -"unknown_vulkan_extension"() { - // expected-error @+1 {{unknown Vulkan extension: VK_KHR_something}} - target_env = #vk.target_env: vector<3xi32> - >> -} : () -> () - -// ----- - -"wrong_vendor_id"() { - // expected-error @+1 {{unknown vendor: AVendor}} - target_env = #vk.target_env: vector<3xi32>, - subgroupFeatures = 63: i32, - subgroupSize = 4 - >> -} : () -> () - -// ----- - -"wrong_device_type"() { - // expected-error @+1 {{unknown device type: ADeviceType}} - target_env = #vk.target_env: vector<3xi32>, - subgroupFeatures = 63: i32, - subgroupSize = 4 - >> -} : () -> () - -// ----- - -"missing_core_1_1_properties_field"() { - target_env = #vk.target_env> -} : () -> () - -// ----- - -"unknown_core_1_1_properties_field"() { - target_env = #vk.target_env: vector<3xi32>, - // expected-error @+1 {{duplicate or unknown struct parameter name: moreStuff}} - moreStuff = 8: i32 - >> -} : () -> () - -// ----- - -"wrong_subgroup_bit"() { - target_env = #vk.target_env: vector<3xi32>, - // expected-error @+2 {{invalid kind of attribute specified}} - // expected-error @+1 {{failed to parse VK_CapabilitiesAttr parameter 'subgroupFeatures'}} - subgroupFeatures = 0xffffffff: i32, - subgroupSize = 4 - >> -} : () -> () diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/BUILD.bazel b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/BUILD.bazel deleted file mode 100644 index cbbd06fe8b33..000000000000 --- a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/BUILD.bazel +++ /dev/null @@ -1,32 +0,0 @@ -# Copyright 2019 The IREE Authors -# -# Licensed under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception - -load("//build_tools/bazel:build_defs.oss.bzl", "iree_compiler_cc_library") - -package( - default_visibility = ["//visibility:public"], - features = ["layering_check"], - licenses = ["notice"], # Apache 2.0 -) - -iree_compiler_cc_library( - name = "Utils", - srcs = [ - "TargetEnvironment.cpp", - "TargetTriple.cpp", - ], - hdrs = [ - "TargetEnvironment.h", - "TargetTriple.h", - ], - deps = [ - "//compiler/src/iree/compiler/Dialect/Vulkan/IR", - "@llvm-project//llvm:Support", - "@llvm-project//mlir:IR", - "@llvm-project//mlir:SPIRVDialect", - "@llvm-project//mlir:Support", - ], -) diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/CMakeLists.txt b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/CMakeLists.txt deleted file mode 100644 index 8435767a6a4a..000000000000 --- a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/CMakeLists.txt +++ /dev/null @@ -1,31 +0,0 @@ -################################################################################ -# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from # -# compiler/src/iree/compiler/Dialect/Vulkan/Utils/BUILD.bazel # -# # -# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary # -# CMake-only content. # -# # -# To disable autogeneration for this file entirely, delete this header. # -################################################################################ - -iree_add_all_subdirs() - -iree_cc_library( - NAME - Utils - HDRS - "TargetEnvironment.h" - "TargetTriple.h" - SRCS - "TargetEnvironment.cpp" - "TargetTriple.cpp" - DEPS - LLVMSupport - MLIRIR - MLIRSPIRVDialect - MLIRSupport - iree::compiler::Dialect::Vulkan::IR - PUBLIC -) - -### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ### diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetEnvironment.cpp b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetEnvironment.cpp deleted file mode 100644 index bcf3b55b2c2c..000000000000 --- a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetEnvironment.cpp +++ /dev/null @@ -1,222 +0,0 @@ -// Copyright 2020 The IREE Authors -// -// Licensed under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception - -#include "iree/compiler/Dialect/Vulkan/Utils/TargetEnvironment.h" - -#include "iree/compiler/Dialect/Vulkan/IR/VulkanTypes.h" -#include "iree/compiler/Dialect/Vulkan/Utils/TargetTriple.h" -#include "llvm/ADT/STLExtras.h" -#include "mlir/Dialect/SPIRV/IR/SPIRVEnums.h" -#include "mlir/Dialect/SPIRV/IR/SPIRVTypes.h" -#include "mlir/IR/Builders.h" -#include "mlir/IR/BuiltinAttributes.h" - -namespace mlir::iree_compiler::IREE::Vulkan { - -namespace { - -/// Gets the corresponding SPIR-V version for the ggiven Vulkan target -/// environment. -spirv::Version convertVersion(Vulkan::TargetEnvAttr vkTargetEnv) { - // Special extension to enable SPIR-V 1.4. - const bool has14Ext = (llvm::is_contained(vkTargetEnv.getExtensions(), - Extension::VK_KHR_spirv_1_4)); - - switch (vkTargetEnv.getVersion()) { - case Version::V_1_0: - // Vulkan 1.0 only supports SPIR-V 1.0 by default. - return has14Ext ? spirv::Version::V_1_4 : spirv::Version::V_1_0; - case Version::V_1_1: - // Vulkan 1.1 supports up to SPIR-V 1.3 by default. - return has14Ext ? spirv::Version::V_1_4 : spirv::Version::V_1_3; - case Version::V_1_2: - // Vulkan 1.1 supports up to SPIR-V 1.5 by default. - return spirv::Version::V_1_5; - case Version::V_1_3: - // Vulkan 1.1 supports up to SPIR-V 1.6 by default. - return spirv::Version::V_1_6; - } - return spirv::Version::V_1_0; -} - -/// Gets the corresponding SPIR-V extensions for the given Vulkan target -/// environment. -void convertExtensions(Vulkan::TargetEnvAttr vkTargetEnv, - SmallVectorImpl &extensions) { - extensions.clear(); - - for (Extension ext : vkTargetEnv.getExtensions()) { - switch (ext) { - case Extension::VK_KHR_16bit_storage: - extensions.push_back(spirv::Extension::SPV_KHR_16bit_storage); - break; - case Extension::VK_KHR_8bit_storage: - extensions.push_back(spirv::Extension::SPV_KHR_8bit_storage); - break; - case Extension::VK_KHR_shader_float16_int8: - // This extension allows using certain SPIR-V capabilities. - break; - case Extension::VK_KHR_shader_integer_dot_product: - extensions.push_back(spirv::Extension::SPV_KHR_integer_dot_product); - break; - case Extension::VK_KHR_spirv_1_4: - // This extension only affects SPIR-V version. - break; - case Extension::VK_KHR_storage_buffer_storage_class: - extensions.push_back( - spirv::Extension::SPV_KHR_storage_buffer_storage_class); - break; - case Extension::VK_KHR_variable_pointers: - extensions.push_back(spirv::Extension::SPV_KHR_variable_pointers); - break; - case Extension::VK_EXT_subgroup_size_control: - // This extension allows specifying min/max subgroup size. - break; - case Extension::VK_KHR_cooperative_matrix: - extensions.push_back(spirv::Extension::SPV_KHR_cooperative_matrix); - break; - case Extension::VK_KHR_buffer_device_address: - extensions.push_back(spirv::Extension::SPV_KHR_physical_storage_buffer); - } - } -} - -/// Gets the corresponding SPIR-V capabilities for the given Vulkan target -/// environment. -void convertCapabilities(Vulkan::TargetEnvAttr vkTargetEnv, - SmallVectorImpl &capabilities) { - // Add unconditionally supported capabilities. - // Note that "Table 54. List of SPIR-V Capabilities and enabling features or - // extensions" in the Vulkan spec contains the full list. Right now omit those - // implicitly declared or not useful for us. - capabilities.assign({spirv::Capability::Shader}); - - auto vkCapabilities = vkTargetEnv.getCapabilitiesAttr(); - -#define MAP_PRIMITIVE_TYPE(type) \ - if (vkCapabilities.getShader##type()) \ - capabilities.push_back(spirv::Capability::type) - - MAP_PRIMITIVE_TYPE(Float64); - MAP_PRIMITIVE_TYPE(Float16); - MAP_PRIMITIVE_TYPE(Int64); - MAP_PRIMITIVE_TYPE(Int16); - MAP_PRIMITIVE_TYPE(Int8); -#undef MAP_PRIMITIVE_TYPE - -#define MAP_8_16_BIT_STORAGE(vkFeature, spvCap) \ - if (vkCapabilities.vkFeature()) \ - capabilities.push_back(spirv::Capability::spvCap) - - MAP_8_16_BIT_STORAGE(getStorageBuffer16BitAccess, StorageBuffer16BitAccess); - MAP_8_16_BIT_STORAGE(getUniformAndStorageBuffer16BitAccess, StorageUniform16); - MAP_8_16_BIT_STORAGE(getStoragePushConstant16, StoragePushConstant16); - MAP_8_16_BIT_STORAGE(getStorageBuffer8BitAccess, StorageBuffer8BitAccess); - MAP_8_16_BIT_STORAGE(getUniformAndStorageBuffer8BitAccess, - UniformAndStorageBuffer8BitAccess); - MAP_8_16_BIT_STORAGE(getStoragePushConstant8, StoragePushConstant8); -#undef MAP_8_16_BIT_STORAGE - - auto subgroupFeatures = vkCapabilities.getSubgroupFeatures().getValue(); - -#define MAP_SUBGROUP_FEATURE(featureBit) \ - if ((subgroupFeatures & SubgroupFeature::featureBit) == \ - SubgroupFeature::featureBit) \ - capabilities.push_back(spirv::Capability::GroupNonUniform##featureBit) - - if ((subgroupFeatures & SubgroupFeature::Basic) == SubgroupFeature::Basic) { - capabilities.push_back(spirv::Capability::GroupNonUniform); - } - MAP_SUBGROUP_FEATURE(Vote); - MAP_SUBGROUP_FEATURE(Arithmetic); - MAP_SUBGROUP_FEATURE(Ballot); - MAP_SUBGROUP_FEATURE(Shuffle); - MAP_SUBGROUP_FEATURE(ShuffleRelative); - MAP_SUBGROUP_FEATURE(Clustered); - MAP_SUBGROUP_FEATURE(Quad); - MAP_SUBGROUP_FEATURE(PartitionedNV); -#undef MAP_SUBGROUP_FEATURE - if (vkCapabilities.getPhysicalDeviceBufferAddresses()) { - capabilities.push_back(spirv::Capability::PhysicalStorageBufferAddresses); - } - if (vkCapabilities.getVariablePointers()) { - capabilities.push_back(spirv::Capability::VariablePointers); - } - if (vkCapabilities.getVariablePointersStorageBuffer()) { - capabilities.push_back(spirv::Capability::VariablePointersStorageBuffer); - } - if (vkCapabilities.getShaderIntegerDotProduct()) { - llvm::append_values(capabilities, spirv::Capability::DotProduct, - spirv::Capability::DotProductInputAll, - spirv::Capability::DotProductInput4x8BitPacked); - if (vkCapabilities.getShaderInt8()) { - capabilities.push_back(spirv::Capability::DotProductInput4x8Bit); - } - } - if (ArrayAttr attr = vkCapabilities.getCooperativeMatrixPropertiesKHR()) { - if (!attr.empty()) { - capabilities.push_back(spirv::Capability::CooperativeMatrixKHR); - } - } -} - -/// Gets the corresponding SPIR-V resource limits for the given Vulkan target -/// environment. -spirv::ResourceLimitsAttr -convertResourceLimits(Vulkan::TargetEnvAttr vkTargetEnv) { - MLIRContext *context = vkTargetEnv.getContext(); - Builder builder(context); - auto vkCapabilities = vkTargetEnv.getCapabilitiesAttr(); - SmallVector khrCoopAttrs; - if (ArrayAttr attr = vkCapabilities.getCooperativeMatrixPropertiesKHR()) { - for (auto props : - attr.getAsRange()) { - auto scope = static_cast(props.getScope().getValue()); - khrCoopAttrs.push_back(spirv::CooperativeMatrixPropertiesKHRAttr::get( - context, props.getMSize(), props.getNSize(), props.getKSize(), - props.getAType(), props.getBType(), props.getCType(), - props.getResultType(), props.getAccSat(), - spirv::ScopeAttr::get(context, scope))); - } - } - auto sizeValues = - vkCapabilities.getMaxComputeWorkGroupSize().getValues(); - SmallVector sizes; - sizes.insert(sizes.end(), sizeValues.begin(), sizeValues.end()); - return spirv::ResourceLimitsAttr::get( - context, vkCapabilities.getMaxComputeSharedMemorySize(), - vkCapabilities.getMaxComputeWorkGroupInvocations(), - builder.getI64ArrayAttr(sizes), vkCapabilities.getSubgroupSize(), - vkCapabilities.getMinSubgroupSize(), vkCapabilities.getMaxSubgroupSize(), - ArrayAttr::get(context, khrCoopAttrs), ArrayAttr{}); -} - -} // namespace - -Vulkan::TargetEnvAttr getTargetEnvForTriple(MLIRContext *context, - llvm::StringRef triple) { - return TargetTriple::get(triple.data()).getTargetEnv(context); -} - -spirv::TargetEnvAttr convertTargetEnv(Vulkan::TargetEnvAttr vkTargetEnv) { - auto spvVersion = convertVersion(vkTargetEnv); - - SmallVector spvExtensions; - convertExtensions(vkTargetEnv, spvExtensions); - - SmallVector spvCapabilities; - convertCapabilities(vkTargetEnv, spvCapabilities); - - auto spvLimits = convertResourceLimits(vkTargetEnv); - - auto triple = spirv::VerCapExtAttr::get( - spvVersion, spvCapabilities, spvExtensions, vkTargetEnv.getContext()); - return spirv::TargetEnvAttr::get( - triple, spvLimits, spirv::ClientAPI::Vulkan, vkTargetEnv.getVendorID(), - vkTargetEnv.getDeviceType(), vkTargetEnv.getDeviceID()); -} - -} // namespace mlir::iree_compiler::IREE::Vulkan diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetEnvironment.h b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetEnvironment.h deleted file mode 100644 index cc1d62aa9321..000000000000 --- a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetEnvironment.h +++ /dev/null @@ -1,36 +0,0 @@ -// Copyright 2020 The IREE Authors -// -// Licensed under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception - -#ifndef IREE_COMPILER_DIALECT_VULKAN_UTILS_TARGETENVIRONMENT_H_ -#define IREE_COMPILER_DIALECT_VULKAN_UTILS_TARGETENVIRONMENT_H_ - -#include "iree/compiler/Dialect/Vulkan/IR/VulkanAttributes.h" -#include "mlir/Dialect/SPIRV/IR/TargetAndABI.h" - -namespace mlir::iree_compiler::IREE::Vulkan { - -/// Returns the Vulkan target environment attribute for the given GPU triple. -Vulkan::TargetEnvAttr getTargetEnvForTriple(MLIRContext *context, - llvm::StringRef triple); - -/// Converts the given Vulkan target environment into the corresponding SPIR-V -/// target environment. -/// -/// Vulkan and SPIR-V are two different domains working closely. A Vulkan target -/// environment specifies the Vulkan version, extensions, features, and resource -/// limits queried from a Vulkan implementation. These properties typically have -/// corresponding SPIR-V bits, directly or indirectly. For example, by default, -/// Vulkan 1.0 supports SPIR-V 1.0 and Vulkan 1.1 supports up to SPIR-V 1.3. -/// If the VK_KHR_spirv_1_4 extension is available, then SPIR-V 1.4 can be used. -/// Similarly, if the VK_KHR_variable_pointers extension is available, then -/// the VariablePointersStorageBuffer capabilities on SPIR-V side can be -/// activated. The function handles the mapping relationship between tese two -/// domains. -spirv::TargetEnvAttr convertTargetEnv(Vulkan::TargetEnvAttr vkTargetEnv); - -} // namespace mlir::iree_compiler::IREE::Vulkan - -#endif // IREE_COMPILER_DIALECT_VULKAN_UTILS_TARGETENVIRONMENT_H_ diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetTriple.cpp b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetTriple.cpp deleted file mode 100644 index 9564bf7b60f7..000000000000 --- a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetTriple.cpp +++ /dev/null @@ -1,539 +0,0 @@ -// Copyright 2021 The IREE Authors -// -// Licensed under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception - -#include "iree/compiler/Dialect/Vulkan/Utils/TargetTriple.h" - -#include "iree/compiler/Dialect/Vulkan/IR/VulkanTypes.h" -#include "llvm/ADT/SmallVector.h" -#include "llvm/ADT/StringExtras.h" -#include "llvm/ADT/StringRef.h" -#include "llvm/Support/FormatVariadic.h" -#include "mlir/Dialect/SPIRV/IR/SPIRVAttributes.h" -#include "mlir/Dialect/SPIRV/IR/SPIRVEnums.h" -#include "mlir/IR/Attributes.h" -#include "mlir/IR/Builders.h" -#include "mlir/IR/BuiltinAttributes.h" -#include "mlir/IR/BuiltinTypes.h" - -namespace mlir::iree_compiler::IREE::Vulkan { - -namespace { - -/// Returns the GPU vendor for the given target `triple`. -spirv::Vendor getVendor(const TargetTriple &triple) { - switch (triple.getArch()) { - case TargetTripleArch::Unknown: - return spirv::Vendor::Unknown; - case TargetTripleArch::AMD_RDNAv1: - case TargetTripleArch::AMD_RDNAv2: - case TargetTripleArch::AMD_RDNAv3: - return spirv::Vendor::AMD; - case TargetTripleArch::ARM_Valhall: - return spirv::Vendor::ARM; - case TargetTripleArch::Apple_M1: - return spirv::Vendor::Apple; - case TargetTripleArch::Intel_Arc: - return spirv::Vendor::Intel; - case TargetTripleArch::NV_Turing: - case TargetTripleArch::NV_Ampere: - case TargetTripleArch::NV_Pascal: - return spirv::Vendor::NVIDIA; - case TargetTripleArch::QC_Adreno: - return spirv::Vendor::Qualcomm; - case TargetTripleArch::CPU: - switch (triple.getProduct()) { - case TargetTripleProduct::SwiftShader: - return spirv::Vendor::SwiftShader; - default: - return spirv::Vendor::Unknown; - } - default: - assert(false && "unhandled vendor"); - return spirv::Vendor::Unknown; - } -} - -/// Returns the GPU device type for the given target `triple`. -spirv::DeviceType getDeviceType(const TargetTriple &triple) { - switch (triple.getArch()) { - case TargetTripleArch::Unknown: - return spirv::DeviceType::Unknown; - case TargetTripleArch::CPU: - return spirv::DeviceType::CPU; - case TargetTripleArch::AMD_RDNAv1: - case TargetTripleArch::AMD_RDNAv2: - case TargetTripleArch::AMD_RDNAv3: - case TargetTripleArch::NV_Turing: - case TargetTripleArch::NV_Ampere: - case TargetTripleArch::NV_Pascal: - case TargetTripleArch::Intel_Arc: - return spirv::DeviceType::DiscreteGPU; - case TargetTripleArch::Apple_M1: - case TargetTripleArch::ARM_Valhall: - case TargetTripleArch::QC_Adreno: - return spirv::DeviceType::IntegratedGPU; - default: - assert(false && "unhandled device type"); - return spirv::DeviceType::Unknown; - } -} - -/// Returns the Vulkan version for the given target `triple`. -Vulkan::Version getVersion(const TargetTriple &triple) { - // Android 11/12 (API level 30/31) stays at Vulkan 1.1. - if (triple.getOS() == TargetTripleOS::Android30 || - triple.getOS() == TargetTripleOS::Android31) { - return Version::V_1_1; - } - - // SwiftShader and MoltenVK stays at Vulkan 1.1. - if (triple.getProduct() == TargetTripleProduct::SwiftShader || - triple.getProduct() == TargetTripleProduct::MoltenVK) { - return Version::V_1_1; - } - - // For unknown architecture, be conservative and use a reasonable lowest - // denominator. - if (triple.getArch() == TargetTripleArch::Unknown) { - return Version::V_1_1; - } - - return Version::V_1_3; -} - -/// Writes the Vulkan extensions supported by the given `triple` into -/// `extensions`. -/// -/// Note that this is an "approximation": Android compatibility will provide -/// some minimal guarantee but still different Android devices can have -/// different set of extensions, depending on the Android and GPU driver -/// version. The GPU triple is a handy way to specify the target but we cannot -/// encode all the information in the triple. -void getExtensions(const TargetTriple &triple, - llvm::SmallVectorImpl &extensions) { - // Mobile GPUs need to take Android version into consideration. - switch (triple.getArch()) { - case TargetTripleArch::Apple_M1: { - // Example: https://vulkan.gpuinfo.org/displayreport.php?id=14673 - return append_values(extensions, Extension::VK_KHR_16bit_storage, - Extension::VK_KHR_8bit_storage, - Extension::VK_KHR_shader_float16_int8, - Extension::VK_KHR_storage_buffer_storage_class, - Extension::VK_KHR_buffer_device_address, - Extension::VK_KHR_variable_pointers); - } - case TargetTripleArch::ARM_Valhall: { - // Example: https://vulkan.gpuinfo.org/displayreport.php?id=10312 - return append_values(extensions, Extension::VK_KHR_16bit_storage, - Extension::VK_KHR_8bit_storage, - Extension::VK_KHR_shader_float16_int8, - Extension::VK_KHR_shader_integer_dot_product, - Extension::VK_KHR_spirv_1_4, - Extension::VK_KHR_storage_buffer_storage_class, - Extension::VK_KHR_variable_pointers); - } - case TargetTripleArch::QC_Adreno: { - // Example: https://vulkan.gpuinfo.org/displayreport.php?id=10983 (11) - // Example: https://vulkan.gpuinfo.org/displayreport.php?id=16312 (12) - append_values(extensions, Extension::VK_KHR_16bit_storage, - Extension::VK_KHR_shader_float16_int8, - Extension::VK_KHR_spirv_1_4, - Extension::VK_KHR_storage_buffer_storage_class, - Extension::VK_KHR_variable_pointers); - if (triple.getOS() == TargetTripleOS::Android31) { - extensions.push_back(Extension::VK_KHR_8bit_storage); - } - return; - } - default: - break; - } - - // SwiftShader is very limited regarding functionalities. - if (getVendor(triple) == spirv::Vendor::SwiftShader) { - extensions.push_back(Extension::VK_KHR_storage_buffer_storage_class); - return; - } - - // For unknown architecture, be conservative and use a reasonable lowest - // denominator. - if (triple.getArch() == TargetTripleArch::Unknown) { - // The following extensions have 90%+ device coverage from - // https://vulkan.gpuinfo.org/listextensions.php. - const Extension list[] = { - Extension::VK_KHR_storage_buffer_storage_class, - Extension::VK_KHR_variable_pointers, - }; - return append_range(extensions, list); - } - - llvm::append_values( - extensions, // Desktop GPUs typically support all extensions we care. - Extension::VK_KHR_16bit_storage, Extension::VK_KHR_8bit_storage, - Extension::VK_KHR_shader_float16_int8, - Extension::VK_KHR_shader_integer_dot_product, Extension::VK_KHR_spirv_1_4, - Extension::VK_KHR_storage_buffer_storage_class, - Extension::VK_KHR_buffer_device_address, - Extension::VK_KHR_variable_pointers, - Extension::VK_EXT_subgroup_size_control); - if (getVendor(triple) == spirv::Vendor::NVIDIA || - triple.getArch() == TargetTripleArch::AMD_RDNAv3) { - extensions.push_back(Extension::VK_KHR_cooperative_matrix); - } -} - -/// Returns the Vulkan features/limits/capabilities supported by the given -/// `triple`. -/// -/// Note that this is an "approximation": Android compatibility will provide -/// some minimal guarantee but still different Android devices can have -/// different set of extensions, depending on the Android and GPU driver -/// version. The GPU triple is a handy way to specify the target but we cannot -/// encode all the information in the triple. -CapabilitiesAttr getCapabilities(const TargetTriple &triple, - MLIRContext *context) { - // Default to Vulkan required limits. - int maxComputeSharedMemorySize = 16384; - int maxComputeWorkGroupInvocations = 128; - std::array maxComputeWorkGroupSize = {128, 128, 64}; - - int subgroupSize = 32; - SubgroupFeature subgroupFeatures = SubgroupFeature::Basic; - std::optional minSubgroupSize, maxSubgroupSize; - - bool shaderFloat16 = false, shaderFloat64 = false; - bool shaderInt8 = false, shaderInt16 = false, shaderInt64 = false; - - bool shaderIntegerDotProduct = false; - - bool storageBuffer16BitAccess = false, storagePushConstant16 = false; - bool uniformAndStorageBuffer16BitAccess = false; - bool storageBuffer8BitAccess = false, storagePushConstant8 = false; - bool uniformAndStorageBuffer8BitAccess = false; - bool physicalStorageBufferAddresses = false; - - bool variablePointers = false, variablePointersStorageBuffer = false; - - SmallVector coopmatCases; - - Builder builder(context); - - switch (triple.getArch()) { - case TargetTripleArch::AMD_RDNAv3: { - auto i8t = builder.getIntegerType(8); - auto i32t = builder.getIntegerType(32); - auto f16t = builder.getF16Type(); - auto f32t = builder.getF32Type(); - auto scope = ScopeKHRAttr::get(context, ScopeKHR::Subgroup); - - // Note: The driver also advertises saturating arithmetic, so we can - // declare this when needed. - coopmatCases.push_back(CooperativeMatrixPropertiesKHRAttr::get( - context, - /*mSize=*/16, /*nSize=*/16, /*kSize=*/16, /*aType=*/i8t, - /*bType=*/i8t, /*cType=*/i32t, /*resultType=*/i32t, /*accSat=*/false, - /*scope=*/scope)); - coopmatCases.push_back(CooperativeMatrixPropertiesKHRAttr::get( - context, - /*mSize=*/16, /*nSize=*/16, /*kSize=*/16, /*aType=*/f16t, - /*bType=*/f16t, /*cType=*/f16t, /*resultType=*/f16t, /*accSat=*/false, - /*scope=*/scope)); - coopmatCases.push_back(CooperativeMatrixPropertiesKHRAttr::get( - context, - /*mSize=*/16, /*nSize=*/16, /*kSize=*/16, /*aType=*/f16t, - /*bType=*/f16t, /*cType=*/f32t, /*resultType=*/f32t, /*accSat=*/false, - /*scope=*/scope)); - } - LLVM_FALLTHROUGH; - case TargetTripleArch::AMD_RDNAv1: - case TargetTripleArch::AMD_RDNAv2: - // Example: https://vulkan.gpuinfo.org/displayreport.php?id=10906 - maxComputeSharedMemorySize = 65536; - maxComputeWorkGroupInvocations = 1024; - maxComputeWorkGroupSize = {1024, 1024, 1024}; - - subgroupSize = 64, minSubgroupSize = 32, maxSubgroupSize = 64; - subgroupFeatures = SubgroupFeature::Basic | SubgroupFeature::Vote | - SubgroupFeature::Arithmetic | SubgroupFeature::Ballot | - SubgroupFeature::Shuffle | - SubgroupFeature::ShuffleRelative | - SubgroupFeature::Clustered | SubgroupFeature::Quad; - - shaderFloat16 = shaderFloat64 = true; - shaderInt8 = shaderInt16 = shaderInt64 = true; - - shaderIntegerDotProduct = true; - - storageBuffer16BitAccess = storagePushConstant16 = true; - uniformAndStorageBuffer16BitAccess = true; - storageBuffer8BitAccess = true, storagePushConstant8 = true; - uniformAndStorageBuffer8BitAccess = true; - physicalStorageBufferAddresses = true; - - variablePointers = variablePointersStorageBuffer = true; - break; - case TargetTripleArch::Apple_M1: - // Example: https://vulkan.gpuinfo.org/displayreport.php?id=14673 - maxComputeSharedMemorySize = 32768; - maxComputeWorkGroupInvocations = 1024; - maxComputeWorkGroupSize = {1024, 1024, 1024}; - - subgroupSize = 32; - subgroupFeatures = SubgroupFeature::Basic | SubgroupFeature::Vote | - SubgroupFeature::Arithmetic | SubgroupFeature::Ballot | - SubgroupFeature::Shuffle | - SubgroupFeature::ShuffleRelative | SubgroupFeature::Quad; - - shaderFloat16 = true; - shaderFloat64 = false; - shaderInt8 = shaderInt16 = shaderInt64 = true; - - storageBuffer16BitAccess = storagePushConstant16 = true; - uniformAndStorageBuffer16BitAccess = true; - storageBuffer8BitAccess = true, storagePushConstant8 = true; - uniformAndStorageBuffer8BitAccess = true; - physicalStorageBufferAddresses = true; - - variablePointers = variablePointersStorageBuffer = true; - break; - case TargetTripleArch::ARM_Valhall: - // Example: https://vulkan.gpuinfo.org/displayreport.php?id=10312 (11) - // Example: https://vulkan.gpuinfo.org/displayreport.php?id=15142 (12) - maxComputeSharedMemorySize = 32768; - maxComputeWorkGroupInvocations = 512; - maxComputeWorkGroupSize = {512, 512, 512}; - - subgroupSize = 16; - subgroupFeatures = SubgroupFeature::Basic | SubgroupFeature::Vote | - SubgroupFeature::Arithmetic | SubgroupFeature::Ballot | - SubgroupFeature::Clustered | SubgroupFeature::Quad; - - if (triple.getOS() == TargetTripleOS::Android31) { - subgroupFeatures = subgroupFeatures | SubgroupFeature::Shuffle | - SubgroupFeature::ShuffleRelative; - } - - shaderFloat16 = shaderInt8 = shaderInt16 = true; - - shaderIntegerDotProduct = true; - - storageBuffer16BitAccess = storagePushConstant16 = true; - uniformAndStorageBuffer16BitAccess = true; - storageBuffer8BitAccess = true, storagePushConstant8 = true; - uniformAndStorageBuffer8BitAccess = true; - - variablePointers = variablePointersStorageBuffer = true; - break; - case TargetTripleArch::CPU: - if (triple.getProduct() == TargetTripleProduct::SwiftShader) { - // Example: https://vulkan.gpuinfo.org/displayreport.php?id=11023 - maxComputeSharedMemorySize = 16384; - - subgroupSize = 4; - subgroupFeatures = SubgroupFeature::Basic | SubgroupFeature::Vote | - SubgroupFeature::Arithmetic | SubgroupFeature::Ballot | - SubgroupFeature::Shuffle | - SubgroupFeature::ShuffleRelative; - } - break; - case TargetTripleArch::NV_Turing: - case TargetTripleArch::NV_Ampere: { - // Example: https://vulkan.gpuinfo.org/displayreport.php?id=11252 - maxComputeSharedMemorySize = 49152; - maxComputeWorkGroupInvocations = 1024; - maxComputeWorkGroupSize = {1024, 1024, 64}; - - subgroupSize = 32, minSubgroupSize = 32, maxSubgroupSize = 32; - subgroupFeatures = SubgroupFeature::Basic | SubgroupFeature::Vote | - SubgroupFeature::Arithmetic | SubgroupFeature::Ballot | - SubgroupFeature::Shuffle | - SubgroupFeature::ShuffleRelative | - SubgroupFeature::Clustered | SubgroupFeature::Quad; - - shaderFloat16 = shaderFloat64 = true; - shaderInt8 = shaderInt16 = shaderInt64 = true; - - shaderIntegerDotProduct = true; - - storageBuffer16BitAccess = storagePushConstant16 = true; - uniformAndStorageBuffer16BitAccess = true; - storageBuffer8BitAccess = true, storagePushConstant8 = true; - uniformAndStorageBuffer8BitAccess = true; - physicalStorageBufferAddresses = true; - - variablePointers = variablePointersStorageBuffer = true; - - auto i8t = builder.getIntegerType(8); - auto i32t = builder.getIntegerType(32); - auto f16t = builder.getF16Type(); - auto f32t = builder.getF32Type(); - auto scope = ScopeKHRAttr::get(context, ScopeKHR::Subgroup); - - // Note: the driver also advertises other shapes that can enabled when - // needed. - coopmatCases.push_back(CooperativeMatrixPropertiesKHRAttr::get( - context, - /*mSize=*/8, /*nSize=*/8, /*kSize=*/32, /*aType=*/i8t, - /*bType=*/i8t, /*cType=*/i32t, /*resultType=*/i32t, /*accSat=*/false, - /*scope=*/scope)); - coopmatCases.push_back(CooperativeMatrixPropertiesKHRAttr::get( - context, - /*mSize=*/16, /*nSize=*/16, /*kSize=*/16, /*aType=*/f16t, - /*bType=*/f16t, /*cType=*/f16t, /*resultType=*/f16t, /*accSat=*/false, - /*scope=*/scope)); - coopmatCases.push_back(CooperativeMatrixPropertiesKHRAttr::get( - context, - /*mSize=*/16, /*nSize=*/16, /*kSize=*/16, /*aType=*/f16t, - /*bType=*/f16t, /*cType=*/f32t, /*resultType=*/f32t, /*accSat=*/false, - /*scope=*/scope)); - } break; - case TargetTripleArch::NV_Pascal: - // Example: https://vulkan.gpuinfo.org/displayreport.php?id=17937 - maxComputeSharedMemorySize = 49152; - maxComputeWorkGroupInvocations = 1536; - maxComputeWorkGroupSize = {1536, 1024, 64}; - - subgroupSize = 32, minSubgroupSize = 32, maxSubgroupSize = 32; - subgroupFeatures = SubgroupFeature::Basic | SubgroupFeature::Vote | - SubgroupFeature::Arithmetic | SubgroupFeature::Ballot | - SubgroupFeature::Shuffle | - SubgroupFeature::ShuffleRelative | - SubgroupFeature::Clustered | SubgroupFeature::Quad; - - shaderFloat16 = shaderFloat64 = true; - shaderInt8 = shaderInt16 = shaderInt64 = true; - - shaderIntegerDotProduct = true; - - storageBuffer16BitAccess = storagePushConstant16 = true; - uniformAndStorageBuffer16BitAccess = true; - storageBuffer8BitAccess = true, storagePushConstant8 = true; - uniformAndStorageBuffer8BitAccess = true; - physicalStorageBufferAddresses = true; - - variablePointers = variablePointersStorageBuffer = true; - break; - case TargetTripleArch::QC_Adreno: - // Example: https://vulkan.gpuinfo.org/displayreport.php?id=10983 (11) - // Example: https://vulkan.gpuinfo.org/displayreport.php?id=16312 (12) - maxComputeSharedMemorySize = 32768; - maxComputeWorkGroupInvocations = 1024; - maxComputeWorkGroupSize = {1024, 1024, 64}; - - subgroupSize = 64; - subgroupFeatures = SubgroupFeature::Basic | SubgroupFeature::Vote | - SubgroupFeature::Arithmetic | SubgroupFeature::Ballot | - SubgroupFeature::Shuffle | - SubgroupFeature::ShuffleRelative | SubgroupFeature::Quad; - - shaderFloat16 = shaderInt8 = shaderInt16 = true; - - storageBuffer16BitAccess = true; - if (triple.getOS() == TargetTripleOS::Android31) { - storageBuffer8BitAccess = true; - } - - variablePointers = variablePointersStorageBuffer = true; - break; - case TargetTripleArch::Intel_Arc: - // Example: https://vulkan.gpuinfo.org/displayreport.php?id=19818 - maxComputeSharedMemorySize = 32768; - maxComputeWorkGroupInvocations = 1024; - maxComputeWorkGroupSize = {1024, 1024, 64}; - - subgroupSize = 32, minSubgroupSize = 8, maxSubgroupSize = 32; - subgroupFeatures = SubgroupFeature::Basic | SubgroupFeature::Vote | - SubgroupFeature::Arithmetic | SubgroupFeature::Ballot | - SubgroupFeature::Shuffle | - SubgroupFeature::ShuffleRelative | - SubgroupFeature::Clustered | SubgroupFeature::Quad; - - shaderFloat16 = true; - shaderFloat64 = false; - shaderInt8 = shaderInt16 = true; - shaderInt64 = false; - - shaderIntegerDotProduct = true; - - storageBuffer16BitAccess = storagePushConstant16 = true; - uniformAndStorageBuffer16BitAccess = true; - storageBuffer8BitAccess = true, storagePushConstant8 = true; - uniformAndStorageBuffer8BitAccess = true; - physicalStorageBufferAddresses = true; - - variablePointers = variablePointersStorageBuffer = true; - break; - case TargetTripleArch::Unknown: - // Use the largest subgroup size we can find across various vendors. - subgroupSize = 64; - // The following capabilities have 90%+ device coverage (Vulkan 1.1+) - // from https://vulkan.gpuinfo.org/listfeaturesextensions.php. - variablePointers = variablePointersStorageBuffer = false; - // Use Vulkan default for others. - break; - } - - auto getBoolAttr = [context](bool value) { - return value ? UnitAttr::get(context) : UnitAttr(); - }; - - return CapabilitiesAttr::get( - context, maxComputeSharedMemorySize, maxComputeWorkGroupInvocations, - builder.getI32VectorAttr(maxComputeWorkGroupSize), - getBoolAttr(shaderFloat64), getBoolAttr(shaderInt16), - getBoolAttr(shaderInt64), - SubgroupFeatureAttr::get(context, subgroupFeatures), subgroupSize, - minSubgroupSize, maxSubgroupSize, getBoolAttr(storageBuffer16BitAccess), - getBoolAttr(storagePushConstant16), - getBoolAttr(uniformAndStorageBuffer16BitAccess), - getBoolAttr(storageBuffer8BitAccess), getBoolAttr(storagePushConstant8), - getBoolAttr(uniformAndStorageBuffer8BitAccess), - getBoolAttr(physicalStorageBufferAddresses), getBoolAttr(shaderFloat16), - getBoolAttr(shaderInt8), getBoolAttr(shaderIntegerDotProduct), - getBoolAttr(variablePointersStorageBuffer), getBoolAttr(variablePointers), - builder.getArrayAttr(coopmatCases)); -} -} // namespace - -TargetTriple TargetTriple::get(const char *triple) { - llvm::SmallVector fragments; - llvm::SplitString(triple, fragments, "-"); - TargetTripleArch arch = TargetTripleArch::Unknown; - if (auto symbol = symbolizeTargetTripleArch(fragments[0])) { - arch = symbol.value(); - } - TargetTripleProduct product = TargetTripleProduct::Unknown; - if (auto symbol = symbolizeTargetTripleProduct(fragments[1])) { - product = symbol.value(); - } - TargetTripleOS os = TargetTripleOS::Unknown; - if (auto symbol = symbolizeTargetTripleOS(fragments[2])) { - os = symbol.value(); - } - return TargetTriple(arch, product, os); -} - -TargetTriple::TargetTriple(TargetTripleArch arch, TargetTripleProduct product, - TargetTripleOS os) - : arch(arch), product(product), os(os) {} - -std::string TargetTriple::getTriple() const { - llvm::StringRef archStr = stringifyTargetTripleArch(arch); - llvm::StringRef productStr = stringifyTargetTripleProduct(product); - llvm::StringRef osStr = stringifyTargetTripleOS(os); - return llvm::formatv("{0}-{1}-{2}", archStr, productStr, osStr); -} - -TargetEnvAttr TargetTriple::getTargetEnv(MLIRContext *context) const { - SmallVector extensions; - getExtensions(*this, extensions); - return TargetEnvAttr::get(getVersion(*this), /*revision=*/0, extensions, - getVendor(*this), getDeviceType(*this), - spirv::TargetEnvAttr::kUnknownDeviceID, - getCapabilities(*this, context)); -} - -} // namespace mlir::iree_compiler::IREE::Vulkan diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetTriple.h b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetTriple.h deleted file mode 100644 index 7ea5e0df0168..000000000000 --- a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetTriple.h +++ /dev/null @@ -1,67 +0,0 @@ -// Copyright 2021 The IREE Authors -// -// Licensed under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception - -#ifndef IREE_COMPILER_DIALECT_VULKAN_UTILS_TARGETTRIPLE_H_ -#define IREE_COMPILER_DIALECT_VULKAN_UTILS_TARGETTRIPLE_H_ - -#include - -#include "iree/compiler/Dialect/Vulkan/IR/VulkanAttributes.h" -#include "iree/compiler/Dialect/Vulkan/IR/VulkanTypes.h" -#include "mlir/IR/MLIRContext.h" - -namespace mlir::iree_compiler::IREE::Vulkan { - -/// GPU triple definitions to describe GPU targets for compilers. -/// -/// We use "triple" here to match common compiler language: historically one -/// would describe a CPU compiler target as a string containing exactly three -/// fields. But here the configuration is for GPU and there can exist a lot of -/// architectures/vendors/products/systems. What matters differ from CPU -/// triples. We define it in the form of: -/// -/// -- -/// -/// For example: -/// ampere-rtx3080-windows -/// rdna1-5700xt-linux -/// adreno-a650-android30 -/// valhall-unknown-android30 -/// cpu-swiftshader-unknown -/// -/// Vendor and architecture are combined together because: -/// * Typically each GPU vendor has its own set of architectures. So given the -/// architecture we know which vendor it is from. This is different from CPU -/// land where the the same architecture can be implemented by mulitple -/// vendors. -/// * There are vendors that we don't have public information regarding its -/// architectures. -/// We need a field for product to differentiate the cases where the -/// architecture is unknown or ambiguous. -class TargetTriple { -public: - static TargetTriple get(const char *triple); - - TargetTriple(TargetTripleArch, TargetTripleProduct, TargetTripleOS); - - TargetTripleArch getArch() const { return arch; } - TargetTripleProduct getProduct() const { return product; } - TargetTripleOS getOS() const { return os; } - - /// Returns the triple string. - std::string getTriple() const; - - TargetEnvAttr getTargetEnv(MLIRContext *context) const; - -private: - TargetTripleArch arch; - TargetTripleProduct product; - TargetTripleOS os; -}; - -} // namespace mlir::iree_compiler::IREE::Vulkan - -#endif // IREE_COMPILER_DIALECT_VULKAN_UTILS_TARGETTRIPLE_H_ diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/BUILD.bazel b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/BUILD.bazel deleted file mode 100644 index 687fa49cba44..000000000000 --- a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/BUILD.bazel +++ /dev/null @@ -1,37 +0,0 @@ -# Copyright 2020 The IREE Authors -# -# Licensed under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception - -load("//build_tools/bazel:build_defs.oss.bzl", "iree_cmake_extra_content") -load("//build_tools/bazel:enforce_glob.bzl", "enforce_glob") -load("//build_tools/bazel:iree_lit_test.bzl", "iree_lit_test_suite") - -package( - features = ["layering_check"], - licenses = ["notice"], # Apache 2.0 -) - -iree_cmake_extra_content( - content = """ -if(NOT IREE_TARGET_BACKEND_VULKAN_SPIRV) - return() -endif() -""", -) - -iree_lit_test_suite( - name = "lit", - srcs = enforce_glob( - [ - "target_env_conversion.mlir", - ], - include = ["*.mlir"], - ), - cfg = "//compiler:lit.cfg.py", - tools = [ - "//tools:iree-opt", - "@llvm-project//llvm:FileCheck", - ], -) diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/CMakeLists.txt b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/CMakeLists.txt deleted file mode 100644 index bb5cbe5bbd69..000000000000 --- a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/CMakeLists.txt +++ /dev/null @@ -1,27 +0,0 @@ -################################################################################ -# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from # -# compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/BUILD.bazel # -# # -# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary # -# CMake-only content. # -# # -# To disable autogeneration for this file entirely, delete this header. # -################################################################################ - -if(NOT IREE_TARGET_BACKEND_VULKAN_SPIRV) - return() -endif() - -iree_add_all_subdirs() - -iree_lit_test_suite( - NAME - lit - SRCS - "target_env_conversion.mlir" - TOOLS - FileCheck - iree-opt -) - -### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ### diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/target_env_conversion.mlir b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/target_env_conversion.mlir deleted file mode 100644 index 3a23031eb1ad..000000000000 --- a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/target_env_conversion.mlir +++ /dev/null @@ -1,86 +0,0 @@ -// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetBackends=vulkan-spirv},iree-hal-transformation-pipeline{serialize-executables=false})' %s | FileCheck %s --check-prefix=DEFAULT -// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetBackends=vulkan-spirv},iree-hal-transformation-pipeline{serialize-executables=false})' --iree-vulkan-target-triple=adreno-a650-android30 %s | FileCheck %s --check-prefix=ADRENO -// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetBackends=vulkan-spirv},iree-hal-transformation-pipeline{serialize-executables=false})' --iree-vulkan-target-triple=valhall-unknown-android31 %s | FileCheck %s --check-prefix=VALHALL -// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetBackends=vulkan-spirv},iree-hal-transformation-pipeline{serialize-executables=false})' --iree-vulkan-target-triple=turing-t4-linux %s | FileCheck %s --check-prefix=TURING -// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetBackends=vulkan-spirv},iree-hal-transformation-pipeline{serialize-executables=false})' --iree-vulkan-target-triple=rdna1-5700xt-windows %s | FileCheck %s --check-prefix=RDNA1 -// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetBackends=vulkan-spirv},iree-hal-transformation-pipeline{serialize-executables=false})' --iree-vulkan-target-triple=rdna3-6900xtx-windows %s | FileCheck %s --check-prefix=RDNA3 -// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetBackends=vulkan-spirv},iree-hal-transformation-pipeline{serialize-executables=false})' --iree-vulkan-target-triple=m1-moltenvk-macos %s | FileCheck %s --check-prefix=M1 -// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetBackends=vulkan-spirv},iree-hal-transformation-pipeline{serialize-executables=false})' --iree-vulkan-target-triple=arc-770-windows %s | FileCheck %s --check-prefix=ARC -// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetBackends=vulkan-spirv},iree-hal-transformation-pipeline{serialize-executables=false})' --iree-vulkan-target-triple=pascal-1080-windows %s | FileCheck %s --check-prefix=PASCAL -// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetBackends=vulkan-spirv},iree-hal-transformation-pipeline{serialize-executables=false})' --iree-vulkan-target-env="#vk.target_env: vector<3xi32>, subgroupFeatures = 63 : i32, subgroupSize = 4 >>" %s | FileCheck %s --check-prefix=ENV - -// TODO(antiagainst): Passing in lenghty strings as command-line options is not -// optimal. We should consider creating a dedicated test pass to pick up -// #vk.target_env in input assembly and convert them. - -// DEFAULT: #spirv.target_env<#spirv.vce, -// DEFAULT-SAME: api=Vulkan, #spirv.resource_limits> - -// ADRENO: #spirv.target_env<#spirv.vce, -// ADRENO-SAME: api=Vulkan, Qualcomm:IntegratedGPU, #spirv.resource_limits> - -// VALHALL: #spirv.target_env<#spirv.vce, -// VALHALL-SAME: api=Vulkan, ARM:IntegratedGPU, #spirv.resource_limits> - -// TURING: #spirv.target_env<#spirv.vce, -// TURING-SAME: api=Vulkan, NVIDIA:DiscreteGPU, #spirv.resource_limits>, #spirv.coop_matrix_props_khr>, #spirv.coop_matrix_props_khr>]>> - -// RDNA1: #spirv.target_env<#spirv.vce, -// RDNA1-SAME: api=Vulkan, AMD:DiscreteGPU, #spirv.resource_limits> - -// RDNA3: #spirv.target_env<#spirv.vce, -// RDNA3-SAME: api=Vulkan, AMD:DiscreteGPU, #spirv.resource_limits>, #spirv.coop_matrix_props_khr>, #spirv.coop_matrix_props_khr>]>> - -// M1: #spirv.target_env<#spirv.vce, -// M1-SAME: api=Vulkan, Apple:IntegratedGPU, #spirv.resource_limits> - -// ARC: #spirv.target_env<#spirv.vce, -// ARC-SAME: api=Vulkan, Intel:DiscreteGPU, #spirv.resource_limits>}> - -// PASCAL: #spirv.target_env<#spirv.vce, -// PASCAL-SAME: api=Vulkan, NVIDIA:DiscreteGPU, #spirv.resource_limits>}> - -// ENV: #spirv.target_env<#spirv.vce, -// ENV-SAME: api=Vulkan, AMD:DiscreteGPU, #spirv.resource_limits> - -stream.executable public @reduce_dispatch { - stream.executable.export @reduce_dispatch workgroups(%arg0: index) -> (index, index, index) { - %x, %y, %z = flow.dispatch.workgroup_count_from_dag_root %arg0 - stream.return %x, %y, %z : index, index, index - } - builtin.module { - func.func @reduce_dispatch(%arg0_binding: !stream.binding, %arg1_binding: !stream.binding) { - %c0 = arith.constant 0 : index - %arg0 = stream.binding.subspan %arg0_binding[%c0] : !stream.binding -> !flow.dispatch.tensor> - %arg1 = stream.binding.subspan %arg1_binding[%c0] : !stream.binding -> !flow.dispatch.tensor> - %0 = tensor.empty() : tensor - %1 = flow.dispatch.tensor.load %arg0, offsets=[0], sizes=[16], strides=[1] : !flow.dispatch.tensor> -> tensor<16xf32> - %3 = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> ()>], iterator_types = ["reduction"]} ins(%1 : tensor<16xf32>) outs(%0 : tensor) { - ^bb0(%arg2: f32, %arg3: f32): - %4 = arith.addf %arg2, %arg3 : f32 - linalg.yield %4 : f32 - } -> tensor - flow.dispatch.tensor.store %3, %arg1, offsets=[], sizes=[], strides=[] : tensor -> !flow.dispatch.tensor> - return - } - } -} diff --git a/compiler/src/iree/compiler/Tools/BUILD.bazel b/compiler/src/iree/compiler/Tools/BUILD.bazel index 314a09c6fa5e..95c620ca1e1b 100644 --- a/compiler/src/iree/compiler/Tools/BUILD.bazel +++ b/compiler/src/iree/compiler/Tools/BUILD.bazel @@ -56,7 +56,6 @@ iree_compiler_cc_library( "//compiler/src/iree/compiler/Dialect/VM/Transforms", "//compiler/src/iree/compiler/Dialect/VMVX/IR:VMVXDialect", "//compiler/src/iree/compiler/Dialect/VMVX/Transforms", - "//compiler/src/iree/compiler/Dialect/Vulkan/IR", "//compiler/src/iree/compiler/ExternalInterfaces:ExternalModels", "//compiler/src/iree/compiler/GlobalOptimization/Interfaces", "//compiler/src/iree/compiler/InputConversion/Common", diff --git a/compiler/src/iree/compiler/Tools/CMakeLists.txt b/compiler/src/iree/compiler/Tools/CMakeLists.txt index a38a4db0ccf3..ee8d8201785a 100644 --- a/compiler/src/iree/compiler/Tools/CMakeLists.txt +++ b/compiler/src/iree/compiler/Tools/CMakeLists.txt @@ -53,7 +53,6 @@ iree_cc_library( iree::compiler::Dialect::VM::Transforms iree::compiler::Dialect::VMVX::IR::VMVXDialect iree::compiler::Dialect::VMVX::Transforms - iree::compiler::Dialect::Vulkan::IR iree::compiler::ExternalInterfaces::ExternalModels iree::compiler::GlobalOptimization::Interfaces::Interfaces iree::compiler::InputConversion::Common diff --git a/compiler/src/iree/compiler/Tools/init_iree_dialects.h b/compiler/src/iree/compiler/Tools/init_iree_dialects.h index 0472723d819d..4d333879ae11 100644 --- a/compiler/src/iree/compiler/Tools/init_iree_dialects.h +++ b/compiler/src/iree/compiler/Tools/init_iree_dialects.h @@ -22,20 +22,17 @@ #include "iree/compiler/Dialect/Flow/IR/FlowDialect.h" #include "iree/compiler/Dialect/HAL/IR/HALDialect.h" #include "iree/compiler/Dialect/LinalgExt/IR/LinalgExtDialect.h" -#include "iree/compiler/Dialect/LinalgExt/Transforms/Passes.h" #include "iree/compiler/Dialect/Stream/IR/StreamDialect.h" #include "iree/compiler/Dialect/Util/IR/UtilDialect.h" #include "iree/compiler/Dialect/Util/TransformOps/UtilTransformOps.h" #include "iree/compiler/Dialect/VM/IR/VMDialect.h" #include "iree/compiler/Dialect/VMVX/IR/VMVXDialect.h" -#include "iree/compiler/Dialect/Vulkan/IR/VulkanDialect.h" #include "iree/compiler/ExternalInterfaces/Interfaces.h" #include "iree/compiler/GlobalOptimization/Interfaces/Interfaces.h" #include "iree/compiler/Modules/HAL/Inline/IR/HALInlineDialect.h" #include "iree/compiler/Modules/HAL/Loader/IR/HALLoaderDialect.h" #include "iree/compiler/Modules/IO/Parameters/IR/IOParametersDialect.h" #include "iree/compiler/Preprocessing/TransformExtensions/PreprocessingExtensions.h" -#include "mlir/IR/Dialect.h" namespace mlir::iree_compiler { @@ -56,8 +53,7 @@ inline void registerIreeDialects(DialectRegistry ®istry) { IREE::Util::UtilDialect, IREE::VM::VMDialect, IREE::VMVX::VMVXDialect, - IREE::VectorExt::IREEVectorExtDialect, - IREE::Vulkan::VulkanDialect>(); + IREE::VectorExt::IREEVectorExtDialect>(); // clang-format on // External models. diff --git a/samples/custom_dispatch/vulkan/shaders/example.mlir b/samples/custom_dispatch/vulkan/shaders/example.mlir index d655b9884f40..ef10fb7b7dbd 100644 --- a/samples/custom_dispatch/vulkan/shaders/example.mlir +++ b/samples/custom_dispatch/vulkan/shaders/example.mlir @@ -14,9 +14,12 @@ // and compilation options (architectures, etc) can be embedded for runtime // selection. #spirv_target = #hal.executable.target<"vulkan-spirv", "vulkan-spirv-fb", { - spirv.target_env = #spirv.target_env< - #spirv.vce, - #spirv.resource_limits + 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 = [64, 64], + max_workgroup_sizes = [128, 128, 64], max_thread_count_per_workgroup = 128, + max_workgroup_memory_bytes = 16384> > }> diff --git a/samples/custom_dispatch/vulkan/shaders/example_inline.mlir b/samples/custom_dispatch/vulkan/shaders/example_inline.mlir index 5cdbcac7e280..36912bb35df9 100644 --- a/samples/custom_dispatch/vulkan/shaders/example_inline.mlir +++ b/samples/custom_dispatch/vulkan/shaders/example_inline.mlir @@ -14,9 +14,12 @@ // and compilation options (architectures, etc) can be embedded for runtime // selection. #spirv_target = #hal.executable.target<"vulkan-spirv", "vulkan-spirv-fb", { - spirv.target_env = #spirv.target_env< - #spirv.vce, - #spirv.resource_limits + 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 = [64, 64], + max_workgroup_sizes = [128, 128, 64], max_thread_count_per_workgroup = 128, + max_workgroup_memory_bytes = 16384> > }> diff --git a/samples/custom_dispatch/vulkan/shaders/example_transform.mlir b/samples/custom_dispatch/vulkan/shaders/example_transform.mlir index 3766a3032928..b4885a03081d 100644 --- a/samples/custom_dispatch/vulkan/shaders/example_transform.mlir +++ b/samples/custom_dispatch/vulkan/shaders/example_transform.mlir @@ -18,10 +18,12 @@ // custom kernel. For things to be truly portable, we need to be able to compare // executable configurations. #spirv_target = #hal.executable.target<"vulkan-spirv", "vulkan-spirv-fb", { - spirv.target_env = #spirv.target_env< - #spirv.vce, - #spirv.resource_limits + iree.gpu.target = #iree_gpu.target< + arch = "", features = "spirv:v1.3,cap:Shader", wgp = < + compute = fp32|int32, storage = b32, subgroup = shuffle|arithmetic, + 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> > }> diff --git a/samples/custom_dispatch/vulkan/shaders/example_transform_spec.mlir b/samples/custom_dispatch/vulkan/shaders/example_transform_spec.mlir index 70ad8989c501..5bcdafe7fba1 100644 --- a/samples/custom_dispatch/vulkan/shaders/example_transform_spec.mlir +++ b/samples/custom_dispatch/vulkan/shaders/example_transform_spec.mlir @@ -7,10 +7,12 @@ // The configuration used for executable compilation. // This specifies the device configurations that support this custom kernel. #spirv_target = #hal.executable.target<"vulkan-spirv", "vulkan-spirv-fb", { - spirv.target_env = #spirv.target_env< - #spirv.vce, - #spirv.resource_limits + iree.gpu.target = #iree_gpu.target< + arch = "", features = "spirv:v1.3,cap:Shader", wgp = < + compute = fp32|int32, storage = b32, subgroup = shuffle|arithmetic, + 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> > }> diff --git a/samples/transform_dialect/example_module.mlir b/samples/transform_dialect/example_module.mlir index 13128e16c3a8..585bb2591534 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, api=Vulkan, #spirv.resource_limits> +#target = #iree_gpu.target> 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, <1, storage_buffer>]>]>) { ^bb0(%arg0: !hal.device): @@ -63,7 +65,7 @@ module attributes { } } 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, <1, storage_buffer, ReadOnly>, <2, storage_buffer>]>]>) { ^bb0(%arg0: !hal.device): @@ -87,7 +89,7 @@ module attributes { } } 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, <1, storage_buffer>]>]>) { ^bb0(%arg0: !hal.device):