[clang-tools-extra] [MLIR] Update convert-gpu-to-spirv pass to prepare using GPU compilat… (PR #69941)
Sang Ik Lee via cfe-commits
cfe-commits at lists.llvm.org
Fri Oct 27 18:10:47 PDT 2023
https://github.com/silee2 updated https://github.com/llvm/llvm-project/pull/69941
>From 7ca3f97b5ee6e5cefd94afd3b090d0dba2120cea Mon Sep 17 00:00:00 2001
From: "Lee, Sang Ik" <sang.ik.lee at intel.com>
Date: Mon, 23 Oct 2023 16:25:15 +0000
Subject: [PATCH 1/5] [MLIR] Update convert-gpu-to-spirv pass to prepare using
GPU compilation pipeline for OpenCL kernels.
This includes a couple of changes to pass behavior for OpenCL kernels.
Vulkan shaders are not impacted by the changes.
1. SPIRV module is placed inside GPU module. This change is required for
gpu-module-to-binary to work correctly as it expects kernel function to be
inside the GPU module.
2. A dummy func.func with same kernel name as gpu.func is created.
GPU compilation pipeline defers lowering of gpu launch kernel op.
Since spirv.func is not directly tied to gpu launch kernel,
a dummy func.func is required to avoid legalization issues.
3. Use correct mapping when mapping MemRef memory space to SPIR-V storage class for OpenCL kernels.
---
mlir/include/mlir/Conversion/Passes.td | 5 +-
.../Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp | 76 ++++++++++++++++---
.../Conversion/GPUToSPIRV/module-opencl.mlir | 4 +
3 files changed, 75 insertions(+), 10 deletions(-)
diff --git a/mlir/include/mlir/Conversion/Passes.td b/mlir/include/mlir/Conversion/Passes.td
index 274784fe4a7b29c..652ef5ad95158ca 100644
--- a/mlir/include/mlir/Conversion/Passes.td
+++ b/mlir/include/mlir/Conversion/Passes.td
@@ -578,7 +578,10 @@ def ConvertGPUToSPIRV : Pass<"convert-gpu-to-spirv", "ModuleOp"> {
to control the set and binding if wanted.
}];
let constructor = "mlir::createConvertGPUToSPIRVPass()";
- let dependentDialects = ["spirv::SPIRVDialect"];
+ let dependentDialects = [
+ "spirv::SPIRVDialect",
+ "func::FuncDialect",
+ ];
let options = [
Option<"use64bitIndex", "use-64bit-index",
"bool", /*default=*/"false",
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
index 272e3de8723aeb6..35ee0d7038a2c9a 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
@@ -17,6 +17,7 @@
#include "mlir/Conversion/FuncToSPIRV/FuncToSPIRV.h"
#include "mlir/Conversion/GPUToSPIRV/GPUToSPIRV.h"
#include "mlir/Conversion/MemRefToSPIRV/MemRefToSPIRV.h"
+#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVOps.h"
@@ -54,22 +55,63 @@ void GPUToSPIRVPass::runOnOperation() {
SmallVector<Operation *, 1> gpuModules;
OpBuilder builder(context);
+
+ auto getTargetEnvFromGPUModuleOp = [=](gpu::GPUModuleOp moduleOp) {
+ Operation *gpuModule = moduleOp.getOperation();
+ auto targetAttr = spirv::lookupTargetEnvOrDefault(gpuModule);
+ std::unique_ptr<ConversionTarget> target =
+ SPIRVConversionTarget::get(targetAttr);
+
+ SPIRVConversionOptions options;
+ options.use64bitIndex = this->use64bitIndex;
+ SPIRVTypeConverter typeConverter(targetAttr, options);
+ const spirv::TargetEnv &targetEnv = typeConverter.getTargetEnv();
+ return targetEnv;
+ };
+
module.walk([&](gpu::GPUModuleOp moduleOp) {
// Clone each GPU kernel module for conversion, given that the GPU
// launch op still needs the original GPU kernel module.
- builder.setInsertionPoint(moduleOp.getOperation());
+ // SPIRV module insertion point by is after original GPU module.
+ // This works fine for Vulkan shader that has a dedicated runner.
+ // But OpenCL kernel needs SPIRV module placed inside original GPU module as
+ // OpenCL uses GPU compilation pipeline.
+ auto targetEnv = getTargetEnvFromGPUModuleOp(moduleOp);
+ FailureOr<spirv::MemoryModel> memoryModel =
+ spirv::getMemoryModel(targetEnv);
+ if (failed(memoryModel))
+ return signalPassFailure();
+ (memoryModel == spirv::MemoryModel::OpenCL)
+ ? builder.setInsertionPoint(moduleOp.getBody(),
+ moduleOp.getBody()->begin())
+ : builder.setInsertionPoint(moduleOp.getOperation());
gpuModules.push_back(builder.clone(*moduleOp.getOperation()));
});
// Run conversion for each module independently as they can have different
// TargetEnv attributes.
for (Operation *gpuModule : gpuModules) {
+ auto targetAttr = spirv::lookupTargetEnvOrDefault(gpuModule);
+ std::unique_ptr<ConversionTarget> target =
+ SPIRVConversionTarget::get(targetAttr);
+
+ SPIRVConversionOptions options;
+ options.use64bitIndex = this->use64bitIndex;
+ SPIRVTypeConverter typeConverter(targetAttr, options);
+ const spirv::TargetEnv &targetEnv = typeConverter.getTargetEnv();
+ FailureOr<spirv::MemoryModel> memoryModel =
+ spirv::getMemoryModel(targetEnv);
+ if (failed(memoryModel))
+ return signalPassFailure();
+
// Map MemRef memory space to SPIR-V storage class first if requested.
if (mapMemorySpace) {
std::unique_ptr<ConversionTarget> target =
spirv::getMemorySpaceToStorageClassTarget(*context);
spirv::MemorySpaceToStorageClassMap memorySpaceMap =
- spirv::mapMemorySpaceToVulkanStorageClass;
+ (memoryModel == spirv::MemoryModel::OpenCL)
+ ? spirv::mapMemorySpaceToOpenCLStorageClass
+ : spirv::mapMemorySpaceToVulkanStorageClass;
spirv::MemorySpaceToStorageClassConverter converter(memorySpaceMap);
RewritePatternSet patterns(context);
@@ -79,13 +121,6 @@ void GPUToSPIRVPass::runOnOperation() {
return signalPassFailure();
}
- auto targetAttr = spirv::lookupTargetEnvOrDefault(gpuModule);
- std::unique_ptr<ConversionTarget> target =
- SPIRVConversionTarget::get(targetAttr);
-
- SPIRVConversionOptions options;
- options.use64bitIndex = this->use64bitIndex;
- SPIRVTypeConverter typeConverter(targetAttr, options);
populateMMAToSPIRVCoopMatrixTypeConversion(typeConverter,
this->useCoopMatrixNV);
@@ -108,6 +143,29 @@ void GPUToSPIRVPass::runOnOperation() {
if (failed(applyFullConversion(gpuModule, *target, std::move(patterns))))
return signalPassFailure();
}
+ // In case of OpenCL, gpu.func in original gpu.module needs to replaced with
+ // an empty func.func with same arguments as gpu.func. And it also needs
+ // gpu.kernel attribute set.
+ module.walk([&](gpu::GPUModuleOp moduleOp) {
+ auto targetEnv = getTargetEnvFromGPUModuleOp(moduleOp);
+ FailureOr<spirv::MemoryModel> memoryModel =
+ spirv::getMemoryModel(targetEnv);
+ if (failed(memoryModel))
+ return signalPassFailure();
+ if (memoryModel == spirv::MemoryModel::OpenCL) {
+ moduleOp.walk([&](gpu::GPUFuncOp funcOp) {
+ builder.setInsertionPoint(funcOp);
+ auto newFuncOp = builder.create<func::FuncOp>(
+ funcOp.getLoc(), funcOp.getName(), funcOp.getFunctionType());
+ auto entryBlock = newFuncOp.addEntryBlock();
+ builder.setInsertionPointToEnd(entryBlock);
+ builder.create<func::ReturnOp>(funcOp.getLoc());
+ newFuncOp->setAttr(gpu::GPUDialect::getKernelFuncAttrName(),
+ builder.getUnitAttr());
+ funcOp.erase();
+ });
+ }
+ });
}
} // namespace
diff --git a/mlir/test/Conversion/GPUToSPIRV/module-opencl.mlir b/mlir/test/Conversion/GPUToSPIRV/module-opencl.mlir
index 0aa50cc1e25294d..4b8d17cd6449389 100644
--- a/mlir/test/Conversion/GPUToSPIRV/module-opencl.mlir
+++ b/mlir/test/Conversion/GPUToSPIRV/module-opencl.mlir
@@ -12,6 +12,8 @@ module attributes {
// CHECK-SAME: {{%.*}}: !spirv.ptr<!spirv.array<12 x f32>, CrossWorkgroup>
// CHECK-NOT: spirv.interface_var_abi
// CHECK-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>
+ // CHECK-LABEL: func.func @basic_module_structure
+ // CHECK-SAME: attributes {gpu.kernel}
gpu.func @basic_module_structure(%arg0 : f32, %arg1 : memref<12xf32, #spirv.storage_class<CrossWorkgroup>>) kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
gpu.return
@@ -45,6 +47,8 @@ module attributes {
// CHECK-SAME: {{%.*}}: !spirv.ptr<!spirv.array<12 x f32>, CrossWorkgroup>
// CHECK-NOT: spirv.interface_var_abi
// CHECK-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>
+ // CHECK-LABEL: func.func @basic_module_structure
+ // CHECK-SAME: attributes {gpu.kernel}
gpu.func @basic_module_structure(%arg0 : f32, %arg1 : memref<12xf32, #spirv.storage_class<CrossWorkgroup>>) kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
gpu.return
>From 02ff96c8185e30b730e5bf0230699a1abbc4dcc6 Mon Sep 17 00:00:00 2001
From: "Lee, Sang Ik" <sang.ik.lee at intel.com>
Date: Fri, 27 Oct 2023 16:53:45 +0000
Subject: [PATCH 2/5] Address reviewer comments. Sort alphabetically. Use
if-else instead of ternary expression. List lambda captures explicitly.
Replace auto with actually type. Add newline before big block.
---
mlir/include/mlir/Conversion/Passes.td | 2 +-
.../Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp | 19 +++++++++++--------
2 files changed, 12 insertions(+), 9 deletions(-)
diff --git a/mlir/include/mlir/Conversion/Passes.td b/mlir/include/mlir/Conversion/Passes.td
index 652ef5ad95158ca..cc9b7d9b4f8a871 100644
--- a/mlir/include/mlir/Conversion/Passes.td
+++ b/mlir/include/mlir/Conversion/Passes.td
@@ -579,8 +579,8 @@ def ConvertGPUToSPIRV : Pass<"convert-gpu-to-spirv", "ModuleOp"> {
}];
let constructor = "mlir::createConvertGPUToSPIRVPass()";
let dependentDialects = [
- "spirv::SPIRVDialect",
"func::FuncDialect",
+ "spirv::SPIRVDialect",
];
let options = [
Option<"use64bitIndex", "use-64bit-index",
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
index 35ee0d7038a2c9a..e1d3268bca08942 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
@@ -56,7 +56,7 @@ void GPUToSPIRVPass::runOnOperation() {
SmallVector<Operation *, 1> gpuModules;
OpBuilder builder(context);
- auto getTargetEnvFromGPUModuleOp = [=](gpu::GPUModuleOp moduleOp) {
+ auto getTargetEnvFromGPUModuleOp = [*this](gpu::GPUModuleOp moduleOp) {
Operation *gpuModule = moduleOp.getOperation();
auto targetAttr = spirv::lookupTargetEnvOrDefault(gpuModule);
std::unique_ptr<ConversionTarget> target =
@@ -76,22 +76,24 @@ void GPUToSPIRVPass::runOnOperation() {
// This works fine for Vulkan shader that has a dedicated runner.
// But OpenCL kernel needs SPIRV module placed inside original GPU module as
// OpenCL uses GPU compilation pipeline.
- auto targetEnv = getTargetEnvFromGPUModuleOp(moduleOp);
+ const mlir::spirv::TargetEnv& targetEnv = getTargetEnvFromGPUModuleOp(moduleOp);
FailureOr<spirv::MemoryModel> memoryModel =
spirv::getMemoryModel(targetEnv);
if (failed(memoryModel))
return signalPassFailure();
- (memoryModel == spirv::MemoryModel::OpenCL)
- ? builder.setInsertionPoint(moduleOp.getBody(),
- moduleOp.getBody()->begin())
- : builder.setInsertionPoint(moduleOp.getOperation());
+ if (memoryModel == spirv::MemoryModel::OpenCL) {
+ builder.setInsertionPoint(moduleOp.getBody(),
+ moduleOp.getBody()->begin());
+ } else {
+ builder.setInsertionPoint(moduleOp.getOperation());
+ }
gpuModules.push_back(builder.clone(*moduleOp.getOperation()));
});
// Run conversion for each module independently as they can have different
// TargetEnv attributes.
for (Operation *gpuModule : gpuModules) {
- auto targetAttr = spirv::lookupTargetEnvOrDefault(gpuModule);
+ mlir::spirv::TargetEnvAttr targetAttr = spirv::lookupTargetEnvOrDefault(gpuModule);
std::unique_ptr<ConversionTarget> target =
SPIRVConversionTarget::get(targetAttr);
@@ -143,11 +145,12 @@ void GPUToSPIRVPass::runOnOperation() {
if (failed(applyFullConversion(gpuModule, *target, std::move(patterns))))
return signalPassFailure();
}
+
// In case of OpenCL, gpu.func in original gpu.module needs to replaced with
// an empty func.func with same arguments as gpu.func. And it also needs
// gpu.kernel attribute set.
module.walk([&](gpu::GPUModuleOp moduleOp) {
- auto targetEnv = getTargetEnvFromGPUModuleOp(moduleOp);
+ const mlir::spirv::TargetEnv& targetEnv = getTargetEnvFromGPUModuleOp(moduleOp);
FailureOr<spirv::MemoryModel> memoryModel =
spirv::getMemoryModel(targetEnv);
if (failed(memoryModel))
>From 561af5dc6dd106d47eb1c065ee701893d117957e Mon Sep 17 00:00:00 2001
From: "Lee, Sang Ik" <sang.ik.lee at intel.com>
Date: Fri, 27 Oct 2023 17:13:25 +0000
Subject: [PATCH 3/5] Run clang formatter.
---
mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp | 11 +++++++----
1 file changed, 7 insertions(+), 4 deletions(-)
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
index e1d3268bca08942..0ee0bc2bd02e33f 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
@@ -76,14 +76,15 @@ void GPUToSPIRVPass::runOnOperation() {
// This works fine for Vulkan shader that has a dedicated runner.
// But OpenCL kernel needs SPIRV module placed inside original GPU module as
// OpenCL uses GPU compilation pipeline.
- const mlir::spirv::TargetEnv& targetEnv = getTargetEnvFromGPUModuleOp(moduleOp);
+ const mlir::spirv::TargetEnv &targetEnv =
+ getTargetEnvFromGPUModuleOp(moduleOp);
FailureOr<spirv::MemoryModel> memoryModel =
spirv::getMemoryModel(targetEnv);
if (failed(memoryModel))
return signalPassFailure();
if (memoryModel == spirv::MemoryModel::OpenCL) {
builder.setInsertionPoint(moduleOp.getBody(),
- moduleOp.getBody()->begin());
+ moduleOp.getBody()->begin());
} else {
builder.setInsertionPoint(moduleOp.getOperation());
}
@@ -93,7 +94,8 @@ void GPUToSPIRVPass::runOnOperation() {
// Run conversion for each module independently as they can have different
// TargetEnv attributes.
for (Operation *gpuModule : gpuModules) {
- mlir::spirv::TargetEnvAttr targetAttr = spirv::lookupTargetEnvOrDefault(gpuModule);
+ mlir::spirv::TargetEnvAttr targetAttr =
+ spirv::lookupTargetEnvOrDefault(gpuModule);
std::unique_ptr<ConversionTarget> target =
SPIRVConversionTarget::get(targetAttr);
@@ -150,7 +152,8 @@ void GPUToSPIRVPass::runOnOperation() {
// an empty func.func with same arguments as gpu.func. And it also needs
// gpu.kernel attribute set.
module.walk([&](gpu::GPUModuleOp moduleOp) {
- const mlir::spirv::TargetEnv& targetEnv = getTargetEnvFromGPUModuleOp(moduleOp);
+ const mlir::spirv::TargetEnv &targetEnv =
+ getTargetEnvFromGPUModuleOp(moduleOp);
FailureOr<spirv::MemoryModel> memoryModel =
spirv::getMemoryModel(targetEnv);
if (failed(memoryModel))
>From 6417cb77714bb9902c69629bafb489410fc57798 Mon Sep 17 00:00:00 2001
From: Sang Ik Lee <sang.ik.lee at intel.com>
Date: Fri, 27 Oct 2023 17:53:26 -0700
Subject: [PATCH 4/5] Address reviewers suggestions.
---
.../Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp | 44 ++++++-------------
1 file changed, 14 insertions(+), 30 deletions(-)
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
index 0ee0bc2bd02e33f..aedd2fdd6668519 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
@@ -56,33 +56,23 @@ void GPUToSPIRVPass::runOnOperation() {
SmallVector<Operation *, 1> gpuModules;
OpBuilder builder(context);
- auto getTargetEnvFromGPUModuleOp = [*this](gpu::GPUModuleOp moduleOp) {
+ auto targetEnvSupportsKernelCapability = [](gpu::GPUModuleOp moduleOp) {
Operation *gpuModule = moduleOp.getOperation();
auto targetAttr = spirv::lookupTargetEnvOrDefault(gpuModule);
- std::unique_ptr<ConversionTarget> target =
- SPIRVConversionTarget::get(targetAttr);
-
- SPIRVConversionOptions options;
- options.use64bitIndex = this->use64bitIndex;
- SPIRVTypeConverter typeConverter(targetAttr, options);
- const spirv::TargetEnv &targetEnv = typeConverter.getTargetEnv();
- return targetEnv;
+ spirv::TargetEnv targetEnv(targetAttr);
+ return targetEnv.allows(spirv::Capability::Kernel);
};
module.walk([&](gpu::GPUModuleOp moduleOp) {
// Clone each GPU kernel module for conversion, given that the GPU
// launch op still needs the original GPU kernel module.
- // SPIRV module insertion point by is after original GPU module.
- // This works fine for Vulkan shader that has a dedicated runner.
- // But OpenCL kernel needs SPIRV module placed inside original GPU module as
- // OpenCL uses GPU compilation pipeline.
- const mlir::spirv::TargetEnv &targetEnv =
- getTargetEnvFromGPUModuleOp(moduleOp);
- FailureOr<spirv::MemoryModel> memoryModel =
- spirv::getMemoryModel(targetEnv);
- if (failed(memoryModel))
- return signalPassFailure();
- if (memoryModel == spirv::MemoryModel::OpenCL) {
+ // For Vulkan Shader capabilities, we insert the newly converted SPIR-V
+ // module right after the original GPU module, as that's the expectation of
+ // the in-tree Vulkan runner.
+ // For OpenCL Kernel capabilities, we insert the newly converted SPIR-V
+ // module inside the original GPU module, as that's the expectaion of the
+ // normal GPU compilation pipeline.
+ if (targetEnvSupportsKernelCapability(moduleOp)) {
builder.setInsertionPoint(moduleOp.getBody(),
moduleOp.getBody()->begin());
} else {
@@ -148,17 +138,11 @@ void GPUToSPIRVPass::runOnOperation() {
return signalPassFailure();
}
- // In case of OpenCL, gpu.func in original gpu.module needs to replaced with
- // an empty func.func with same arguments as gpu.func. And it also needs
- // gpu.kernel attribute set.
+ // For OpenCL, the gpu.func op in the original gpu.module op needs to be
+ // replaced with an empty func.func op with the same arguments as the gpu.func
+ // op. The func.func op needs gpu.kernel attribute set.
module.walk([&](gpu::GPUModuleOp moduleOp) {
- const mlir::spirv::TargetEnv &targetEnv =
- getTargetEnvFromGPUModuleOp(moduleOp);
- FailureOr<spirv::MemoryModel> memoryModel =
- spirv::getMemoryModel(targetEnv);
- if (failed(memoryModel))
- return signalPassFailure();
- if (memoryModel == spirv::MemoryModel::OpenCL) {
+ if (targetEnvSupportsKernelCapability(moduleOp)) {
moduleOp.walk([&](gpu::GPUFuncOp funcOp) {
builder.setInsertionPoint(funcOp);
auto newFuncOp = builder.create<func::FuncOp>(
>From 57c0d761c597b8566885c9c4bd246c6b9d46e9fb Mon Sep 17 00:00:00 2001
From: Sang Ik Lee <sang.ik.lee at intel.com>
Date: Fri, 27 Oct 2023 18:09:58 -0700
Subject: [PATCH 5/5] Check for spirv::MemoryModel only if mapping MemRef
memory space to SPIR-V storage class.
---
.../Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp | 23 ++++++++++---------
1 file changed, 12 insertions(+), 11 deletions(-)
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
index aedd2fdd6668519..d5aa5e37a7cbfab 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
@@ -86,20 +86,15 @@ void GPUToSPIRVPass::runOnOperation() {
for (Operation *gpuModule : gpuModules) {
mlir::spirv::TargetEnvAttr targetAttr =
spirv::lookupTargetEnvOrDefault(gpuModule);
- std::unique_ptr<ConversionTarget> target =
- SPIRVConversionTarget::get(targetAttr);
-
- SPIRVConversionOptions options;
- options.use64bitIndex = this->use64bitIndex;
- SPIRVTypeConverter typeConverter(targetAttr, options);
- const spirv::TargetEnv &targetEnv = typeConverter.getTargetEnv();
- FailureOr<spirv::MemoryModel> memoryModel =
- spirv::getMemoryModel(targetEnv);
- if (failed(memoryModel))
- return signalPassFailure();
// Map MemRef memory space to SPIR-V storage class first if requested.
if (mapMemorySpace) {
+ spirv::TargetEnv targetEnv(targetAttr);
+ FailureOr<spirv::MemoryModel> memoryModel =
+ spirv::getMemoryModel(targetEnv);
+ if (failed(memoryModel))
+ return signalPassFailure();
+
std::unique_ptr<ConversionTarget> target =
spirv::getMemorySpaceToStorageClassTarget(*context);
spirv::MemorySpaceToStorageClassMap memorySpaceMap =
@@ -115,6 +110,12 @@ void GPUToSPIRVPass::runOnOperation() {
return signalPassFailure();
}
+ std::unique_ptr<ConversionTarget> target =
+ SPIRVConversionTarget::get(targetAttr);
+
+ SPIRVConversionOptions options;
+ options.use64bitIndex = this->use64bitIndex;
+ SPIRVTypeConverter typeConverter(targetAttr, options);
populateMMAToSPIRVCoopMatrixTypeConversion(typeConverter,
this->useCoopMatrixNV);
More information about the cfe-commits
mailing list