9dbb6e1978
There are execution modes need target width as their extra operands. SignedZeroInfNanPreserve is one of them. This patch adds `target width` as one of SPIR-V ABI attributes.
27 lines
1.3 KiB
MLIR
27 lines
1.3 KiB
MLIR
// RUN: mlir-opt -test-spirv-entry-point-abi %s | FileCheck %s -check-prefix=DEFAULT
|
|
// RUN: mlir-opt -test-spirv-entry-point-abi="workgroup-size=32" %s | FileCheck %s -check-prefix=WG32
|
|
// RUN: mlir-opt -test-spirv-entry-point-abi="subgroup-size=4" %s | FileCheck %s -check-prefix=SG4
|
|
// RUN: mlir-opt -test-spirv-entry-point-abi="target-width=32" %s | FileCheck %s -check-prefix=TW32
|
|
// RUN: mlir-opt -test-spirv-entry-point-abi="workgroup-size=32,8 subgroup-size=4 target-width=32" %s | FileCheck %s -check-prefix=WG32_8-SG4-TW32
|
|
|
|
// DEFAULT: gpu.func @foo()
|
|
// DEFAULT-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [1, 1, 1]>
|
|
|
|
// WG32: gpu.func @foo()
|
|
// WG32-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 1, 1]>
|
|
|
|
// SG4: gpu.func @foo()
|
|
// SG4-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [1, 1, 1], subgroup_size = 4>
|
|
|
|
// TW32: gpu.func @foo()
|
|
// TW32-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [1, 1, 1], target_width = 32>
|
|
|
|
// WG32_8-SG4-TW32: gpu.func @foo()
|
|
// WG32_8-SG4-TW32-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 8, 1], subgroup_size = 4, target_width = 32>
|
|
|
|
gpu.module @kernels {
|
|
gpu.func @foo() kernel {
|
|
gpu.return
|
|
}
|
|
}
|