[Mlir-commits] [mlir] 80bb947 - [mlir] NFC: Prepare GPUToSPIRV tests for supporting 64bit index
Lei Zhang
llvmlistbot at llvm.org
Sun Feb 26 14:13:32 PST 2023
Author: Lei Zhang
Date: 2023-02-26T22:13:15Z
New Revision: 80bb9477a97376a1c0a25a8c249db5518a9feb94
URL: https://github.com/llvm/llvm-project/commit/80bb9477a97376a1c0a25a8c249db5518a9feb94
DIFF: https://github.com/llvm/llvm-project/commit/80bb9477a97376a1c0a25a8c249db5518a9feb94.diff
LOG: [mlir] NFC: Prepare GPUToSPIRV tests for supporting 64bit index
This commit just adds options to control index type bitwidth in
GPUToSPIRV conversion, and updates tests to prepare for 64bit
index conversion.
Reviewed By: kuhar
Differential Revision: https://reviews.llvm.org/D144826
Added:
Modified:
mlir/include/mlir/Conversion/Passes.td
mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
mlir/test/Conversion/GPUToSPIRV/builtins.mlir
Removed:
################################################################################
diff --git a/mlir/include/mlir/Conversion/Passes.td b/mlir/include/mlir/Conversion/Passes.td
index 1417232f5d7f3..ed7282a824abe 100644
--- a/mlir/include/mlir/Conversion/Passes.td
+++ b/mlir/include/mlir/Conversion/Passes.td
@@ -476,6 +476,11 @@ def ConvertGPUToSPIRV : Pass<"convert-gpu-to-spirv", "ModuleOp"> {
}];
let constructor = "mlir::createConvertGPUToSPIRVPass()";
let dependentDialects = ["spirv::SPIRVDialect"];
+ let options = [
+ Option<"use64bitIndex", "use-64bit-index",
+ "bool", /*default=*/"false",
+ "Use 64-bit integers to convert index types">
+ ];
}
//===----------------------------------------------------------------------===//
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
index f1c4e32da827f..f37c70a771f59 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
@@ -85,7 +85,9 @@ void GPUToSPIRVPass::runOnOperation() {
std::unique_ptr<ConversionTarget> target =
SPIRVConversionTarget::get(targetAttr);
- SPIRVTypeConverter typeConverter(targetAttr);
+ SPIRVConversionOptions options;
+ options.use64bitIndex = this->use64bitIndex;
+ SPIRVTypeConverter typeConverter(targetAttr, options);
typeConverter.addConversion([&](gpu::MMAMatrixType type) -> Type {
return convertMMAToSPIRVType(type);
});
diff --git a/mlir/test/Conversion/GPUToSPIRV/builtins.mlir b/mlir/test/Conversion/GPUToSPIRV/builtins.mlir
index 76496875827a9..df2efbed50c90 100644
--- a/mlir/test/Conversion/GPUToSPIRV/builtins.mlir
+++ b/mlir/test/Conversion/GPUToSPIRV/builtins.mlir
@@ -1,6 +1,9 @@
-// RUN: mlir-opt -split-input-file -convert-gpu-to-spirv %s -o - | FileCheck %s
+// RUN: mlir-opt -split-input-file -convert-gpu-to-spirv="use-64bit-index=false" %s -o - | FileCheck %s --check-prefix=INDEX32
-module attributes {gpu.container_module} {
+module attributes {
+ gpu.container_module,
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
func.func @builtin() {
%c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_workgroup_id_x
@@ -8,14 +11,14 @@ module attributes {gpu.container_module} {
return
}
- // CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450
- // CHECK: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
+ // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
+ // INDEX32: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") : !spirv.ptr<vector<3xi32>, Input>
gpu.module @kernels {
gpu.func @builtin_workgroup_id_x() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
- // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPID]]
- // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
- // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
+ // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPID]]
+ // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
+ // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
%0 = gpu.block_id x
gpu.return
}
@@ -24,7 +27,10 @@ module attributes {gpu.container_module} {
// -----
-module attributes {gpu.container_module} {
+module attributes {
+ gpu.container_module,
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
func.func @builtin() {
%c0 = arith.constant 1 : index
%c256 = arith.constant 256 : i32
@@ -34,14 +40,14 @@ module attributes {gpu.container_module} {
return
}
- // CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450
- // CHECK: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
+ // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
+ // INDEX32: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") : !spirv.ptr<vector<3xi32>, Input>
gpu.module @kernels {
gpu.func @builtin_workgroup_id_y() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
- // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPID]]
- // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
- // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
+ // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPID]]
+ // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
+ // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
%0 = gpu.block_id y
gpu.return
}
@@ -50,7 +56,10 @@ module attributes {gpu.container_module} {
// -----
-module attributes {gpu.container_module} {
+module attributes {
+ gpu.container_module,
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
func.func @builtin() {
%c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_workgroup_id_z
@@ -58,14 +67,14 @@ module attributes {gpu.container_module} {
return
}
- // CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450
- // CHECK: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
+ // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
+ // INDEX32: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") : !spirv.ptr<vector<3xi32>, Input>
gpu.module @kernels {
gpu.func @builtin_workgroup_id_z() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
- // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPID]]
- // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
- // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
+ // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPID]]
+ // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
+ // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
%0 = gpu.block_id z
gpu.return
}
@@ -74,7 +83,10 @@ module attributes {gpu.container_module} {
// -----
-module attributes {gpu.container_module} {
+module attributes {
+ gpu.container_module,
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
func.func @builtin() {
%c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_workgroup_size_x
@@ -82,7 +94,7 @@ module attributes {gpu.container_module} {
return
}
- // CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450
+ // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
gpu.module @kernels {
gpu.func @builtin_workgroup_size_x() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 1, 1]>} {
@@ -90,7 +102,7 @@ module attributes {gpu.container_module} {
// Note that this ignores the workgroup size specification in gpu.launch.
// We may want to define gpu.workgroup_size and convert it to the entry
// point ABI we want here.
- // CHECK: spirv.Constant 32 : i32
+ // INDEX32: spirv.Constant 32 : i32
%0 = gpu.block_dim x
gpu.return
}
@@ -99,7 +111,10 @@ module attributes {gpu.container_module} {
// -----
-module attributes {gpu.container_module} {
+module attributes {
+ gpu.container_module,
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
func.func @builtin() {
%c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_workgroup_size_y
@@ -107,12 +122,12 @@ module attributes {gpu.container_module} {
return
}
- // CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450
+ // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
gpu.module @kernels {
gpu.func @builtin_workgroup_size_y() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
// The constant value is obtained from the spirv.entry_point_abi.
- // CHECK: spirv.Constant 4 : i32
+ // INDEX32: spirv.Constant 4 : i32
%0 = gpu.block_dim y
gpu.return
}
@@ -121,7 +136,10 @@ module attributes {gpu.container_module} {
// -----
-module attributes {gpu.container_module} {
+module attributes {
+ gpu.container_module,
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
func.func @builtin() {
%c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_workgroup_size_z
@@ -129,12 +147,12 @@ module attributes {gpu.container_module} {
return
}
- // CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450
+ // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
gpu.module @kernels {
gpu.func @builtin_workgroup_size_z() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
// The constant value is obtained from the spirv.entry_point_abi.
- // CHECK: spirv.Constant 1 : i32
+ // INDEX32: spirv.Constant 1 : i32
%0 = gpu.block_dim z
gpu.return
}
@@ -143,7 +161,10 @@ module attributes {gpu.container_module} {
// -----
-module attributes {gpu.container_module} {
+module attributes {
+ gpu.container_module,
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
func.func @builtin() {
%c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_local_id_x
@@ -151,14 +172,14 @@ module attributes {gpu.container_module} {
return
}
- // CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450
- // CHECK: spirv.GlobalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId")
+ // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
+ // INDEX32: spirv.GlobalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId") : !spirv.ptr<vector<3xi32>, Input>
gpu.module @kernels {
gpu.func @builtin_local_id_x() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
- // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[LOCALINVOCATIONID]]
- // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
- // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
+ // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[LOCALINVOCATIONID]]
+ // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
+ // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
%0 = gpu.thread_id x
gpu.return
}
@@ -167,7 +188,10 @@ module attributes {gpu.container_module} {
// -----
-module attributes {gpu.container_module} {
+module attributes {
+ gpu.container_module,
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
func.func @builtin() {
%c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_num_workgroups_x
@@ -175,14 +199,14 @@ module attributes {gpu.container_module} {
return
}
- // CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450
- // CHECK: spirv.GlobalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups")
+ // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
+ // INDEX32: spirv.GlobalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups") : !spirv.ptr<vector<3xi32>, Input>
gpu.module @kernels {
gpu.func @builtin_num_workgroups_x() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
- // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[NUMWORKGROUPS]]
- // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
- // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
+ // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[NUMWORKGROUPS]]
+ // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
+ // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
%0 = gpu.grid_dim x
gpu.return
}
@@ -191,14 +215,17 @@ module attributes {gpu.container_module} {
// -----
-module attributes {gpu.container_module} {
- // CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450
- // CHECK: spirv.GlobalVariable [[SUBGROUPID:@.*]] built_in("SubgroupId")
+module attributes {
+ gpu.container_module,
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
+ // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
+ // INDEX32: spirv.GlobalVariable [[SUBGROUPID:@.*]] built_in("SubgroupId") : !spirv.ptr<i32, Input>
gpu.module @kernels {
gpu.func @builtin_subgroup_id() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
- // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[SUBGROUPID]]
- // CHECK-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]]
+ // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[SUBGROUPID]]
+ // INDEX32-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]]
%0 = gpu.subgroup_id : index
gpu.return
}
@@ -207,14 +234,17 @@ module attributes {gpu.container_module} {
// -----
-module attributes {gpu.container_module} {
- // CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450
- // CHECK: spirv.GlobalVariable [[NUMSUBGROUPS:@.*]] built_in("NumSubgroups")
+module attributes {
+ gpu.container_module,
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
+ // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
+ // INDEX32: spirv.GlobalVariable [[NUMSUBGROUPS:@.*]] built_in("NumSubgroups") : !spirv.ptr<i32, Input>
gpu.module @kernels {
gpu.func @builtin_num_subgroups() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
- // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[NUMSUBGROUPS]]
- // CHECK-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]]
+ // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[NUMSUBGROUPS]]
+ // INDEX32-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]]
%0 = gpu.num_subgroups : index
gpu.return
}
@@ -223,7 +253,10 @@ module attributes {gpu.container_module} {
// -----
-module attributes {gpu.container_module} {
+module attributes {
+ gpu.container_module,
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
func.func @builtin() {
%c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_workgroup_size_x
@@ -231,14 +264,14 @@ module attributes {gpu.container_module} {
return
}
- // CHECK-LABEL: spirv.module @{{.*}}
- // CHECK: spirv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize")
+ // INDEX32-LABEL: spirv.module @{{.*}}
+ // INDEX32: spirv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize") : !spirv.ptr<vector<3xi32>, Input>
gpu.module @kernels {
gpu.func @builtin_workgroup_size_x() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<>} {
- // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPSIZE]]
- // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
- // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
+ // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPSIZE]]
+ // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
+ // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
%0 = gpu.block_dim x
gpu.return
}
@@ -247,7 +280,10 @@ module attributes {gpu.container_module} {
// -----
-module attributes {gpu.container_module} {
+module attributes {
+ gpu.container_module,
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
func.func @builtin() {
%c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_workgroup_size_y
@@ -255,14 +291,14 @@ module attributes {gpu.container_module} {
return
}
- // CHECK-LABEL: spirv.module @{{.*}}
- // CHECK: spirv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize")
+ // INDEX32-LABEL: spirv.module @{{.*}}
+ // INDEX32: spirv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize") : !spirv.ptr<vector<3xi32>, Input>
gpu.module @kernels {
gpu.func @builtin_workgroup_size_y() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<>} {
- // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPSIZE]]
- // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
- // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
+ // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPSIZE]]
+ // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
+ // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
%0 = gpu.block_dim y
gpu.return
}
@@ -271,7 +307,10 @@ module attributes {gpu.container_module} {
// -----
-module attributes {gpu.container_module} {
+module attributes {
+ gpu.container_module,
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
func.func @builtin() {
%c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_workgroup_size_z
@@ -279,14 +318,14 @@ module attributes {gpu.container_module} {
return
}
- // CHECK-LABEL: spirv.module @{{.*}}
- // CHECK: spirv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize")
+ // INDEX32-LABEL: spirv.module @{{.*}}
+ // INDEX32: spirv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize") : !spirv.ptr<vector<3xi32>, Input>
gpu.module @kernels {
gpu.func @builtin_workgroup_size_z() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<>} {
- // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPSIZE]]
- // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
- // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
+ // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPSIZE]]
+ // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
+ // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
%0 = gpu.block_dim z
gpu.return
}
@@ -295,7 +334,10 @@ module attributes {gpu.container_module} {
// -----
-module attributes {gpu.container_module} {
+module attributes {
+ gpu.container_module,
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
func.func @builtin() {
%c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_global_id_x
@@ -303,14 +345,14 @@ module attributes {gpu.container_module} {
return
}
- // CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450
- // CHECK: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId")
+ // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
+ // INDEX32: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId") : !spirv.ptr<vector<3xi32>, Input>
gpu.module @kernels {
gpu.func @builtin_global_id_x() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
- // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[GLOBALINVOCATIONID]]
- // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
- // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
+ // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[GLOBALINVOCATIONID]]
+ // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
+ // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
%0 = gpu.global_id x
gpu.return
}
@@ -319,7 +361,10 @@ module attributes {gpu.container_module} {
// -----
-module attributes {gpu.container_module} {
+module attributes {
+ gpu.container_module,
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
func.func @builtin() {
%c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_global_id_y
@@ -327,14 +372,14 @@ module attributes {gpu.container_module} {
return
}
- // CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450
- // CHECK: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId")
+ // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
+ // INDEX32: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId") : !spirv.ptr<vector<3xi32>, Input>
gpu.module @kernels {
gpu.func @builtin_global_id_y() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
- // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[GLOBALINVOCATIONID]]
- // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
- // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
+ // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[GLOBALINVOCATIONID]]
+ // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
+ // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
%0 = gpu.global_id y
gpu.return
}
@@ -343,7 +388,10 @@ module attributes {gpu.container_module} {
// -----
-module attributes {gpu.container_module} {
+module attributes {
+ gpu.container_module,
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
func.func @builtin() {
%c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_global_id_z
@@ -351,14 +399,14 @@ module attributes {gpu.container_module} {
return
}
- // CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450
- // CHECK: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId")
+ // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
+ // INDEX32: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId") : !spirv.ptr<vector<3xi32>, Input>
gpu.module @kernels {
gpu.func @builtin_global_id_z() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
- // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[GLOBALINVOCATIONID]]
- // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
- // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
+ // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[GLOBALINVOCATIONID]]
+ // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
+ // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
%0 = gpu.global_id z
gpu.return
}
@@ -368,14 +416,17 @@ module attributes {gpu.container_module} {
// -----
-module attributes {gpu.container_module} {
- // CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450
- // CHECK: spirv.GlobalVariable [[SUBGROUPSIZE:@.*]] built_in("SubgroupSize")
+module attributes {
+ gpu.container_module,
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
+ // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
+ // INDEX32: spirv.GlobalVariable [[SUBGROUPSIZE:@.*]] built_in("SubgroupSize") : !spirv.ptr<i32, Input>
gpu.module @kernels {
gpu.func @builtin_subgroup_size() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
- // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[SUBGROUPSIZE]]
- // CHECK-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]]
+ // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[SUBGROUPSIZE]]
+ // INDEX32-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]]
%0 = gpu.subgroup_size : index
gpu.return
}
More information about the Mlir-commits
mailing list