This is an archive of the discontinued LLVM Phabricator instance.

Differential D144826

[mlir] NFC: Prepare GPUToSPIRV tests for supporting 64bit index
ClosedPublic

Authored by antiagainst on Feb 26 2023, 7:31 AM.

Details

Reviewers
kuhar
ThomasRaoux
herhut
Commits
rG80bb9477a973: [mlir] NFC: Prepare GPUToSPIRV tests for supporting 64bit index
Summary

This commit just adds options to control index type bitwidth in
GPUToSPIRV conversion, and updates tests to prepare for 64bit
index conversion.

Diff Detail

Event Timeline

antiagainst created this revision.Feb 26 2023, 7:31 AM
Herald added a project: Restricted Project. · View Herald TranscriptFeb 26 2023, 7:31 AM
Herald added subscribers: Moerafaat, zero9178, bzcheeseman and 23 others. · View Herald Transcript
antiagainst requested review of this revision.Feb 26 2023, 7:31 AM
Herald added a reviewer: herhut. · View Herald TranscriptFeb 26 2023, 7:31 AM
Herald added a project: Restricted Project. · View Herald Transcript
Herald added subscribers: stephenneuendorffer, nicolasvasilache. · View Herald Transcript
antiagainst added a child revision: D144819: [mlir][spirv] Respect client API requirements for 64-bit index.Feb 26 2023, 7:31 AM
Harbormaster completed remote builds in B216072: Diff 500581.Feb 26 2023, 7:50 AM
kuhar accepted this revision.Feb 26 2023, 12:22 PM
This revision is now accepted and ready to land.Feb 26 2023, 12:22 PM
Closed by commit rG80bb9477a973: [mlir] NFC: Prepare GPUToSPIRV tests for supporting 64bit index (authored by antiagainst). · Explain WhyFeb 26 2023, 2:13 PM
This revision was automatically updated to reflect the committed changes.
antiagainst added a commit: rG80bb9477a973: [mlir] NFC: Prepare GPUToSPIRV tests for supporting 64bit index.

Revision Contents

PathSize
mlir/
include/
mlir/
Conversion/
5 lines
lib/
Conversion/
GPUToSPIRV/
4 lines
test/
Conversion/
GPUToSPIRV/
233 lines

Diff 500581

mlir/include/mlir/Conversion/Passes.td

Show First 20 Lines • Show All 470 Lines • ▼ Show 20 Lines let description = [{
resources. By default, parameters to a `gpu.func` op will be converted to resources. By default, parameters to a `gpu.func` op will be converted to
global variables. These global variables will be assigned sequential binding global variables. These global variables will be assigned sequential binding
numbers following their order in the original `gpu.func` op, starting from numbers following their order in the original `gpu.func` op, starting from
0, in set 0. One can attach `spirv.interface_var_abi` to those parameters 0, in set 0. One can attach `spirv.interface_var_abi` to those parameters
to control the set and binding if wanted. to control the set and binding if wanted.
}]; }];
let constructor = "mlir::createConvertGPUToSPIRVPass()"; let constructor = "mlir::createConvertGPUToSPIRVPass()";
let dependentDialects = ["spirv::SPIRVDialect"]; let dependentDialects = ["spirv::SPIRVDialect"];
let options = [
Option<"use64bitIndex", "use-64bit-index",
"bool", /*default=*/"false",
"Use 64-bit integers to convert index types">
];
} }
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// GPUToVulkan // GPUToVulkan
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
def ConvertGpuLaunchFuncToVulkanLaunchFunc def ConvertGpuLaunchFuncToVulkanLaunchFunc
: Pass<"convert-gpu-launch-to-vulkan-launch", "ModuleOp"> { : Pass<"convert-gpu-launch-to-vulkan-launch", "ModuleOp"> {
▲ Show 20 Lines • Show All 605 Lines • Show Last 20 Lines

mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp

Show First 20 Lines • Show All 79 Lines • ▼ Show 20 Lines if (mapMemorySpace) {
if (failed(applyFullConversion(gpuModule, *target, std::move(patterns)))) if (failed(applyFullConversion(gpuModule, *target, std::move(patterns))))
return signalPassFailure(); return signalPassFailure();
} }
auto targetAttr = spirv::lookupTargetEnvOrDefault(gpuModule); auto targetAttr = spirv::lookupTargetEnvOrDefault(gpuModule);
std::unique_ptr<ConversionTarget> target = std::unique_ptr<ConversionTarget> target =
SPIRVConversionTarget::get(targetAttr); SPIRVConversionTarget::get(targetAttr);
SPIRVTypeConverter typeConverter(targetAttr); SPIRVConversionOptions options;
options.use64bitIndex = this->use64bitIndex;
SPIRVTypeConverter typeConverter(targetAttr, options);
typeConverter.addConversion([&](gpu::MMAMatrixType type) -> Type { typeConverter.addConversion([&](gpu::MMAMatrixType type) -> Type {
return convertMMAToSPIRVType(type); return convertMMAToSPIRVType(type);
}); });
RewritePatternSet patterns(context); RewritePatternSet patterns(context);
populateGPUToSPIRVPatterns(typeConverter, patterns); populateGPUToSPIRVPatterns(typeConverter, patterns);
populateGpuWMMAToSPIRVConversionPatterns(typeConverter, patterns); populateGpuWMMAToSPIRVConversionPatterns(typeConverter, patterns);
// TODO: Change SPIR-V conversion to be progressive and remove the following // TODO: Change SPIR-V conversion to be progressive and remove the following
// patterns. // patterns.
Show All 13 Lines

mlir/test/Conversion/GPUToSPIRV/builtins.mlir

// 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() { func.func @builtin() {
%c0 = arith.constant 1 : index %c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_workgroup_id_x gpu.launch_func @kernels::@builtin_workgroup_id_x
blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
return return
} }
// CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
// CHECK: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") // INDEX32: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") : !spirv.ptr<vector<3xi32>, Input>
gpu.module @kernels { gpu.module @kernels {
gpu.func @builtin_workgroup_id_x() kernel gpu.func @builtin_workgroup_id_x() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} { attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
// CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPID]] // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPID]]
// CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]] // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
// CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}} // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
%0 = gpu.block_id x %0 = gpu.block_id x
gpu.return gpu.return
} }
} }
} }
// ----- // -----
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() { func.func @builtin() {
%c0 = arith.constant 1 : index %c0 = arith.constant 1 : index
%c256 = arith.constant 256 : i32 %c256 = arith.constant 256 : i32
gpu.launch_func @kernels::@builtin_workgroup_id_y gpu.launch_func @kernels::@builtin_workgroup_id_y
blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
dynamic_shared_memory_size %c256 dynamic_shared_memory_size %c256
return return
} }
// CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
// CHECK: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") // INDEX32: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") : !spirv.ptr<vector<3xi32>, Input>
gpu.module @kernels { gpu.module @kernels {
gpu.func @builtin_workgroup_id_y() kernel gpu.func @builtin_workgroup_id_y() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} { attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
// CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPID]] // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPID]]
// CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]] // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
// CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}} // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
%0 = gpu.block_id y %0 = gpu.block_id y
gpu.return gpu.return
} }
} }
} }
// ----- // -----
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() { func.func @builtin() {
%c0 = arith.constant 1 : index %c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_workgroup_id_z gpu.launch_func @kernels::@builtin_workgroup_id_z
blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
return return
} }
// CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
// CHECK: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") // INDEX32: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") : !spirv.ptr<vector<3xi32>, Input>
gpu.module @kernels { gpu.module @kernels {
gpu.func @builtin_workgroup_id_z() kernel gpu.func @builtin_workgroup_id_z() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} { attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
// CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPID]] // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPID]]
// CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]] // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
// CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}} // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
%0 = gpu.block_id z %0 = gpu.block_id z
gpu.return gpu.return
} }
} }
} }
// ----- // -----
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() { func.func @builtin() {
%c0 = arith.constant 1 : index %c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_workgroup_size_x gpu.launch_func @kernels::@builtin_workgroup_size_x
blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
return return
} }
// CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
gpu.module @kernels { gpu.module @kernels {
gpu.func @builtin_workgroup_size_x() kernel gpu.func @builtin_workgroup_size_x() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 1, 1]>} { attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 1, 1]>} {
// The constant value is obtained from the spirv.entry_point_abi. // The constant value is obtained from the spirv.entry_point_abi.
// Note that this ignores the workgroup size specification in gpu.launch. // 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 // We may want to define gpu.workgroup_size and convert it to the entry
// point ABI we want here. // point ABI we want here.
// CHECK: spirv.Constant 32 : i32 // INDEX32: spirv.Constant 32 : i32
%0 = gpu.block_dim x %0 = gpu.block_dim x
gpu.return gpu.return
} }
} }
} }
// ----- // -----
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() { func.func @builtin() {
%c0 = arith.constant 1 : index %c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_workgroup_size_y gpu.launch_func @kernels::@builtin_workgroup_size_y
blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
return return
} }
// CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
gpu.module @kernels { gpu.module @kernels {
gpu.func @builtin_workgroup_size_y() kernel gpu.func @builtin_workgroup_size_y() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} { 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. // 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 %0 = gpu.block_dim y
gpu.return gpu.return
} }
} }
} }
// ----- // -----
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() { func.func @builtin() {
%c0 = arith.constant 1 : index %c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_workgroup_size_z gpu.launch_func @kernels::@builtin_workgroup_size_z
blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
return return
} }
// CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
gpu.module @kernels { gpu.module @kernels {
gpu.func @builtin_workgroup_size_z() kernel gpu.func @builtin_workgroup_size_z() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} { 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. // 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 %0 = gpu.block_dim z
gpu.return gpu.return
} }
} }
} }
// ----- // -----
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() { func.func @builtin() {
%c0 = arith.constant 1 : index %c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_local_id_x gpu.launch_func @kernels::@builtin_local_id_x
blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
return return
} }
// CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
// CHECK: spirv.GlobalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId") // INDEX32: spirv.GlobalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId") : !spirv.ptr<vector<3xi32>, Input>
gpu.module @kernels { gpu.module @kernels {
gpu.func @builtin_local_id_x() kernel gpu.func @builtin_local_id_x() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} { attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
// CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[LOCALINVOCATIONID]] // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[LOCALINVOCATIONID]]
// CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]] // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
// CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}} // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
%0 = gpu.thread_id x %0 = gpu.thread_id x
gpu.return gpu.return
} }
} }
} }
// ----- // -----
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() { func.func @builtin() {
%c0 = arith.constant 1 : index %c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_num_workgroups_x gpu.launch_func @kernels::@builtin_num_workgroups_x
blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
return return
} }
// CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
// CHECK: spirv.GlobalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups") // INDEX32: spirv.GlobalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups") : !spirv.ptr<vector<3xi32>, Input>
gpu.module @kernels { gpu.module @kernels {
gpu.func @builtin_num_workgroups_x() kernel gpu.func @builtin_num_workgroups_x() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} { attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
// CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[NUMWORKGROUPS]] // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[NUMWORKGROUPS]]
// CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]] // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
// CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}} // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
%0 = gpu.grid_dim x %0 = gpu.grid_dim x
gpu.return gpu.return
} }
} }
} }
// ----- // -----
module attributes {gpu.container_module} { module attributes {
// CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450 gpu.container_module,
// CHECK: spirv.GlobalVariable [[SUBGROUPID:@.*]] built_in("SubgroupId") 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.module @kernels {
gpu.func @builtin_subgroup_id() kernel gpu.func @builtin_subgroup_id() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} { attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
// CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[SUBGROUPID]] // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[SUBGROUPID]]
// CHECK-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]] // INDEX32-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]]
%0 = gpu.subgroup_id : index %0 = gpu.subgroup_id : index
gpu.return gpu.return
} }
} }
} }
// ----- // -----
module attributes {gpu.container_module} { module attributes {
// CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450 gpu.container_module,
// CHECK: spirv.GlobalVariable [[NUMSUBGROUPS:@.*]] built_in("NumSubgroups") 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.module @kernels {
gpu.func @builtin_num_subgroups() kernel gpu.func @builtin_num_subgroups() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} { attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
// CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[NUMSUBGROUPS]] // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[NUMSUBGROUPS]]
// CHECK-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]] // INDEX32-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]]
%0 = gpu.num_subgroups : index %0 = gpu.num_subgroups : index
gpu.return gpu.return
} }
} }
} }
// ----- // -----
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() { func.func @builtin() {
%c0 = arith.constant 1 : index %c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_workgroup_size_x gpu.launch_func @kernels::@builtin_workgroup_size_x
blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
return return
} }
// CHECK-LABEL: spirv.module @{{.*}} // INDEX32-LABEL: spirv.module @{{.*}}
// CHECK: spirv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize") // INDEX32: spirv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize") : !spirv.ptr<vector<3xi32>, Input>
gpu.module @kernels { gpu.module @kernels {
gpu.func @builtin_workgroup_size_x() kernel gpu.func @builtin_workgroup_size_x() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<>} { attributes {spirv.entry_point_abi = #spirv.entry_point_abi<>} {
// CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPSIZE]] // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPSIZE]]
// CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]] // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
// CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}} // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
%0 = gpu.block_dim x %0 = gpu.block_dim x
gpu.return gpu.return
} }
} }
} }
// ----- // -----
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() { func.func @builtin() {
%c0 = arith.constant 1 : index %c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_workgroup_size_y gpu.launch_func @kernels::@builtin_workgroup_size_y
blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
return return
} }
// CHECK-LABEL: spirv.module @{{.*}} // INDEX32-LABEL: spirv.module @{{.*}}
// CHECK: spirv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize") // INDEX32: spirv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize") : !spirv.ptr<vector<3xi32>, Input>
gpu.module @kernels { gpu.module @kernels {
gpu.func @builtin_workgroup_size_y() kernel gpu.func @builtin_workgroup_size_y() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<>} { attributes {spirv.entry_point_abi = #spirv.entry_point_abi<>} {
// CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPSIZE]] // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPSIZE]]
// CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]] // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
// CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}} // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
%0 = gpu.block_dim y %0 = gpu.block_dim y
gpu.return gpu.return
} }
} }
} }
// ----- // -----
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() { func.func @builtin() {
%c0 = arith.constant 1 : index %c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_workgroup_size_z gpu.launch_func @kernels::@builtin_workgroup_size_z
blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
return return
} }
// CHECK-LABEL: spirv.module @{{.*}} // INDEX32-LABEL: spirv.module @{{.*}}
// CHECK: spirv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize") // INDEX32: spirv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize") : !spirv.ptr<vector<3xi32>, Input>
gpu.module @kernels { gpu.module @kernels {
gpu.func @builtin_workgroup_size_z() kernel gpu.func @builtin_workgroup_size_z() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<>} { attributes {spirv.entry_point_abi = #spirv.entry_point_abi<>} {
// CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPSIZE]] // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPSIZE]]
// CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]] // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
// CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}} // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
%0 = gpu.block_dim z %0 = gpu.block_dim z
gpu.return gpu.return
} }
} }
} }
// ----- // -----
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() { func.func @builtin() {
%c0 = arith.constant 1 : index %c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_global_id_x gpu.launch_func @kernels::@builtin_global_id_x
blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
return return
} }
// CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
// CHECK: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId") // INDEX32: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId") : !spirv.ptr<vector<3xi32>, Input>
gpu.module @kernels { gpu.module @kernels {
gpu.func @builtin_global_id_x() kernel gpu.func @builtin_global_id_x() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} { attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
// CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[GLOBALINVOCATIONID]] // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[GLOBALINVOCATIONID]]
// CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]] // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
// CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}} // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
%0 = gpu.global_id x %0 = gpu.global_id x
gpu.return gpu.return
} }
} }
} }
// ----- // -----
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() { func.func @builtin() {
%c0 = arith.constant 1 : index %c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_global_id_y gpu.launch_func @kernels::@builtin_global_id_y
blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
return return
} }
// CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
// CHECK: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId") // INDEX32: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId") : !spirv.ptr<vector<3xi32>, Input>
gpu.module @kernels { gpu.module @kernels {
gpu.func @builtin_global_id_y() kernel gpu.func @builtin_global_id_y() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} { attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
// CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[GLOBALINVOCATIONID]] // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[GLOBALINVOCATIONID]]
// CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]] // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
// CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}} // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
%0 = gpu.global_id y %0 = gpu.global_id y
gpu.return gpu.return
} }
} }
} }
// ----- // -----
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() { func.func @builtin() {
%c0 = arith.constant 1 : index %c0 = arith.constant 1 : index
gpu.launch_func @kernels::@builtin_global_id_z gpu.launch_func @kernels::@builtin_global_id_z
blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0) blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
return return
} }
// CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450 // INDEX32-LABEL: spirv.module @{{.*}} Logical GLSL450
// CHECK: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId") // INDEX32: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId") : !spirv.ptr<vector<3xi32>, Input>
gpu.module @kernels { gpu.module @kernels {
gpu.func @builtin_global_id_z() kernel gpu.func @builtin_global_id_z() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} { attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
// CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[GLOBALINVOCATIONID]] // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[GLOBALINVOCATIONID]]
// CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]] // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
// CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}} // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
%0 = gpu.global_id z %0 = gpu.global_id z
gpu.return gpu.return
} }
} }
} }
// ----- // -----
module attributes {gpu.container_module} { module attributes {
// CHECK-LABEL: spirv.module @{{.*}} Logical GLSL450 gpu.container_module,
// CHECK: spirv.GlobalVariable [[SUBGROUPSIZE:@.*]] built_in("SubgroupSize") 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.module @kernels {
gpu.func @builtin_subgroup_size() kernel gpu.func @builtin_subgroup_size() kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} { attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
// CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[SUBGROUPSIZE]] // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[SUBGROUPSIZE]]
// CHECK-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]] // INDEX32-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]]
%0 = gpu.subgroup_size : index %0 = gpu.subgroup_size : index
gpu.return gpu.return
} }
} }
} }