Index: mlir/include/mlir/Dialect/SPIRV/ExecutionEnvironment.h =================================================================== --- /dev/null +++ mlir/include/mlir/Dialect/SPIRV/ExecutionEnvironment.h @@ -0,0 +1,43 @@ +//===-- ExecutionEnvironment.h - SPIR-V execution environment utilities --===// +// +// Part of the LLVM Project, 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 +// +//===----------------------------------------------------------------------===// +// +// This file defines utilities used to select values describing SPIR-V +// execution environment based on target environment description. +// +//===----------------------------------------------------------------------===// +#ifndef MLIR_DIALECT_SPIRV_EXECUTIONENVIRONMENT_H_ +#define MLIR_DIALECT_SPIRV_EXECUTIONENVIRONMENT_H_ + +#include "mlir/Support/LogicalResult.h" +#include + +namespace mlir { +namespace spirv { + +enum class AddressingModel : uint32_t; +enum class MemoryModel : uint32_t; +enum class ExecutionModel : uint32_t; +class TargetEnvAttr; + +/// Returns addressing model selected based on target environment. +AddressingModel getAddressingModel(TargetEnvAttr targetAttr); + +/// Returns memory model selected based on target environment. +FailureOr getMemoryModel(TargetEnvAttr targetAttr); + +/// Returns execution model selected based on target environment. +FailureOr getExecutionModel(TargetEnvAttr targetAttr); + +/// Returns whether SPIR-V target (described by TargetEnvAttr) needs +/// ABI attributes for interface variables (spv.interface_var_abi). +bool needsInterfaceVarABIAttrs(TargetEnvAttr targetAttr); + +} // namespace spirv +} // namespace mlir + +#endif // MLIR_DIALECT_SPIRV_EXECUTIONENVIRONMENT_H_ Index: mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp =================================================================== --- mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp +++ mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp @@ -11,9 +11,11 @@ //===----------------------------------------------------------------------===// #include "mlir/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.h" #include "mlir/Dialect/GPU/GPUDialect.h" +#include "mlir/Dialect/SPIRV/ExecutionEnvironment.h" #include "mlir/Dialect/SPIRV/SPIRVDialect.h" #include "mlir/Dialect/SPIRV/SPIRVLowering.h" #include "mlir/Dialect/SPIRV/SPIRVOps.h" +#include "mlir/Dialect/SPIRV/TargetAndABI.h" #include "mlir/IR/Module.h" using namespace mlir; @@ -170,9 +172,10 @@ "with no return values right now"); return nullptr; } - if (fnType.getNumInputs() != argABIInfo.size()) { + if (!argABIInfo.empty() && fnType.getNumInputs() != argABIInfo.size()) { funcOp.emitError( - "lowering as entry functions requires ABI info for all arguments"); + "lowering as entry functions requires ABI info for all arguments " + "or none of them"); return nullptr; } // Update the signature to valid SPIR-V types and add the ABI @@ -202,7 +205,6 @@ &signatureConverter))) return nullptr; rewriter.eraseOp(funcOp); - spirv::setABIAttrs(newFuncOp, entryPointInfo, argABIInfo); return newFuncOp; } @@ -213,6 +215,10 @@ static LogicalResult getDefaultABIAttrs(MLIRContext *context, gpu::GPUFuncOp funcOp, SmallVectorImpl &argABI) { + auto targetEnv = spirv::lookupTargetEnvOrDefault(funcOp); + if (!spirv::needsInterfaceVarABIAttrs(targetEnv)) + return success(); + for (auto argIndex : llvm::seq(0, funcOp.getNumArguments())) { if (funcOp.getArgAttrOfType( argIndex, spirv::getInterfaceVarABIAttrName())) @@ -272,9 +278,15 @@ LogicalResult GPUModuleConversion::matchAndRewrite( gpu::GPUModuleOp moduleOp, ArrayRef operands, ConversionPatternRewriter &rewriter) const { + auto targetEnv = spirv::lookupTargetEnvOrDefault(moduleOp); + auto addressingModel = spirv::getAddressingModel(targetEnv); + auto memoryModel = spirv::getMemoryModel(targetEnv); + if (failed(memoryModel)) + return moduleOp.emitRemark("match failure: could not selected memory model " + "based on 'spv.target_env'"); + auto spvModule = rewriter.create( - moduleOp.getLoc(), spirv::AddressingModel::Logical, - spirv::MemoryModel::GLSL450); + moduleOp.getLoc(), addressingModel, memoryModel.getValue()); // Move the region from the module op into the SPIR-V module. Region &spvModuleRegion = spvModule.body(); Index: mlir/lib/Dialect/SPIRV/CMakeLists.txt =================================================================== --- mlir/lib/Dialect/SPIRV/CMakeLists.txt +++ mlir/lib/Dialect/SPIRV/CMakeLists.txt @@ -4,6 +4,7 @@ add_public_tablegen_target(MLIRSPIRVCanonicalizationIncGen) add_mlir_dialect_library(MLIRSPIRV + ExecutionEnvironment.cpp LayoutUtils.cpp SPIRVAttributes.cpp SPIRVCanonicalization.cpp Index: mlir/lib/Dialect/SPIRV/ExecutionEnvironment.cpp =================================================================== --- /dev/null +++ mlir/lib/Dialect/SPIRV/ExecutionEnvironment.cpp @@ -0,0 +1,60 @@ +//===-- ExecutionEnvironment.cpp - SPIR-V execution environment utilities -===// +// +// Part of the LLVM Project, 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 +// +//===----------------------------------------------------------------------===// +// +// This file implements utilities used to select values describing SPIR-V +// execution environment based on target environment description. +// +//===----------------------------------------------------------------------===// +#include "mlir/Dialect/SPIRV/ExecutionEnvironment.h" +#include "mlir/Dialect/SPIRV/SPIRVAttributes.h" +#include "mlir/Dialect/SPIRV/SPIRVTypes.h" + +namespace mlir { +namespace spirv { + +AddressingModel getAddressingModel(TargetEnvAttr targetAttr) { + for (auto cap : targetAttr.getCapabilities()) { + if (cap == Capability::Kernel) + return AddressingModel::Physical64; + } + // Logical addressing doesn't need any capabilities so return it as default. + return AddressingModel::Logical; +} + +FailureOr getMemoryModel(TargetEnvAttr targetAttr) { + for (auto cap : targetAttr.getCapabilities()) { + if (cap == Capability::Addresses) + return MemoryModel::OpenCL; + if (cap == Capability::Shader) + return MemoryModel::GLSL450; + } + return failure(); +} + +FailureOr getExecutionModel(TargetEnvAttr targetAttr) { + for (auto cap : targetAttr.getCapabilities()) { + if (cap == Capability::Kernel) + return ExecutionModel::Kernel; + if (cap == Capability::Shader) + return ExecutionModel::GLCompute; + } + return failure(); +} + +bool needsInterfaceVarABIAttrs(TargetEnvAttr targetAttr) { + for (auto cap : targetAttr.getCapabilities()) { + if (cap == Capability::Kernel) + return false; + if (cap == Capability::Shader) + return true; + } + return false; +} + +} // namespace spirv +} // namespace mlir Index: mlir/lib/Dialect/SPIRV/SPIRVLowering.cpp =================================================================== --- mlir/lib/Dialect/SPIRV/SPIRVLowering.cpp +++ mlir/lib/Dialect/SPIRV/SPIRVLowering.cpp @@ -653,7 +653,7 @@ ArrayRef argABIInfo) { // Set the attributes for argument and the function. StringRef argABIAttrName = spirv::getInterfaceVarABIAttrName(); - for (auto argIndex : llvm::seq(0, funcOp.getNumArguments())) { + for (auto argIndex : llvm::seq(0, argABIInfo.size())) { funcOp.setArgAttr(argIndex, argABIAttrName, argABIInfo[argIndex]); } funcOp.setAttr(spirv::getEntryPointABIAttrName(), entryPointInfo); Index: mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp =================================================================== --- mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp +++ mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp @@ -12,6 +12,7 @@ //===----------------------------------------------------------------------===// #include "PassDetail.h" +#include "mlir/Dialect/SPIRV/ExecutionEnvironment.h" #include "mlir/Dialect/SPIRV/LayoutUtils.h" #include "mlir/Dialect/SPIRV/Passes.h" #include "mlir/Dialect/SPIRV/SPIRVDialect.h" @@ -119,8 +120,16 @@ if (failed(getInterfaceVariables(funcOp, interfaceVars))) { return failure(); } + + auto targetEnv = spirv::lookupTargetEnv(funcOp); + auto executionModel = spirv::getExecutionModel(targetEnv); + if (failed(executionModel)) + return funcOp.emitRemark("lower entry point failure: could not select " + "execution model based on 'spv.target_env'"); + builder.create( - funcOp.getLoc(), spirv::ExecutionModel::GLCompute, funcOp, interfaceVars); + funcOp.getLoc(), executionModel.getValue(), funcOp, interfaceVars); + // Specifies the spv.ExecutionModeOp. auto localSizeAttr = entryPointAttr.local_size(); SmallVector localSize(localSizeAttr.getValues()); Index: mlir/test/Conversion/GPUToSPIRV/test_opencl_spirv.mlir =================================================================== --- /dev/null +++ mlir/test/Conversion/GPUToSPIRV/test_opencl_spirv.mlir @@ -0,0 +1,33 @@ +// RUN: mlir-opt -allow-unregistered-dialect -convert-gpu-to-spirv -verify-diagnostics %s -o - | FileCheck %s + +module attributes { + gpu.container_module, + spv.target_env = #spv.target_env< + #spv.vce, + {max_compute_workgroup_invocations = 128 : i32, + max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> +} { + gpu.module @kernels { + // CHECK: spv.module Physical64 OpenCL { + // CHECK-LABEL: spv.func @basic_module_structure + // CHECK-SAME: {{%.*}}: f32 + // CHECK-NOT: spv.interface_var_abi + // CHECK-SAME: {{%.*}}: !spv.ptr [0]>, CrossWorkgroup> + // CHECK-NOT: spv.interface_var_abi + // CHECK-SAME: spv.entry_point_abi = {local_size = dense<[32, 4, 1]> : vector<3xi32>} + gpu.func @basic_module_structure(%arg0 : f32, %arg1 : memref<12xf32, 11>) kernel + attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} { + // CHECK: spv.Return + gpu.return + } + } + + func @main() { + %0 = "op"() : () -> (f32) + %1 = "op"() : () -> (memref<12xf32, 11>) + %cst = constant 1 : index + "gpu.launch_func"(%cst, %cst, %cst, %cst, %cst, %cst, %0, %1) { kernel = @kernels::@basic_module_structure } + : (index, index, index, index, index, index, f32, memref<12xf32, 11>) -> () + return + } +} Index: mlir/test/Dialect/SPIRV/Transforms/abi-interface-opencl.mlir =================================================================== --- /dev/null +++ mlir/test/Dialect/SPIRV/Transforms/abi-interface-opencl.mlir @@ -0,0 +1,26 @@ +// RUN: mlir-opt -spirv-lower-abi-attrs -verify-diagnostics %s -o - | FileCheck %s + +module attributes { + spv.target_env = #spv.target_env< + #spv.vce, + {max_compute_workgroup_invocations = 128 : i32, + max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> +} { + +// CHECK: spv.module Physical64 OpenCL +spv.module Physical64 OpenCL { + // CHECK-LABEL: spv.func @kernel + // CHECK-SAME: {{%.*}}: f32 + // CHECK-SAME: {{%.*}}: !spv.ptr>, CrossWorkgroup> + spv.func @kernel( + %arg0: f32, + %arg1: !spv.ptr>, CrossWorkgroup>) "None" + attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} { + // CHECK: spv.Return + spv.Return + } + // CHECK: spv.EntryPoint "Kernel" @kernel + // CHECK: spv.ExecutionMode @kernel "LocalSize", 32, 1, 1 +} // end spv.module + +} // end module