Reland "[spirv] Switch to use common target description" (#17699)
This relands https://github.com/iree-org/iree/pull/17623.
This commit switches SPIR-V side to use the common `#iree_gpu.target` to
describe the GPU characteristics. With it we can now remove the ad-hoc
Vulkan attributes and dialects and unify how GPU are described across
various GPU compiler backends in IREE.
SPIR-V has some additional requirements that we need to account for:
We have many vendors and APIs to handle there so this commit adds
various AMD/ARM/NVIDIA/Qualcomm targets for
development purposes so that we can specify them with a shorthand.
In order to be extensible, leverage the `feature` field in
`#iree_gpu.target` to specify additional capabilities with `cap:` prefix
and extensions with `ext:` prefix. We also use the `feature` field to
specify what SPIR-V version to target with the `spirv:v1.x` format.
Right now the `SPIRVConvertGPUTarget` pass is
invoked immediately before configuration. This is to stage the changes.
As a next step we need to move
it immediately before `ConvertToSPIRV` pass.
`--iree-vulkan-target-env` is dropped given now we removed the whole
Vulkan dialect and cannot control with a `#vk.target_env` attribute
anymore.
The default `--iree-vulkan-target-triple` now becomes
`vp_android_baseline_2022`, which is a a good lowest common denominator
to guarantee the generated SPIR-V is widely accepted. We are not
considering SwiftShader now anymore like previously due to testing
purposes.
The `--iree-vulkan-target-triple` should be renamed given it's not a
triple anymore--that will happen later together with other GPU backends
(i.e., cuda/hip) to be consistent.
In order to support cooperative matrix conversion, we added
`WMMA_F16_16x16x16_F16`. For NVIDIA GPUs
we are abusing it right now without considering the concrete explicit
layout--that is fine given in Vulkan they are opaque anyway. But this
need to be fixed if we are targeting WMMA in CUDA.
We now contruct a `#iree_gpu.target` to specify
the target to drive SPIR-V CodeGen.
Progress towards https://github.com/iree-org/iree/issues/16341
ci-extra: test_nvidia_gpu,test_nvidia_a100,test_amd_mi250,
build_test_all_macos_arm64,build_and_test_android,test_on_moto-edge-x30
---------
Signed-off-by: Lei Zhang <antiagainst@gmail.com>
diff --git a/compiler/plugins/target/MetalSPIRV/BUILD.bazel b/compiler/plugins/target/MetalSPIRV/BUILD.bazel
index 9773eff..ede5566 100644
--- a/compiler/plugins/target/MetalSPIRV/BUILD.bazel
+++ b/compiler/plugins/target/MetalSPIRV/BUILD.bazel
@@ -26,6 +26,7 @@
":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 4dd1b06..678a37a 100644
--- a/compiler/plugins/target/MetalSPIRV/CMakeLists.txt
+++ b/compiler/plugins/target/MetalSPIRV/CMakeLists.txt
@@ -36,6 +36,7 @@
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 6ea1cbf..25e8e51 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 @@
};
} // namespace
-static spirv::TargetEnvAttr getMetalTargetEnv(MLIRContext *context) {
- using spirv::Capability;
- using spirv::Extension;
-
- // Capabilities and limits according to Metal 3 devices.
- const std::array<Extension, 4> 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<Capability, 21> 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 @@
MLIRContext *context, StringRef deviceID, DictionaryAttr deviceConfigAttr,
SmallVectorImpl<IREE::HAL::ExecutableTargetAttr> &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<NamedAttribute> 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<IREE::HAL::ExecutableTargetAttr>(
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 84dc61e..720e00b 100644
--- a/compiler/plugins/target/MetalSPIRV/test/smoketest.mlir
+++ b/compiler/plugins/target/MetalSPIRV/test/smoketest.mlir
@@ -4,7 +4,9 @@
hal.device.targets = [
#hal.device.target<"metal", [
#hal.executable.target<"metal-spirv", "metal-msl-fb", {
- spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #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 = [32],
+ max_workgroup_sizes = [128, 128, 64], max_thread_count_per_workgroup = 128, max_workgroup_memory_bytes = 16384>>
}>
]>
]
diff --git a/compiler/plugins/target/ROCM/test/target_device_features.mlir b/compiler/plugins/target/ROCM/test/target_device_features.mlir
index 9f01246..0f427c5 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 = [<MFMA_F16_16x16x16_F32>, <MFMA_F16_32x32x8_F32>]
// GFX1100: target = #iree_gpu.target<arch = "gfx1100",
-// GFX1100-SAME: mma = [<WMMA_F16_16x16x16_F32>]
+// GFX1100-SAME: mma = [<WMMA_F16_16x16x16_F32>, <WMMA_F16_16x16x16_F16>]
// GFX1100-SAME: subgroup_size_choices = [32, 64]
// GFX941: target = #iree_gpu.target<arch = "gfx941",
diff --git a/compiler/plugins/target/VulkanSPIRV/BUILD.bazel b/compiler/plugins/target/VulkanSPIRV/BUILD.bazel
index 4e51dee..984bef9 100644
--- a/compiler/plugins/target/VulkanSPIRV/BUILD.bazel
+++ b/compiler/plugins/target/VulkanSPIRV/BUILD.bazel
@@ -25,11 +25,10 @@
deps = [
"//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/HAL/Target",
- "//compiler/src/iree/compiler/Dialect/Vulkan/IR",
- "//compiler/src/iree/compiler/Dialect/Vulkan/Utils",
"//compiler/src/iree/compiler/PluginAPI",
"//compiler/src/iree/compiler/Utils",
"//runtime/src/iree/schemas:spirv_executable_def_c_fbs",
diff --git a/compiler/plugins/target/VulkanSPIRV/CMakeLists.txt b/compiler/plugins/target/VulkanSPIRV/CMakeLists.txt
index c14b76a..958e277 100644
--- a/compiler/plugins/target/VulkanSPIRV/CMakeLists.txt
+++ b/compiler/plugins/target/VulkanSPIRV/CMakeLists.txt
@@ -33,11 +33,10 @@
MLIRSupport
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::HAL::Target
- iree::compiler::Dialect::Vulkan::IR
- iree::compiler::Dialect::Vulkan::Utils
iree::compiler::PluginAPI
iree::compiler::Utils
iree::schemas::spirv_executable_def_c_fbs
diff --git a/compiler/plugins/target/VulkanSPIRV/VulkanSPIRVTarget.cpp b/compiler/plugins/target/VulkanSPIRV/VulkanSPIRVTarget.cpp
index 49bd44d..5fdeb54 100644
--- a/compiler/plugins/target/VulkanSPIRV/VulkanSPIRVTarget.cpp
+++ b/compiler/plugins/target/VulkanSPIRV/VulkanSPIRVTarget.cpp
@@ -5,11 +5,9 @@
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
#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/HAL/Target/TargetRegistry.h"
-#include "iree/compiler/Dialect/Vulkan/IR/VulkanAttributes.h"
-#include "iree/compiler/Dialect/Vulkan/IR/VulkanDialect.h"
-#include "iree/compiler/Dialect/Vulkan/Utils/TargetEnvironment.h"
#include "iree/compiler/PluginAPI/Client.h"
#include "iree/compiler/Utils/FlatbufferUtils.h"
#include "iree/compiler/Utils/ModuleUtils.h"
@@ -34,20 +32,19 @@
namespace {
struct VulkanSPIRVTargetOptions {
- std::string targetTriple = "";
- std::string targetEnv = "";
+ // Use vp_android_baseline_2022 profile as the default target--it's a good
+ // lowest common denominator to guarantee the generated SPIR-V is widely
+ // accepted for now. Eventually we want to use a list for multi-targeting.
+ std::string targetTriple = "vp_android_baseline_2022";
bool indirectBindings = false;
void bindOptions(OptionsBinder &binder) {
static llvm::cl::OptionCategory category("VulkanSPIRV HAL Target");
binder.opt<std::string>(
+ // 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<std::string>(
- "iree-vulkan-target-env", targetEnv,
- llvm::cl::desc(
- "Vulkan target environment as #vk.target_env attribute assembly."));
binder.opt<bool>(
"iree-vulkan-experimental-indirect-bindings", indirectBindings,
llvm::cl::desc(
@@ -56,31 +53,6 @@
};
} // 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<Vulkan::TargetEnvAttr>(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 @@
MLIRContext *context, StringRef deviceID, DictionaryAttr deviceConfigAttr,
SmallVectorImpl<IREE::HAL::ExecutableTargetAttr> &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<NamedAttribute> 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 @@
}
void getDependentDialects(DialectRegistry ®istry) const override {
- registry.insert<IREE::Codegen::IREECodegenDialect, Vulkan::VulkanDialect,
- spirv::SPIRVDialect, gpu::GPUDialect>();
+ registry.insert<IREE::Codegen::IREECodegenDialect, spirv::SPIRVDialect,
+ gpu::GPUDialect>();
}
void
diff --git a/compiler/plugins/target/VulkanSPIRV/test/smoketest.mlir b/compiler/plugins/target/VulkanSPIRV/test/smoketest.mlir
index 68d6542..f8d8159 100644
--- a/compiler/plugins/target/VulkanSPIRV/test/smoketest.mlir
+++ b/compiler/plugins/target/VulkanSPIRV/test/smoketest.mlir
@@ -4,7 +4,9 @@
hal.device.targets = [
#hal.device.target<"vulkan", [
#hal.executable.target<"vulkan-spirv", "vulkan-spirv-fb", {
- spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #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 = [32, 32],
+ max_workgroup_sizes = [128, 128, 64], max_thread_count_per_workgroup = 128, max_workgroup_memory_bytes = 16384>>
}>
]>
]
diff --git a/compiler/plugins/target/WebGPUSPIRV/CMakeLists.txt b/compiler/plugins/target/WebGPUSPIRV/CMakeLists.txt
index d98dcf2..caf4460 100644
--- a/compiler/plugins/target/WebGPUSPIRV/CMakeLists.txt
+++ b/compiler/plugins/target/WebGPUSPIRV/CMakeLists.txt
@@ -48,6 +48,7 @@
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 8397eb1..0a37691 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 @@
}
};
-// 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 @@
MLIRContext *context, StringRef deviceID, DictionaryAttr deviceConfigAttr,
SmallVectorImpl<IREE::HAL::ExecutableTargetAttr> &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<NamedAttribute> 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<IREE::HAL::ExecutableTargetAttr>(
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 1a17240..31f361b 100644
--- a/compiler/plugins/target/WebGPUSPIRV/test/smoketest.mlir
+++ b/compiler/plugins/target/WebGPUSPIRV/test/smoketest.mlir
@@ -5,7 +5,9 @@
hal.device.targets = [
#hal.device.target<"webgpu", [
#hal.executable.target<"webgpu-spirv", "webgpu-wgsl-fb", {
- spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>>
+ iree.gpu.target = #iree_gpu.target<arch = "", features = "spirv:v1.0,cap:Shader,ext:SPV_KHR_storage_buffer_storage_class", wgp = <
+ compute = fp32|int32, storage = b32, subgroup = none, dot = none, mma = [], subgroup_size_choices = [32],
+ max_workgroup_sizes = [128, 128, 64], max_thread_count_per_workgroup = 128, max_workgroup_memory_bytes = 16384>>
}>
]>
]
diff --git a/compiler/src/iree/compiler/API/Internal/IREEReduceToolEntryPoint.cpp b/compiler/src/iree/compiler/API/Internal/IREEReduceToolEntryPoint.cpp
index 7e0d201..4cb61ec 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 LogicalResult ireeReduceMainFromCL(int argc, char **argv,
MLIRContext ®istry) {
- llvm::cl::OptionCategory ireeReduceCategory("iree-reduce options");
+ cl::OptionCategory ireeReduceCategory("iree-reduce options");
- llvm::cl::opt<std::string> testScript(cl::Positional, cl::Required,
- cl::desc("<test script>"),
- cl::cat(ireeReduceCategory));
+ cl::opt<std::string> testScript(cl::Positional, cl::Required,
+ cl::desc("<test script>"),
+ cl::cat(ireeReduceCategory));
cl::opt<std::string> inputFilename(cl::Positional, cl::desc("<input file>"),
cl::init("-"),
@@ -74,12 +74,11 @@
"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 5df8a36..d88dc84 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 <numeric>
#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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
rhs, acc)
.getResult();
}
- case MMAIntrinsic::WMMA_F16_16x16x16_F32: {
+ case MMAIntrinsic::WMMA_F16_16x16x16_F32:
+ case MMAIntrinsic::WMMA_F16_16x16x16_F16: {
return builder.create<amdgpu::WMMAOp>(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 5c1bead..a7abbb6 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 @@
let genSpecializedAttr = 0;
}
+// Format: <kind>_<input-type>_<M>x<N>x<K>_<output-type>
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 e02e759..b6f0aeb 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 <optional>
#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"
@@ -38,6 +39,8 @@
uint32_t mmaCount;
const MMAIntrinsic *mmaOps;
+ // We support two static values here mostly due to AMD RDNA GPUs have two
+ // modes. Use duplicated values if the GPU only have one subgroup size.
std::array<int32_t, 2> subgroupSizeChoices;
std::array<int32_t, 3> maxWorkgroupSizes;
uint32_t maxThreadSize;
@@ -88,9 +91,12 @@
mmaAttrs.push_back(MMAAttr::get(context, wgp->mmaOps[i]));
SmallVector<int32_t, 2> subgroupSizes;
+ assert(wgp->subgroupSizeChoices.front() != 0);
+ assert(wgp->subgroupSizeChoices.back() != 0);
subgroupSizes.push_back(wgp->subgroupSizeChoices.front());
- if (wgp->subgroupSizeChoices.back() != wgp->subgroupSizeChoices.front())
+ if (wgp->subgroupSizeChoices.back() != wgp->subgroupSizeChoices.front()) {
subgroupSizes.push_back(wgp->subgroupSizeChoices.back());
+ }
auto targetWgp = TargetWgpAttr::get(
context, ComputeBitwidthsAttr::get(context, details.wgp->compute),
@@ -156,6 +162,7 @@
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 +172,29 @@
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<TargetDetails> 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 +240,10 @@
.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 +251,140 @@
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<StringRef>(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());
+ .Default("");
+}
+
+//===----------------------------------------------------------------------===//
+// Known Apple target details
+//===----------------------------------------------------------------------===//
+
+std::optional<TargetDetails> 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() {
+ // Recent drivers report support for shaderInt64. Aside from not widely
+ // applicable, we don't know whether that's emulated and how performant it is.
+ // So exclude that for now.
+ ComputeBitwidths computeBitwdiths =
+ ComputeBitwidths::Int32 | ComputeBitwidths::Int16 |
+ ComputeBitwidths::Int8 | ComputeBitwidths::FP32 | ComputeBitwidths::FP16;
+ // clang-format off
+ static const WgpDetails valhallWgp = {
+ computeBitwdiths, allStorageBits, allSubgroupOps, allDotProductOps,
+ /*mmaCount=*/0, /*mmaOps=*/nullptr, {16, 16}, {512, 512, 512},
+ 512, 32 * 1024};
+ // clang-format on
+ return &valhallWgp;
+}
+
+std::optional<TargetDetails> 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<std::optional<TargetDetails>>(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<StringRef>(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;
}
@@ -329,7 +457,101 @@
.Case("turing", "sm_75")
.Case("volta", "sm_70") // Or sm_72; use smaller version
.Case("pascal", "sm_60") // Or sm_61/62; use smaller version
- .Default(StringRef());
+ .Default("");
+}
+
+//===----------------------------------------------------------------------===//
+// 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, 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<TargetDetails> 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<TargetDetails> getAndroidProfileDetails(StringRef target) {
+ const WgpDetails *baseline2022Wgp = getAndroidBaseline2022WgpDetails();
+
+ return llvm::StringSwitch<std::optional<TargetDetails>>(target.lower())
+ .Case("vp_android_baseline_2022", TargetDetails{baseline2022Wgp, nullptr})
+ .Default(std::nullopt);
}
} // namespace
@@ -338,9 +560,26 @@
// 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<TargetDetails> 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<TargetDetails> details = getAMDGPUTargetDetails(target)) {
return createTargetAttr(*details, normalizeAMDGPUTarget(target), features,
context);
}
@@ -351,16 +590,62 @@
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<TargetDetails> details = getAMDGPUTargetDetails(target)) {
+ return createTargetAttr(*details, normalizeAMDGPUTarget(target),
+ /*features=*/"spirv:v1.6,cap:Shader", context);
+ }
+ if (std::optional<TargetDetails> details = getARMGPUTargetDetails(target)) {
+ return createTargetAttr(*details, normalizeARMGPUTarget(target),
+ /*features=*/"spirv:v1.4,cap:Shader", context);
+ }
+ if (std::optional<TargetDetails> details =
+ getNVIDIAGPUTargetDetails(target)) {
return createTargetAttr(*details, normalizeNVIDIAGPUTarget(target),
- features, context);
+ /*features=*/"spirv:v1.6,cap:Shader", context);
+ }
+ if (std::optional<TargetDetails> 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<TargetDetails> 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 ffe9a15..d9698cc 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 @@
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 d73740f..a6415cd 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp
@@ -74,9 +74,6 @@
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 @@
} // namespace
+bool isROCmBackend(IREE::GPU::TargetAttr target) {
+ return target.getArch().starts_with("gfx");
+}
+
//====---------------------------------------------------------------------===//
// Matmul Configuration Helpers
//====---------------------------------------------------------------------===//
@@ -576,6 +577,10 @@
setVectorDistributionConfig(IREE::GPU::TargetAttr target,
mlir::FunctionOpInterface entryPoint,
Operation *computeOp) {
+ // We haven't properly plumbed through MMA op layouts and conversions for CUDA
+ // to target NVIDIA GPUs. So disable the vector distribution pass for it.
+ if (!isROCmBackend(target))
+ return failure();
if (!clGPUEnableVectorDistribution) {
LDBG("Vector Distribution not enabled, skipping...");
@@ -1188,15 +1193,6 @@
// 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 +1363,8 @@
//
// 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 5c80445..e72fdc5 100644
--- a/compiler/src/iree/compiler/Codegen/SPIRV/BUILD.bazel
+++ b/compiler/src/iree/compiler/Codegen/SPIRV/BUILD.bazel
@@ -58,6 +58,7 @@
"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 13632df..1378bbc 100644
--- a/compiler/src/iree/compiler/Codegen/SPIRV/CMakeLists.txt
+++ b/compiler/src/iree/compiler/Codegen/SPIRV/CMakeLists.txt
@@ -57,6 +57,7 @@
"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 538e3ba..75eb17d 100644
--- a/compiler/src/iree/compiler/Codegen/SPIRV/Passes.cpp
+++ b/compiler/src/iree/compiler/Codegen/SPIRV/Passes.cpp
@@ -631,6 +631,8 @@
void buildSPIRVCodegenConfigurationPassPipeline(
OpPassManager &variantPassManager) {
+ // TODO: move the following pass to be immediately before ConvertToSPIRVPass.
+ variantPassManager.addPass(createSPIRVConvertGPUTargetPass());
OpPassManager &modulePassManager = variantPassManager.nest<ModuleOp>();
buildSPIRVCodegenConfigurationPassPipelineImpl(modulePassManager);
}
diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/Passes.h b/compiler/src/iree/compiler/Codegen/SPIRV/Passes.h
index e1dc830..a0b0d16 100644
--- a/compiler/src/iree/compiler/Codegen/SPIRV/Passes.h
+++ b/compiler/src/iree/compiler/Codegen/SPIRV/Passes.h
@@ -86,6 +86,10 @@
std::unique_ptr<InterfacePass<FunctionOpInterface>>
createSPIRVBreakDownLargeVectorPass();
+// Converts #iree_gpu.target into #spirv.target_env.
+std::unique_ptr<OperationPass<IREE::HAL::ExecutableVariantOp>>
+createSPIRVConvertGPUTargetPass();
+
/// Emulates bfloat 16 ops with 32-bit float ops.
std::unique_ptr<InterfacePass<FunctionOpInterface>>
createSPIRVEmulateBf16Pass();
diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/Passes.td b/compiler/src/iree/compiler/Codegen/SPIRV/Passes.td
index 29b396d..dc94eb2 100644
--- a/compiler/src/iree/compiler/Codegen/SPIRV/Passes.td
+++ b/compiler/src/iree/compiler/Codegen/SPIRV/Passes.td
@@ -33,6 +33,13 @@
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 0000000..fc9cfe3
--- /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<Version> deduceVersion(ArrayRef<StringRef> features) {
+ for (StringRef feature : features) {
+ if (feature.consume_front("spirv:v1.")) {
+ return llvm::StringSwitch<std::optional<Version>>(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<Version> processCapabilities(ArrayRef<StringRef> features,
+ SetVector<Capability> &caps) {
+ for (StringRef feature : features) {
+ if (feature.consume_front("cap:")) {
+ if (std::optional<Capability> 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<Version> processExtensions(ArrayRef<StringRef> features,
+ SetVector<Extension> &exts) {
+ for (StringRef feature : features) {
+ if (feature.consume_front("ext:")) {
+ if (std::optional<Extension> ext = spirv::symbolizeExtension(feature))
+ exts.insert(*ext);
+ }
+ }
+ return std::nullopt;
+}
+
+//===----------------------------------------------------------------------===//
+// Client API and vendor
+//===----------------------------------------------------------------------===//
+
+ClientAPI deduceClientAPI(StringRef backend) {
+ return llvm::StringSwitch<ClientAPI>(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("valhall"))
+ 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<Capability> &caps,
+ SetVector<Extension> &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<Capability> &caps,
+ SetVector<Extension> &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<Capability> &caps,
+ SetVector<Extension> &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<Capability> &caps,
+ SetVector<Extension> &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<Capability> &caps, SetVector<Extension> &exts,
+ SetVector<Attribute> &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<Attribute, 4> 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<int> 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<spirv::TargetEnvAttr>
+convertGPUTarget(IREE::HAL::ExecutableVariantOp variant) {
+ IREE::HAL::ExecutableTargetAttr target = variant.getTarget();
+ IREE::GPU::TargetAttr gpuTarget = getGPUTargetAttr(target);
+
+ SmallVector<StringRef> features;
+ llvm::SplitString(gpuTarget.getFeatures(), features, ",");
+
+ SetVector<Capability> caps;
+ SetVector<Extension> exts;
+ SetVector<Attribute> coopMatAttrs;
+
+ std::optional<Version> 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<SPIRVConvertGPUTargetPass> {
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ registry.insert<spirv::SPIRVDialect>();
+ }
+
+ void runOnOperation() override {
+ IREE::HAL::ExecutableVariantOp variant = getOperation();
+ IREE::HAL::ExecutableTargetAttr target = variant.getTarget();
+
+ FailureOr<spirv::TargetEnvAttr> 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<OperationPass<IREE::HAL::ExecutableVariantOp>>
+createSPIRVConvertGPUTargetPass() {
+ return std::make_unique<SPIRVConvertGPUTargetPass>();
+}
+
+} // 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 ed6a8be..eb59474 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 @@
"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 3dc8277..273e581 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 @@
"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 0000000..b1f8092
--- /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<arch = "rdna3", features = "spirv:v1.6,cap:Shader",
+ wgp = <compute = fp64|fp32|fp16|int64|int32|int16|int8, storage = b64|b32|b16|b8, subgroup = shuffle|arithmetic, dot = dp4xi8toi32, mma = [<WMMA_F16_16x16x16_F32>, <WMMA_F16_16x16x16_F16>],
+ 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<push_constants = 0, sets = [<0, bindings = [<0, storage_buffer>]>]>) {
+ ^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<v1.6,
+// CHECK-SAME: [Shader, Float64, Float16, Int64, Int16, Int8,
+// CHECK-SAME: StorageBuffer16BitAccess, StorageUniform16, StoragePushConstant16,
+// CHECK-SMAE: StorageBuffer8BitAccess, UniformAndStorageBuffer8BitAccess, StoragePushConstant8,
+// CHECK-SAME: GroupNonUniformShuffle, GroupNonUniformShuffleRelative, GroupNonUniformArithmetic,
+// CHECK-SAME: DotProduct, DotProductInput4x8BitPacked, DotProductInputAll, DotProductInput4x8Bit,
+// CHECK-SAME: CooperativeMatrixKHR],
+// CHECK-SAME: [SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_integer_dot_product, SPV_KHR_cooperative_matrix]>,
+// CHECK-SAME: AMD,
+// CHECK-SAME: #spirv.resource_limits<max_compute_shared_memory_size = 65536,
+// CHECK-SAME: max_compute_workgroup_invocations = 1024, max_compute_workgroup_size = [1024 : i32, 1024 : i32, 1024 : i32],
+// CHECK-SAME: subgroup_size = 64, min_subgroup_size = 32, max_subgroup_size = 64,
+// CHECK-SAME: cooperative_matrix_properties_khr = [
+// CHECK-SAME: #spirv.coop_matrix_props_khr<m_size = 16, n_size = 16, k_size = 16, a_type = f16, b_type = f16, c_type = f32, result_type = f32, acc_sat = false, scope = <Subgroup>>,
+// CHECK-SAME: #spirv.coop_matrix_props_khr<m_size = 16, n_size = 16, k_size = 16, a_type = f16, b_type = f16, c_type = f16, result_type = f16, acc_sat = false, scope = <Subgroup>>
+// 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 236a474..0000000
--- 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 487e4f1..0000000
--- 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 da4b65e..0000000
--- 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 3b03c56..0000000
--- 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 dc33c2b..0000000
--- 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<Attribute, Attribute, Attribute, spirv::Vendor,
- spirv::DeviceType, uint32_t, Attribute>;
-
- 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>())
- 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<Extension> extensions,
- spirv::Vendor vendorID,
- spirv::DeviceType deviceType,
- uint32_t deviceID,
- CapabilitiesAttr capabilities) {
- mlir::Builder builder(capabilities.getContext());
- llvm::SmallVector<Attribute, 0> extAttrs;
- extAttrs.reserve(extensions.size());
- for (auto ext : extensions) {
- extAttrs.push_back(ExtensionAttr::get(builder.getContext(), ext));
- }
- return get(builder.getI32IntegerAttr(static_cast<uint32_t>(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<Version>(
- llvm::cast<IntegerAttr>(getImpl()->version).getValue().getZExtValue());
-}
-
-unsigned TargetEnvAttr::getRevision() {
- return llvm::cast<IntegerAttr>(getImpl()->revision).getValue().getZExtValue();
-}
-
-TargetEnvAttr::ext_iterator::ext_iterator(ArrayAttr::iterator it)
- : llvm::mapped_iterator<ArrayAttr::iterator, Extension (*)(Attribute)>(
- it, [](Attribute attr) {
- return llvm::cast<ExtensionAttr>(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<ArrayAttr>(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<CapabilitiesAttr>(getImpl()->capabilities);
-}
-
-LogicalResult
-TargetEnvAttr::verify(function_ref<InFlightDiagnostic()> 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<LogicalResult(llvm::SMLoc, StringRef)> 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<uint32_t>(*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<Attribute, 1> extensions;
- llvm::SMLoc errorloc;
- StringRef errorKeyword;
-
- MLIRContext *context = parser.getContext();
- auto processExtension = [&](llvm::SMLoc loc, StringRef extension) {
- if (std::optional<Extension> 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<TargetEnvAttr>(attr))
- return print(targetEnv, printer);
-
- assert(false && "unhandled Vulkan attribute kind");
-}
-
-//===----------------------------------------------------------------------===//
-// Registration
-//===----------------------------------------------------------------------===//
-
-void VulkanDialect::registerAttributes() {
- addAttributes<TargetEnvAttr,
-#define GET_ATTRDEF_LIST
-#include "iree/compiler/Dialect/Vulkan/IR/VulkanAttributes.cpp.inc"
- >();
-}
-
-} // 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 1175db6..0000000
--- 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<TargetEnvAttr, Attribute,
- detail::TargetEnvAttributeStorage> {
-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<Extension> 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<ArrayAttr::iterator,
- Extension (*)(Attribute)> {
- explicit ext_iterator(ArrayAttr::iterator it);
- };
- using ext_range = llvm::iterator_range<ext_iterator>;
-
- /// 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<InFlightDiagnostic()> 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 fcd0ccf..0000000
--- 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<string attrName, string attrMnemonic>
- : AttrDef<VK_Dialect, attrName> {
- 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<int>">:$minSubgroupSize,
- OptionalParameter<"::std::optional<int>">:$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 c256111..0000000
--- 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<string name> :
- CPred<"::mlir::iree_compiler::IREE::Vulkan::symbolize" # name # "("
- "cast<IntegerAttr>($_self).getValue().getZExtValue()).hasValue()">;
-class VK_IsKnownIntEnumCaseFor<string name> :
- CPred<"::mlir::iree_compiler::IREE::Vulkan::symbolize" # name # "("
- "cast<IntegerAttr>($_self).getValue().getZExtValue()).hasValue()">;
-
-// Wrapper over base I32BitEnumAttr to set common fields.
-class VK_BitEnumAttr<string name, string description,
- list<I32BitEnumAttrCase> cases> :
- I32BitEnumAttr<name, description, cases> {
- let predicate = And<[I32Attr.predicate, VK_IsKnownBitEnumCaseFor<name>]>;
- let cppNamespace = "::mlir::iree_compiler::IREE::Vulkan";
-}
-
-class VK_I32Enum<string name, string description, list<I32EnumAttrCase> cases> :
- I32EnumAttr<name, description, cases> {
- let predicate = And<[I32Attr.predicate, VK_IsKnownIntEnumCaseFor<name>]>;
- let cppNamespace = "::mlir::iree_compiler::IREE::Vulkan";
-}
-
-class VK_I32EnumAttr<string name, string description, string mnemonic,
- list<I32EnumAttrCase> cases> :
- EnumAttr<VK_Dialect, I32EnumAttr<name, description, cases>, 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 2e78feb..0000000
--- 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<VulkanDialect>()) {
- 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 9cb3d01..0000000
--- 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 fc67767..0000000
--- 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 2422a85..0000000
--- 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 bbddf7d..0000000
--- 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 cebe847..0000000
--- 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 343f1aa..0000000
--- 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<v1.1, r(120), [VK_KHR_spirv_1_4, VK_KHR_storage_buffer_storage_class], AMD:DiscreteGPU, #vk.caps<
- // CHECK-SAME: maxComputeSharedMemorySize = 16384,
- // CHECK-SAME: maxComputeWorkGroupInvocations = 1024,
- // CHECK-SAME: maxComputeWorkGroupSize = dense<[128, 8, 4]> : vector<3xi32>
- // CHECK-SAME: subgroupFeatures = 63 : i32,
- // CHECK-SAME: subgroupSize = 4
- // CHECK-SAME: >>
- target_env = #vk.target_env<v1.1, r(120), [VK_KHR_spirv_1_4, VK_KHR_storage_buffer_storage_class], AMD:DiscreteGPU, #vk.caps<
- maxComputeSharedMemorySize = 16384,
- maxComputeWorkGroupInvocations = 1024,
- maxComputeWorkGroupSize = dense<[128, 8, 4]>: 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<mSize = 8, nSize = 8, kSize = 32,
- // CHECK-SAME: aType = i8, bType = i8, cType = i32, resultType = i32,
- // CHECK-SAME: accSat = false, scope = <Subgroup>>
- // CHECK-SAME: #vk.coop_matrix_props<mSize = 8, nSize = 8, kSize = 16,
- // CHECK-SAME: aType = f16, bType = f16, cType = f16, resultType = f16,
- // CHECK-SAME: accSat = false, scope = <Subgroup>>
- target_env =
- #vk.target_env<v1.2, r(133),
- [VK_KHR_storage_buffer_storage_class, VK_KHR_cooperative_matrix],
- NVIDIA:DiscreteGPU,
- #vk.caps<maxComputeSharedMemorySize = 49152,
- maxComputeWorkGroupInvocations = 1024,
- maxComputeWorkGroupSize = dense<[2147483647, 65535, 65535]> : 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<Subgroup>>,
- #vk.coop_matrix_props<
- mSize = 8, nSize = 8, kSize = 16,
- aType = f16, bType = f16, cType = f16, resultType = f16,
- accSat = false, scope = #vk.scope<Subgroup>>
- ]
- >>
-} : () -> ()
-
-// -----
-
-"vk_configure_op"() {
- // CHECK: Qualcomm:IntegratedGPU:100925441
- // CHECK-SAME: shaderFloat64
- // CHECK-SAME: shaderInt16
- target_env = #vk.target_env<v1.1, r(120), [VK_KHR_spirv_1_4, VK_KHR_storage_buffer_storage_class], Qualcomm:IntegratedGPU:0x6040001, #vk.caps<
- maxComputeSharedMemorySize = 16384,
- maxComputeWorkGroupInvocations = 1024,
- maxComputeWorkGroupSize = dense<[128, 8, 4]>: 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<v10.8, r(0), [], #vk.caps<
- maxComputeWorkGroupInvocations = 128,
- maxComputeWorkGroupSize = dense<[64, 4, 4]>: vector<3xi32>
- >>
-} : () -> ()
-
-// -----
-
-"unknown_vulkan_extension"() {
- // expected-error @+1 {{unknown Vulkan extension: VK_KHR_something}}
- target_env = #vk.target_env<v1.0, r(10), [VK_KHR_something], #vk.caps<
- maxComputeWorkGroupInvocations = 128,
- maxComputeWorkGroupSize = dense<[64, 4, 4]>: vector<3xi32>
- >>
-} : () -> ()
-
-// -----
-
-"wrong_vendor_id"() {
- // expected-error @+1 {{unknown vendor: AVendor}}
- target_env = #vk.target_env<v1.0, r(10), [], AVendor:Unknown, #vk.caps<
- maxComputeSharedMemorySize = 16384,
- maxComputeWorkGroupInvocations = 1024,
- maxComputeWorkGroupSize = dense<[128, 8, 4]>: vector<3xi32>,
- subgroupFeatures = 63: i32,
- subgroupSize = 4
- >>
-} : () -> ()
-
-// -----
-
-"wrong_device_type"() {
- // expected-error @+1 {{unknown device type: ADeviceType}}
- target_env = #vk.target_env<v1.0, r(10), [], NVIDIA:ADeviceType, #vk.caps<
- maxComputeSharedMemorySize = 16384,
- maxComputeWorkGroupInvocations = 1024,
- maxComputeWorkGroupSize = dense<[128, 8, 4]>: vector<3xi32>,
- subgroupFeatures = 63: i32,
- subgroupSize = 4
- >>
-} : () -> ()
-
-// -----
-
-"missing_core_1_1_properties_field"() {
- target_env = #vk.target_env<v1.0, r(10), [], Unknown:Unknown, #vk.caps<
- maxComputeWorkGroupInvocations = 128
- // expected-error @+1 {{struct is missing required parameter: maxComputeSharedMemorySize}}
- >>
-} : () -> ()
-
-// -----
-
-"unknown_core_1_1_properties_field"() {
- target_env = #vk.target_env<v1.0, r(10), [], Unknown:Unknown, #vk.caps<
- maxComputeSharedMemorySize = 16384,
- maxComputeWorkGroupInvocations = 128,
- maxComputeWorkGroupSize = dense<[64, 4, 4]>: vector<3xi32>,
- // expected-error @+1 {{duplicate or unknown struct parameter name: moreStuff}}
- moreStuff = 8: i32
- >>
-} : () -> ()
-
-// -----
-
-"wrong_subgroup_bit"() {
- target_env = #vk.target_env<v1.0, r(10), [], Unknown:Unknown, #vk.caps<
- maxComputeSharedMemorySize = 16384,
- maxComputeWorkGroupInvocations = 1024,
- maxComputeWorkGroupSize = dense<[128, 8, 4]>: 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 cbbd06f..0000000
--- 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 8435767..0000000
--- 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 bcf3b55..0000000
--- 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<spirv::Extension> &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<spirv::Capability> &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<Attribute, 1> khrCoopAttrs;
- if (ArrayAttr attr = vkCapabilities.getCooperativeMatrixPropertiesKHR()) {
- for (auto props :
- attr.getAsRange<Vulkan::CooperativeMatrixPropertiesKHRAttr>()) {
- auto scope = static_cast<spirv::Scope>(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<int32_t>();
- SmallVector<int64_t> 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<spirv::Extension> spvExtensions;
- convertExtensions(vkTargetEnv, spvExtensions);
-
- SmallVector<spirv::Capability, 8> 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 cc1d62a..0000000
--- 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 9564bf7..0000000
--- 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<Extension> &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<int, 3> maxComputeWorkGroupSize = {128, 128, 64};
-
- int subgroupSize = 32;
- SubgroupFeature subgroupFeatures = SubgroupFeature::Basic;
- std::optional<int> 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<Attribute> 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<llvm::StringRef, 3> 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<Extension> 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 7ea5e0d..0000000
--- 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 <string>
-
-#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:
-///
-/// <vendor/arch>-<product>-<os>
-///
-/// 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 687fa49..0000000
--- 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 bb5cbe5..0000000
--- 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 3a23031..0000000
--- 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<v1.1, r(120), [VK_KHR_spirv_1_4, VK_KHR_storage_buffer_storage_class], AMD:DiscreteGPU, #vk.caps<maxComputeSharedMemorySize = 16384, maxComputeWorkGroupInvocations = 1024, maxComputeWorkGroupSize = dense<[128, 8, 4]>: 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<v1.3,
-// DEFAULT-SAME: [Shader, GroupNonUniform], [SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers]>,
-// DEFAULT-SAME: api=Vulkan, #spirv.resource_limits<max_compute_workgroup_size = [128, 128, 64], subgroup_size = 64, cooperative_matrix_properties_khr = []>>
-
-// ADRENO: #spirv.target_env<#spirv.vce<v1.4,
-// ADRENO-SAME: [Shader, Float16, Int16, Int8, StorageBuffer16BitAccess, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative, GroupNonUniformQuad, VariablePointers, VariablePointersStorageBuffer],
-// ADRENO-SAME: [SPV_KHR_16bit_storage, SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers]>,
-// ADRENO-SAME: api=Vulkan, Qualcomm:IntegratedGPU, #spirv.resource_limits<max_compute_shared_memory_size = 32768, max_compute_workgroup_invocations = 1024, max_compute_workgroup_size = [1024, 1024, 64], subgroup_size = 64, cooperative_matrix_properties_khr = []>>
-
-// VALHALL: #spirv.target_env<#spirv.vce<v1.4,
-// VALHALL-SAME: [Shader, Float16, Int16, Int8, StorageBuffer16BitAccess, StorageUniform16, StoragePushConstant16, StorageBuffer8BitAccess, UniformAndStorageBuffer8BitAccess, StoragePushConstant8, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative, GroupNonUniformClustered, GroupNonUniformQuad, VariablePointers, VariablePointersStorageBuffer, DotProduct, DotProductInputAll, DotProductInput4x8BitPacked, DotProductInput4x8Bit],
-// VALHALL-SAME: [SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_integer_dot_product, SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers]>,
-// VALHALL-SAME: api=Vulkan, ARM:IntegratedGPU, #spirv.resource_limits<max_compute_shared_memory_size = 32768, max_compute_workgroup_invocations = 512, max_compute_workgroup_size = [512, 512, 512], subgroup_size = 16, cooperative_matrix_properties_khr = []>>
-
-// TURING: #spirv.target_env<#spirv.vce<v1.6,
-// TURING-SAME: [Shader, Float64, Float16, Int64, Int16, Int8, StorageBuffer16BitAccess, StorageUniform16, StoragePushConstant16, StorageBuffer8BitAccess, UniformAndStorageBuffer8BitAccess, StoragePushConstant8, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative, GroupNonUniformClustered, GroupNonUniformQuad, PhysicalStorageBufferAddresses, VariablePointers, VariablePointersStorageBuffer, DotProduct, DotProductInputAll, DotProductInput4x8BitPacked, DotProductInput4x8Bit, CooperativeMatrixKHR],
-// TURING-SAME: [SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_integer_dot_product, SPV_KHR_storage_buffer_storage_class, SPV_KHR_physical_storage_buffer, SPV_KHR_variable_pointers, SPV_KHR_cooperative_matrix]>,
-// TURING-SAME: api=Vulkan, NVIDIA:DiscreteGPU, #spirv.resource_limits<max_compute_shared_memory_size = 49152, max_compute_workgroup_invocations = 1024, max_compute_workgroup_size = [1024, 1024, 64], min_subgroup_size = 32, max_subgroup_size = 32, cooperative_matrix_properties_khr = [#spirv.coop_matrix_props_khr<m_size = 8, n_size = 8, k_size = 32, a_type = i8, b_type = i8, c_type = i32, result_type = i32, acc_sat = false, scope = <Subgroup>>, #spirv.coop_matrix_props_khr<m_size = 16, n_size = 16, k_size = 16, a_type = f16, b_type = f16, c_type = f16, result_type = f16, acc_sat = false, scope = <Subgroup>>, #spirv.coop_matrix_props_khr<m_size = 16, n_size = 16, k_size = 16, a_type = f16, b_type = f16, c_type = f32, result_type = f32, acc_sat = false, scope = <Subgroup>>]>>
-
-// RDNA1: #spirv.target_env<#spirv.vce<v1.6,
-// RDNA1-SAME: [Shader, Float64, Float16, Int64, Int16, Int8, StorageBuffer16BitAccess, StorageUniform16, StoragePushConstant16, StorageBuffer8BitAccess, UniformAndStorageBuffer8BitAccess, StoragePushConstant8, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative, GroupNonUniformClustered, GroupNonUniformQuad, PhysicalStorageBufferAddresses, VariablePointers, VariablePointersStorageBuffer, DotProduct, DotProductInputAll, DotProductInput4x8BitPacked, DotProductInput4x8Bit],
-// RDNA1-SAME: [SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_integer_dot_product, SPV_KHR_storage_buffer_storage_class, SPV_KHR_physical_storage_buffer, SPV_KHR_variable_pointers]>,
-// RDNA1-SAME: api=Vulkan, AMD:DiscreteGPU, #spirv.resource_limits<max_compute_shared_memory_size = 65536, max_compute_workgroup_invocations = 1024, max_compute_workgroup_size = [1024, 1024, 1024], subgroup_size = 64, min_subgroup_size = 32, max_subgroup_size = 64, cooperative_matrix_properties_khr = []>>
-
-// RDNA3: #spirv.target_env<#spirv.vce<v1.6,
-// RDNA3-SAME: [Shader, Float64, Float16, Int64, Int16, Int8, StorageBuffer16BitAccess, StorageUniform16, StoragePushConstant16, StorageBuffer8BitAccess, UniformAndStorageBuffer8BitAccess, StoragePushConstant8, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative, GroupNonUniformClustered, GroupNonUniformQuad, PhysicalStorageBufferAddresses, VariablePointers, VariablePointersStorageBuffer, DotProduct, DotProductInputAll, DotProductInput4x8BitPacked, DotProductInput4x8Bit, CooperativeMatrixKHR],
-// RDNA3-SAME: [SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_integer_dot_product, SPV_KHR_storage_buffer_storage_class, SPV_KHR_physical_storage_buffer, SPV_KHR_variable_pointers, SPV_KHR_cooperative_matrix]>,
-// RDNA3-SAME: api=Vulkan, AMD:DiscreteGPU, #spirv.resource_limits<max_compute_shared_memory_size = 65536, max_compute_workgroup_invocations = 1024, max_compute_workgroup_size = [1024, 1024, 1024], subgroup_size = 64, min_subgroup_size = 32, max_subgroup_size = 64, cooperative_matrix_properties_khr = [#spirv.coop_matrix_props_khr<m_size = 16, n_size = 16, k_size = 16, a_type = i8, b_type = i8, c_type = i32, result_type = i32, acc_sat = false, scope = <Subgroup>>, #spirv.coop_matrix_props_khr<m_size = 16, n_size = 16, k_size = 16, a_type = f16, b_type = f16, c_type = f16, result_type = f16, acc_sat = false, scope = <Subgroup>>, #spirv.coop_matrix_props_khr<m_size = 16, n_size = 16, k_size = 16, a_type = f16, b_type = f16, c_type = f32, result_type = f32, acc_sat = false, scope = <Subgroup>>]>>
-
-// M1: #spirv.target_env<#spirv.vce<v1.3,
-// M1-SAME: [Shader, Float16, Int64, Int16, Int8, StorageBuffer16BitAccess, StorageUniform16, StoragePushConstant16, StorageBuffer8BitAccess, UniformAndStorageBuffer8BitAccess, StoragePushConstant8, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative, GroupNonUniformQuad, PhysicalStorageBufferAddresses, VariablePointers, VariablePointersStorageBuffer],
-// M1-SAME: [SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_storage_buffer_storage_class, SPV_KHR_physical_storage_buffer, SPV_KHR_variable_pointers]>,
-// M1-SAME: api=Vulkan, Apple:IntegratedGPU, #spirv.resource_limits<max_compute_shared_memory_size = 32768, max_compute_workgroup_invocations = 1024, max_compute_workgroup_size = [1024, 1024, 1024], cooperative_matrix_properties_khr = []>>
-
-// ARC: #spirv.target_env<#spirv.vce<v1.6,
-// ARC-SAME: [Shader, Float16, Int16, Int8, StorageBuffer16BitAccess, StorageUniform16, StoragePushConstant16, StorageBuffer8BitAccess, UniformAndStorageBuffer8BitAccess, StoragePushConstant8, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative, GroupNonUniformClustered, GroupNonUniformQuad, PhysicalStorageBufferAddresses, VariablePointers, VariablePointersStorageBuffer, DotProduct, DotProductInputAll, DotProductInput4x8BitPacked, DotProductInput4x8Bit],
-// ARC-SAME: [SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_integer_dot_product, SPV_KHR_storage_buffer_storage_class, SPV_KHR_physical_storage_buffer, SPV_KHR_variable_pointers]>,
-// ARC-SAME: api=Vulkan, Intel:DiscreteGPU, #spirv.resource_limits<max_compute_shared_memory_size = 32768, max_compute_workgroup_invocations = 1024, max_compute_workgroup_size = [1024, 1024, 64], min_subgroup_size = 8, max_subgroup_size = 32, cooperative_matrix_properties_khr = []>>}>
-
-// PASCAL: #spirv.target_env<#spirv.vce<v1.6,
-// PASCAL-SAME: [Shader, Float64, Float16, Int64, Int16, Int8, StorageBuffer16BitAccess, StorageUniform16, StoragePushConstant16, StorageBuffer8BitAccess, UniformAndStorageBuffer8BitAccess, StoragePushConstant8, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative, GroupNonUniformClustered, GroupNonUniformQuad, PhysicalStorageBufferAddresses, VariablePointers, VariablePointersStorageBuffer, DotProduct, DotProductInputAll, DotProductInput4x8BitPacked, DotProductInput4x8Bit],
-// PASCAL-SAME: [SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_integer_dot_product, SPV_KHR_storage_buffer_storage_class, SPV_KHR_physical_storage_buffer, SPV_KHR_variable_pointers, SPV_KHR_cooperative_matrix]>,
-// PASCAL-SAME: api=Vulkan, NVIDIA:DiscreteGPU, #spirv.resource_limits<max_compute_shared_memory_size = 49152, max_compute_workgroup_invocations = 1536, max_compute_workgroup_size = [1536, 1024, 64], min_subgroup_size = 32, max_subgroup_size = 32, cooperative_matrix_properties_khr = []>>}>
-
-// ENV: #spirv.target_env<#spirv.vce<v1.4,
-// ENV-SAME: [Shader, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative],
-// ENV-SAME: [SPV_KHR_storage_buffer_storage_class]>,
-// ENV-SAME: api=Vulkan, AMD:DiscreteGPU, #spirv.resource_limits<max_compute_workgroup_invocations = 1024, max_compute_workgroup_size = [128, 8, 4], subgroup_size = 4, cooperative_matrix_properties_khr = []>>
-
-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<readonly:tensor<16xf32>>
- %arg1 = stream.binding.subspan %arg1_binding[%c0] : !stream.binding -> !flow.dispatch.tensor<writeonly:tensor<f32>>
- %0 = tensor.empty() : tensor<f32>
- %1 = flow.dispatch.tensor.load %arg0, offsets=[0], sizes=[16], strides=[1] : !flow.dispatch.tensor<readonly:tensor<16xf32>> -> tensor<16xf32>
- %3 = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> ()>], iterator_types = ["reduction"]} ins(%1 : tensor<16xf32>) outs(%0 : tensor<f32>) {
- ^bb0(%arg2: f32, %arg3: f32):
- %4 = arith.addf %arg2, %arg3 : f32
- linalg.yield %4 : f32
- } -> tensor<f32>
- flow.dispatch.tensor.store %3, %arg1, offsets=[], sizes=[], strides=[] : tensor<f32> -> !flow.dispatch.tensor<writeonly:tensor<f32>>
- return
- }
- }
-}
diff --git a/compiler/src/iree/compiler/Tools/BUILD.bazel b/compiler/src/iree/compiler/Tools/BUILD.bazel
index 314a09c..95c620c 100644
--- a/compiler/src/iree/compiler/Tools/BUILD.bazel
+++ b/compiler/src/iree/compiler/Tools/BUILD.bazel
@@ -56,7 +56,6 @@
"//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 a38a4db..ee8d820 100644
--- a/compiler/src/iree/compiler/Tools/CMakeLists.txt
+++ b/compiler/src/iree/compiler/Tools/CMakeLists.txt
@@ -53,7 +53,6 @@
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 0472723..4d33387 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 @@
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 d655b98..ef10fb7 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<v1.3, [Shader, GroupNonUniform], [SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers]>,
- #spirv.resource_limits<max_compute_workgroup_size = [128, 128, 64], subgroup_size = 64>
+ 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 5cdbcac..36912bb 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<v1.3, [Shader, GroupNonUniform], [SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers]>,
- #spirv.resource_limits<max_compute_workgroup_size = [128, 128, 64], subgroup_size = 64>
+ 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 3766a30..b4885a0 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<v1.3, [Shader, GroupNonUniform, GroupNonUniformArithmetic, GroupNonUniformBallot],
- [SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers]>,
- #spirv.resource_limits<max_compute_workgroup_size = [128, 128, 64], subgroup_size = 64>
+ 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 70ad898..5bcdafe 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<v1.3, [Shader, GroupNonUniform, GroupNonUniformArithmetic, GroupNonUniformBallot],
- [SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers]>,
- #spirv.resource_limits<max_compute_workgroup_size = [128, 128, 64], subgroup_size = 64>
+ 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 13128e1..585bb25 100644
--- a/samples/transform_dialect/example_module.mlir
+++ b/samples/transform_dialect/example_module.mlir
@@ -25,19 +25,21 @@
// }
// }
-#target_env = #spirv.target_env<#spirv.vce<v1.3, [Shader, GroupNonUniform], [SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers]>, api=Vulkan, #spirv.resource_limits<max_compute_workgroup_size = [128, 128, 64], subgroup_size = 64, cooperative_matrix_properties_khr = []>>
+#target = #iree_gpu.target<arch = "", features = "spirv:v1.3,cap:Shader", wgp = <
+ compute = fp32|int32, storage = b32, subgroup = none, dot = none, mma = [], subgroup_size_choices = [64, 64],
+ max_workgroup_sizes = [128, 128, 64], max_thread_count_per_workgroup = 128, max_workgroup_memory_bytes = 16384>>
module attributes {
hal.device.targets = [
#hal.device.target<"vulkan", [
#hal.executable.target<"vulkan-spirv", "vulkan-spirv-fb", {
- spirv.target_env = #target_env
+ iree.gpu.target = #target
}>
]>
]
} {
hal.executable private @example_module_dispatch_0 {
- hal.executable.variant public @vulkan_spirv_fb target(<"vulkan-spirv", "vulkan-spirv-fb", {spirv.target_env = #target_env}>) {
+ hal.executable.variant public @vulkan_spirv_fb target(<"vulkan-spirv", "vulkan-spirv-fb", {iree.gpu.target = #target}>) {
hal.executable.export public @example_module_dispatch_0_generic_80_f32 ordinal(0) layout(
#hal.pipeline.layout<push_constants = 0, sets = [<0, bindings = [<0, storage_buffer, ReadOnly>, <1, storage_buffer>]>]>) {
^bb0(%arg0: !hal.device):
@@ -63,7 +65,7 @@
}
}
hal.executable private @example_module_dispatch_1 {
- hal.executable.variant public @vulkan_spirv_fb target(<"vulkan-spirv", "vulkan-spirv-fb", {spirv.target_env = #target_env}>) {
+ hal.executable.variant public @vulkan_spirv_fb target(<"vulkan-spirv", "vulkan-spirv-fb", {iree.gpu.target = #target}>) {
hal.executable.export public @example_module_dispatch_1_matmul_16x16x5_f32 ordinal(0) layout(
#hal.pipeline.layout<push_constants = 0, sets = [<0, bindings = [<0, storage_buffer, ReadOnly>, <1, storage_buffer, ReadOnly>, <2, storage_buffer>]>]>) {
^bb0(%arg0: !hal.device):
@@ -87,7 +89,7 @@
}
}
hal.executable private @example_module_dispatch_2 {
- hal.executable.variant public @vulkan_spirv_fb target(<"vulkan-spirv", "vulkan-spirv-fb", {spirv.target_env = #target_env}>) {
+ hal.executable.variant public @vulkan_spirv_fb target(<"vulkan-spirv", "vulkan-spirv-fb", {iree.gpu.target = #target}>) {
hal.executable.export public @example_module_dispatch_2_generic_16x16_f32 ordinal(0) layout(
#hal.pipeline.layout<push_constants = 0, sets = [<0, bindings = [<0, storage_buffer, ReadOnly>, <1, storage_buffer>]>]>) {
^bb0(%arg0: !hal.device):
