[vulkan] Update various targets to match current status (#10283)
* Used Android API level (e.g., 30) instead of the Android version
(e.g., 11). This makes it consistent with how we specify Android
targets on CPU side (e.g., aarch64-none-linux-android30).
(Android API level to version mapping: https://apilevels.com/)
* Updated capabilties and extensions for mobile GPUs on Android 12.
* Changed benchmarking to track Android 12 (which is what the
devices have for a long time).
diff --git a/benchmarks/TFLite/android-adreno.cmake b/benchmarks/TFLite/android-adreno.cmake
index ff540cc..d737ece 100644
--- a/benchmarks/TFLite/android-adreno.cmake
+++ b/benchmarks/TFLite/android-adreno.cmake
@@ -18,7 +18,7 @@
set(ANDROID_ADRENO_GPU_COMPILATION_FLAGS
"--iree-input-type=tosa"
- "--iree-vulkan-target-triple=adreno-unknown-android11"
+ "--iree-vulkan-target-triple=adreno-unknown-android31"
)
# GPU, Vulkan, Adreno, full-inference
diff --git a/benchmarks/TFLite/android-mali.cmake b/benchmarks/TFLite/android-mali.cmake
index 805db83..766b5ff 100644
--- a/benchmarks/TFLite/android-mali.cmake
+++ b/benchmarks/TFLite/android-mali.cmake
@@ -18,7 +18,7 @@
set(ANDROID_MALI_GPU_COMPILATION_FLAGS
"--iree-input-type=tosa"
- "--iree-vulkan-target-triple=valhall-unknown-android11"
+ "--iree-vulkan-target-triple=valhall-unknown-android31"
)
# GPU, Vulkan, Mali, full-inference
@@ -139,7 +139,7 @@
COMPILATION_FLAGS
"--iree-input-type=tosa"
"--iree-flow-demote-f32-to-f16"
- "--iree-vulkan-target-triple=valhall-unknown-android11"
+ "--iree-vulkan-target-triple=valhall-unknown-android31"
"--iree-flow-enable-fuse-padding-into-consumer-ops"
BENCHMARK_TOOL
iree-benchmark-module
@@ -208,7 +208,7 @@
COMPILATION_FLAGS
"--iree-input-type=tosa"
"--iree-flow-demote-f32-to-f16"
- "--iree-vulkan-target-triple=valhall-unknown-android11"
+ "--iree-vulkan-target-triple=valhall-unknown-android31"
"--iree-flow-enable-fuse-padding-into-consumer-ops"
"--iree-hal-benchmark-dispatch-repeat-count=32"
BENCHMARK_TOOL
diff --git a/build_tools/benchmarks/comparisons/setup_mobile.sh b/build_tools/benchmarks/comparisons/setup_mobile.sh
index 45f8422..9592879 100644
--- a/build_tools/benchmarks/comparisons/setup_mobile.sh
+++ b/build_tools/benchmarks/comparisons/setup_mobile.sh
@@ -103,7 +103,7 @@
"${IREE_COMPILE_PATH}" \
--iree-input-type=tosa \
--iree-hal-target-backends=vulkan-spirv \
- --iree-vulkan-target-triple=valhall-unknown-android11 \
+ --iree-vulkan-target-triple=valhall-unknown-android31 \
--iree-llvm-debug-symbols=false \
--iree-vm-bytecode-module-strip-source-map=true \
--iree-vm-emit-polyglot-zip=false \
@@ -114,7 +114,7 @@
"${IREE_COMPILE_PATH}" \
--iree-input-type=tosa \
--iree-hal-target-backends=vulkan-spirv \
- --iree-vulkan-target-triple=adreno-unknown-android11 \
+ --iree-vulkan-target-triple=adreno-unknown-android31 \
--iree-llvm-debug-symbols=false \
--iree-vm-bytecode-module-strip-source-map=true \
--iree-vm-emit-polyglot-zip=false \
diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanBase.td b/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanBase.td
index 24657e8..a16adc0 100644
--- a/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanBase.td
+++ b/compiler/src/iree/compiler/Dialect/Vulkan/IR/VulkanBase.td
@@ -78,9 +78,11 @@
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_I32EnumAttr<"Version", "valid Vulkan version", [
- VK_V_1_0, VK_V_1_1, VK_V_1_2]>;
+ 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>;
@@ -141,16 +143,19 @@
]>;
def VK_TTOS_Unknown : I32EnumAttrCase<"Unknown", 0, "unknown">;
-def VK_TTOS_Android11 : I32EnumAttrCase<"Android11", 2, "android11">;
-def VK_TTOS_iOS : I32EnumAttrCase<"iOS", 3, "iOS">;
-def VK_TTOS_Linux : I32EnumAttrCase<"Linux", 4, "linux">;
-def VK_TTOS_macOS : I32EnumAttrCase<"macOS", 5, "macos">;
-def VK_TTOS_Windows : I32EnumAttrCase<"Windows", 6, "windows">;
+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_I32EnumAttr<
"TargetTripleOS", "recognized target operating system", [
- VK_TTOS_Unknown, VK_TTOS_Android11, VK_TTOS_iOS, VK_TTOS_Linux,
- VK_TTOS_macOS, VK_TTOS_Windows,
+ VK_TTOS_Unknown, VK_TTOS_Linux, VK_TTOS_iOS, VK_TTOS_macOS,
+ VK_TTOS_Windows, VK_TTOS_Android30, VK_TTOS_Android31,
]>;
//===----------------------------------------------------------------------===//
diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetEnvironment.cpp b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetEnvironment.cpp
index 839b6a6..15a30c6 100644
--- a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetEnvironment.cpp
+++ b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetEnvironment.cpp
@@ -21,25 +21,24 @@
/// Gets the corresponding SPIR-V version for the ggiven Vulkan target
/// environment.
spirv::Version convertVersion(Vulkan::TargetEnvAttr vkTargetEnv) {
- // Vulkan 1.2 supports up to SPIR-V 1.5 by default.
- if (vkTargetEnv.getVersion() == Version::V_1_2) return spirv::Version::V_1_5;
-
// Special extension to enable SPIR-V 1.4.
- if (llvm::is_contained(vkTargetEnv.getExtensions(),
- Extension::VK_KHR_spirv_1_4))
- return spirv::Version::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 spirv::Version::V_1_0;
+ 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 spirv::Version::V_1_3;
- default:
- break;
+ 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;
}
- assert(false && "unhandled Vulkan version!");
return spirv::Version::V_1_0;
}
diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetTriple.cpp b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetTriple.cpp
index 16510c1..d818eea 100644
--- a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetTriple.cpp
+++ b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetTriple.cpp
@@ -80,8 +80,9 @@
/// Returns the Vulkan version for the given target `triple`.
Vulkan::Version getVersion(const TargetTriple &triple) {
- // Android 11 stays at Vulkan 1.1.
- if (triple.getOS() == TargetTripleOS::Android11) {
+ // 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;
}
@@ -91,7 +92,7 @@
return Version::V_1_1;
}
- return Version::V_1_2;
+ return Version::V_1_3;
}
/// Writes the Vulkan extensions supported by the given `triple` into
@@ -130,7 +131,8 @@
return extensions.append(list.begin(), list.end());
}
case TargetTripleArch::QC_Adreno: {
- // Example: https://vulkan.gpuinfo.org/displayreport.php?id=10983
+ // Example: https://vulkan.gpuinfo.org/displayreport.php?id=10983 (11)
+ // Example: https://vulkan.gpuinfo.org/displayreport.php?id=16312 (12)
const std::array<Extension, 5> list = {
Extension::VK_KHR_16bit_storage,
Extension::VK_KHR_shader_float16_int8,
@@ -138,7 +140,11 @@
Extension::VK_KHR_storage_buffer_storage_class,
Extension::VK_KHR_variable_pointers,
};
- return extensions.append(list.begin(), list.end());
+ extensions.append(list.begin(), list.end());
+ if (triple.getOS() == TargetTripleOS::Android31) {
+ extensions.push_back(Extension::VK_KHR_8bit_storage);
+ }
+ return;
}
default:
break;
@@ -249,7 +255,8 @@
variablePointers = variablePointersStorageBuffer = true;
break;
case TargetTripleArch::ARM_Valhall:
- // Example: https://vulkan.gpuinfo.org/displayreport.php?id=10312
+ // 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};
@@ -259,6 +266,11 @@
SubgroupFeature::Arithmetic | SubgroupFeature::Ballot |
SubgroupFeature::Clustered | SubgroupFeature::Quad;
+ if (triple.getOS() == TargetTripleOS::Android31) {
+ subgroupFeatures = subgroupFeatures | SubgroupFeature::Shuffle |
+ SubgroupFeature::ShuffleRelative;
+ }
+
shaderFloat16 = shaderInt8 = shaderInt16 = true;
storageBuffer16BitAccess = storagePushConstant16 = true;
@@ -324,7 +336,8 @@
/*bType=*/f16t, /*cType=*/f32t, /*resultType=*/f32t, scope));
} break;
case TargetTripleArch::QC_Adreno:
- // Example: https://vulkan.gpuinfo.org/displayreport.php?id=10983
+ // 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};
@@ -339,6 +352,10 @@
shaderFloat16 = shaderInt8 = shaderInt16 = true;
storageBuffer16BitAccess = true;
+ if (triple.getOS() == TargetTripleOS::Android31) {
+ storageBuffer8BitAccess = true;
+ }
+
variablePointers = variablePointersStorageBuffer = true;
break;
}
diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetTriple.h b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetTriple.h
index 4012f56..4171a86 100644
--- a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetTriple.h
+++ b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetTriple.h
@@ -31,8 +31,8 @@
/// For example:
/// ampere-rtx3080-windows
/// rdna1-5700xt-linux
-/// adreno-a650-android11
-/// valhall-unknown-android11
+/// adreno-a650-android30
+/// valhall-unknown-android30
/// cpu-swiftshader-unknown
///
/// Vendor and architecture are combined together because:
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
index 6e3fd39..4a1293a 100644
--- 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
@@ -1,6 +1,6 @@
// RUN: iree-opt --pass-pipeline='iree-hal-transformation-pipeline{serialize-executables=false}' --iree-hal-target-backends=vulkan-spirv %s | FileCheck %s --check-prefix=DEFAULT
-// RUN: iree-opt --pass-pipeline='iree-hal-transformation-pipeline{serialize-executables=false}' --iree-hal-target-backends=vulkan-spirv --iree-vulkan-target-triple=adreno-a650-android11 %s | FileCheck %s --check-prefix=ADRENO
-// RUN: iree-opt --pass-pipeline='iree-hal-transformation-pipeline{serialize-executables=false}' --iree-hal-target-backends=vulkan-spirv --iree-vulkan-target-triple=valhall-unknown-android11 %s | FileCheck %s --check-prefix=MALI
+// RUN: iree-opt --pass-pipeline='iree-hal-transformation-pipeline{serialize-executables=false}' --iree-hal-target-backends=vulkan-spirv --iree-vulkan-target-triple=adreno-a650-android30 %s | FileCheck %s --check-prefix=ADRENO
+// RUN: iree-opt --pass-pipeline='iree-hal-transformation-pipeline{serialize-executables=false}' --iree-hal-target-backends=vulkan-spirv --iree-vulkan-target-triple=valhall-unknown-android31 %s | FileCheck %s --check-prefix=MALI
// RUN: iree-opt --pass-pipeline='iree-hal-transformation-pipeline{serialize-executables=false}' --iree-hal-target-backends=vulkan-spirv --iree-vulkan-target-triple=turing-t4-linux %s | FileCheck %s --check-prefix=TURINGT4
// RUN: iree-opt --pass-pipeline='iree-hal-transformation-pipeline{serialize-executables=false}' --iree-hal-target-backends=vulkan-spirv --iree-vulkan-target-triple=rdna1-5700xt-windows %s | FileCheck %s --check-prefix=AMD5700XT
// RUN: iree-opt --pass-pipeline='iree-hal-transformation-pipeline{serialize-executables=false}' --iree-hal-target-backends=vulkan-spirv --iree-vulkan-target-triple=m1-moltenvk-macos %s | FileCheck %s --check-prefix=M1
@@ -11,9 +11,9 @@
// DEFAULT: #spv.target_env<#spv.vce<v1.3, [Shader, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative], [SPV_KHR_storage_buffer_storage_class]>, SwiftShader:CPU, #spv.resource_limits<max_compute_workgroup_size = [128, 128, 64], subgroup_size = 4, cooperative_matrix_properties_nv = []>>
// ADRENO: #spv.target_env<#spv.vce<v1.4, [Shader, Float16, Int16, Int8, StorageBuffer16BitAccess, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative, GroupNonUniformQuad, VariablePointers, VariablePointersStorageBuffer], [SPV_KHR_16bit_storage, SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers]>, Qualcomm:IntegratedGPU, #spv.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_nv = []>>
-// MALI: #spv.target_env<#spv.vce<v1.4, [Shader, Float16, Int16, Int8, StorageBuffer16BitAccess, StorageUniform16, StoragePushConstant16, StorageBuffer8BitAccess, UniformAndStorageBuffer8BitAccess, StoragePushConstant8, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformClustered, GroupNonUniformQuad, VariablePointers, VariablePointersStorageBuffer], [SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers]>, ARM:IntegratedGPU, #spv.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_nv = []>>
-// TURINGT4: #spv.target_env<#spv.vce<v1.5, [Shader, Float64, Float16, Int64, Int16, Int8, StorageBuffer16BitAccess, StorageUniform16, StoragePushConstant16, StorageBuffer8BitAccess, UniformAndStorageBuffer8BitAccess, StoragePushConstant8, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative, GroupNonUniformClustered, GroupNonUniformQuad, VariablePointers, VariablePointersStorageBuffer, CooperativeMatrixNV], [SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers, SPV_NV_cooperative_matrix]>, NVIDIA:DiscreteGPU, #spv.resource_limits<max_compute_shared_memory_size = 49152, max_compute_workgroup_invocations = 1024, max_compute_workgroup_size = [1024, 1024, 64], cooperative_matrix_properties_nv = [#spv.coop_matrix_props<m_size = 8, n_size = 8, k_size = 32, a_type = i8, b_type = i8, c_type = i32, result_type = i32, scope = <Subgroup>>, #spv.coop_matrix_props<m_size = 16, n_size = 16, k_size = 16, a_type = f16, b_type = f16, c_type = f16, result_type = f16, scope = <Subgroup>>, #spv.coop_matrix_props<m_size = 16, n_size = 16, k_size = 16, a_type = f16, b_type = f16, c_type = f32, result_type = f32, scope = <Subgroup>>]>>
-// AMD5700XT: #spv.target_env<#spv.vce<v1.5, [Shader, Float64, Float16, Int64, Int16, Int8, StorageBuffer16BitAccess, StorageUniform16, StoragePushConstant16, StorageBuffer8BitAccess, UniformAndStorageBuffer8BitAccess, StoragePushConstant8, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative, GroupNonUniformClustered, GroupNonUniformQuad, VariablePointers, VariablePointersStorageBuffer], [SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers]>, AMD:DiscreteGPU, #spv.resource_limits<max_compute_shared_memory_size = 65536, max_compute_workgroup_invocations = 1024, max_compute_workgroup_size = [1024, 1024, 1024], subgroup_size = 64, cooperative_matrix_properties_nv = []>>
+// MALI: #spv.target_env<#spv.vce<v1.4, [Shader, Float16, Int16, Int8, StorageBuffer16BitAccess, StorageUniform16, StoragePushConstant16, StorageBuffer8BitAccess, UniformAndStorageBuffer8BitAccess, StoragePushConstant8, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative, GroupNonUniformClustered, GroupNonUniformQuad, VariablePointers, VariablePointersStorageBuffer], [SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers]>, ARM:IntegratedGPU, #spv.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_nv = []>>
+// TURINGT4: #spv.target_env<#spv.vce<v1.6, [Shader, Float64, Float16, Int64, Int16, Int8, StorageBuffer16BitAccess, StorageUniform16, StoragePushConstant16, StorageBuffer8BitAccess, UniformAndStorageBuffer8BitAccess, StoragePushConstant8, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative, GroupNonUniformClustered, GroupNonUniformQuad, VariablePointers, VariablePointersStorageBuffer, CooperativeMatrixNV], [SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers, SPV_NV_cooperative_matrix]>, NVIDIA:DiscreteGPU, #spv.resource_limits<max_compute_shared_memory_size = 49152, max_compute_workgroup_invocations = 1024, max_compute_workgroup_size = [1024, 1024, 64], cooperative_matrix_properties_nv = [#spv.coop_matrix_props<m_size = 8, n_size = 8, k_size = 32, a_type = i8, b_type = i8, c_type = i32, result_type = i32, scope = <Subgroup>>, #spv.coop_matrix_props<m_size = 16, n_size = 16, k_size = 16, a_type = f16, b_type = f16, c_type = f16, result_type = f16, scope = <Subgroup>>, #spv.coop_matrix_props<m_size = 16, n_size = 16, k_size = 16, a_type = f16, b_type = f16, c_type = f32, result_type = f32, scope = <Subgroup>>]>>
+// AMD5700XT: #spv.target_env<#spv.vce<v1.6, [Shader, Float64, Float16, Int64, Int16, Int8, StorageBuffer16BitAccess, StorageUniform16, StoragePushConstant16, StorageBuffer8BitAccess, UniformAndStorageBuffer8BitAccess, StoragePushConstant8, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative, GroupNonUniformClustered, GroupNonUniformQuad, VariablePointers, VariablePointersStorageBuffer], [SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers]>, AMD:DiscreteGPU, #spv.resource_limits<max_compute_shared_memory_size = 65536, max_compute_workgroup_invocations = 1024, max_compute_workgroup_size = [1024, 1024, 1024], subgroup_size = 64, cooperative_matrix_properties_nv = []>>
// M1: #spv.target_env<#spv.vce<v1.3, [Shader, Float16, Int64, Int16, Int8, StorageBuffer16BitAccess, StorageUniform16, StoragePushConstant16, StorageBuffer8BitAccess, UniformAndStorageBuffer8BitAccess, StoragePushConstant8, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative, GroupNonUniformQuad, VariablePointers, VariablePointersStorageBuffer], [SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers]>, Apple:IntegratedGPU, #spv.resource_limits<max_compute_shared_memory_size = 32768, max_compute_workgroup_invocations = 1024, max_compute_workgroup_size = [1024, 1024, 1024], cooperative_matrix_properties_nv = []>>
diff --git a/docs/website/docs/deployment-configurations/gpu-vulkan.md b/docs/website/docs/deployment-configurations/gpu-vulkan.md
index 0ffc239..799cdf9 100644
--- a/docs/website/docs/deployment-configurations/gpu-vulkan.md
+++ b/docs/website/docs/deployment-configurations/gpu-vulkan.md
@@ -161,8 +161,8 @@
GPU Vendor | Target Triple
:--------: | :-----------:
-ARM Mali GPU | `valhall-g78-android11`
-Qualcomm Adreno GPU | `adreno-unknown-android11`
+ARM Mali GPU | `valhall-g78-android30`
+Qualcomm Adreno GPU | `adreno-unknown-android30`
AMD GPU | e.g., `rdna1-5700xt-linux`
NVIDIA GPU | e..g, `ampere-rtx3080-windows`
SwiftShader CPU | `cpu-swiftshader-unknown`
diff --git a/tests/e2e/matmul/BUILD b/tests/e2e/matmul/BUILD
index c19fe3a..743a048 100644
--- a/tests/e2e/matmul/BUILD
+++ b/tests/e2e/matmul/BUILD
@@ -235,7 +235,7 @@
],
trace_runner = "//tools:iree-e2e-matmul-test",
) for vulkan_target in [
- "valhall-unknown-android11",
+ "valhall-unknown-android31",
"ampere-unknown-linux",
]]
diff --git a/tests/e2e/matmul/CMakeLists.txt b/tests/e2e/matmul/CMakeLists.txt
index 5a68326..52b6a39 100644
--- a/tests/e2e/matmul/CMakeLists.txt
+++ b/tests/e2e/matmul/CMakeLists.txt
@@ -320,7 +320,7 @@
iree_generated_trace_runner_test(
NAME
- e2e_matmul_direct_f32_gpu_large_valhall-unknown-android11
+ e2e_matmul_direct_f32_gpu_large_valhall-unknown-android31
GENERATOR
"generate_e2e_matmul_tests.py"
GENERATOR_ARGS
@@ -334,7 +334,7 @@
DRIVERS
"vulkan"
COMPILER_FLAGS
- "--iree-vulkan-target-triple=valhall-unknown-android11"
+ "--iree-vulkan-target-triple=valhall-unknown-android31"
LABELS
"requires-gpu-nvidia"
)
diff --git a/tests/e2e/vulkan_specific/BUILD b/tests/e2e/vulkan_specific/BUILD
index b8421f0..ba2715b 100644
--- a/tests/e2e/vulkan_specific/BUILD
+++ b/tests/e2e/vulkan_specific/BUILD
@@ -21,7 +21,7 @@
],
compiler_flags = [
"--iree-input-type=mhlo",
- "--iree-vulkan-target-triple=valhall-unknown-android11",
+ "--iree-vulkan-target-triple=valhall-unknown-android31",
],
driver = "vulkan",
tags = [
@@ -39,7 +39,7 @@
],
compiler_flags = [
"--iree-input-type=mhlo",
- "--iree-vulkan-target-triple=valhall-unknown-android11",
+ "--iree-vulkan-target-triple=valhall-unknown-android31",
],
driver = "vulkan",
target_backend = "vulkan-spirv",
diff --git a/tests/e2e/vulkan_specific/CMakeLists.txt b/tests/e2e/vulkan_specific/CMakeLists.txt
index 240fcc2..71367ad 100644
--- a/tests/e2e/vulkan_specific/CMakeLists.txt
+++ b/tests/e2e/vulkan_specific/CMakeLists.txt
@@ -22,7 +22,7 @@
"vulkan"
COMPILER_FLAGS
"--iree-input-type=mhlo"
- "--iree-vulkan-target-triple=valhall-unknown-android11"
+ "--iree-vulkan-target-triple=valhall-unknown-android31"
LABELS
"manual"
"notap"
@@ -40,7 +40,7 @@
"vulkan"
COMPILER_FLAGS
"--iree-input-type=mhlo"
- "--iree-vulkan-target-triple=valhall-unknown-android11"
+ "--iree-vulkan-target-triple=valhall-unknown-android31"
)
### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
