[llvm] [mlir] [mlir][bazel] Define MLIR_CUDA_CONVERSIONS_ENABLED in GPU targets. (PR #83004)
Ingo Müller via llvm-commits
llvm-commits at lists.llvm.org
Tue Feb 27 07:43:41 PST 2024
https://github.com/ingomueller-net updated https://github.com/llvm/llvm-project/pull/83004
>From 4c9cb9edb77a8b3bbdbea7129bf36e9c292eebbb Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Ingo=20M=C3=BCller?= <ingomueller at google.com>
Date: Mon, 26 Feb 2024 13:46:18 +0000
Subject: [PATCH] [mlir] Expose MLIR_CUDA_CONVERSIONS_ENABLED in mlir-config.h.
That macro was not defined in some cases and thus yielded warnings if
compiled with `-Wundef`. In particular, they were not defined in the
BUILD files, so the GPU targets were broken when built with Bazel. This
commit exposes mentioned CMake variable through mlir-config.h and uses
the macro that is introduced with the same name. This replaces the macro
MLIR_CUDA_CONVERSIONS_ENABLED, which the CMake files previously defined
manually.
---
mlir/CMakeLists.txt | 2 --
mlir/include/mlir/Config/mlir-config.h.cmake | 4 +++
mlir/include/mlir/InitAllPasses.h | 3 +-
.../GPU/Pipelines/GPUToNVVMPipeline.cpp | 5 +--
.../Dialect/GPU/Transforms/ModuleToBinary.cpp | 3 +-
mlir/lib/Target/LLVM/NVVM/Target.cpp | 9 ++---
.../Target/LLVM/SerializeNVVMTarget.cpp | 3 +-
.../llvm-project-overlay/mlir/BUILD.bazel | 10 ++++--
.../mlir/test/BUILD.bazel | 36 +++++++++----------
9 files changed, 43 insertions(+), 32 deletions(-)
diff --git a/mlir/CMakeLists.txt b/mlir/CMakeLists.txt
index 16c898bdeb6e00..070609c94a3b38 100644
--- a/mlir/CMakeLists.txt
+++ b/mlir/CMakeLists.txt
@@ -111,8 +111,6 @@ if ("NVPTX" IN_LIST LLVM_TARGETS_TO_BUILD)
else()
set(MLIR_ENABLE_CUDA_CONVERSIONS 0)
endif()
-# TODO: we should use a config.h file like LLVM does
-add_definitions(-DMLIR_CUDA_CONVERSIONS_ENABLED=${MLIR_ENABLE_CUDA_CONVERSIONS})
# Build the ROCm conversions and run according tests if the AMDGPU backend
# is available.
diff --git a/mlir/include/mlir/Config/mlir-config.h.cmake b/mlir/include/mlir/Config/mlir-config.h.cmake
index e152a36c0ce0cf..4a7d75e2266823 100644
--- a/mlir/include/mlir/Config/mlir-config.h.cmake
+++ b/mlir/include/mlir/Config/mlir-config.h.cmake
@@ -29,4 +29,8 @@
/* If set, enables PDL usage. */
#cmakedefine01 MLIR_ENABLE_PDL_IN_PATTERNMATCH
+/* If set, enables CUDA-related features in CUDA-related transforms, pipelines,
+ and targets. */
+#cmakedefine01 MLIR_ENABLE_CUDA_CONVERSIONS
+
#endif
diff --git a/mlir/include/mlir/InitAllPasses.h b/mlir/include/mlir/InitAllPasses.h
index e28921619fe582..5d90c197a6cced 100644
--- a/mlir/include/mlir/InitAllPasses.h
+++ b/mlir/include/mlir/InitAllPasses.h
@@ -14,6 +14,7 @@
#ifndef MLIR_INITALLPASSES_H_
#define MLIR_INITALLPASSES_H_
+#include "mlir/Config/mlir-config.h"
#include "mlir/Conversion/Passes.h"
#include "mlir/Dialect/AMDGPU/Transforms/Passes.h"
#include "mlir/Dialect/Affine/Passes.h"
@@ -96,7 +97,7 @@ inline void registerAllPasses() {
bufferization::registerBufferizationPipelines();
sparse_tensor::registerSparseTensorPipelines();
tosa::registerTosaToLinalgPipelines();
-#if MLIR_CUDA_CONVERSIONS_ENABLED
+#if MLIR_ENABLE_CUDA_CONVERSIONS
gpu::registerGPUToNVVMPipeline();
#endif
}
diff --git a/mlir/lib/Dialect/GPU/Pipelines/GPUToNVVMPipeline.cpp b/mlir/lib/Dialect/GPU/Pipelines/GPUToNVVMPipeline.cpp
index 935f0deaf9c8a6..db1974ddb3773b 100644
--- a/mlir/lib/Dialect/GPU/Pipelines/GPUToNVVMPipeline.cpp
+++ b/mlir/lib/Dialect/GPU/Pipelines/GPUToNVVMPipeline.cpp
@@ -11,6 +11,7 @@
//
//===----------------------------------------------------------------------===//
+#include "mlir/Config/mlir-config.h"
#include "mlir/Conversion/AffineToStandard/AffineToStandard.h"
#include "mlir/Conversion/ArithToLLVM/ArithToLLVM.h"
#include "mlir/Conversion/FuncToLLVM/ConvertFuncToLLVMPass.h"
@@ -38,7 +39,7 @@
using namespace mlir;
-#if MLIR_CUDA_CONVERSIONS_ENABLED
+#if MLIR_ENABLE_CUDA_CONVERSIONS
namespace {
//===----------------------------------------------------------------------===//
@@ -127,4 +128,4 @@ void mlir::gpu::registerGPUToNVVMPipeline() {
buildLowerToNVVMPassPipeline);
}
-#endif // MLIR_CUDA_CONVERSIONS_ENABLED
+#endif // MLIR_ENABLE_CUDA_CONVERSIONS
diff --git a/mlir/lib/Dialect/GPU/Transforms/ModuleToBinary.cpp b/mlir/lib/Dialect/GPU/Transforms/ModuleToBinary.cpp
index 0527073da85b69..f379ea8193923d 100644
--- a/mlir/lib/Dialect/GPU/Transforms/ModuleToBinary.cpp
+++ b/mlir/lib/Dialect/GPU/Transforms/ModuleToBinary.cpp
@@ -13,6 +13,7 @@
#include "mlir/Dialect/GPU/Transforms/Passes.h"
+#include "mlir/Config/mlir-config.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
@@ -48,7 +49,7 @@ void GpuModuleToBinaryPass::getDependentDialects(
// Register all GPU related translations.
registry.insert<gpu::GPUDialect>();
registry.insert<LLVM::LLVMDialect>();
-#if MLIR_CUDA_CONVERSIONS_ENABLED == 1
+#if MLIR_ENABLE_CUDA_CONVERSIONS
registry.insert<NVVM::NVVMDialect>();
#endif
#if MLIR_ROCM_CONVERSIONS_ENABLED == 1
diff --git a/mlir/lib/Target/LLVM/NVVM/Target.cpp b/mlir/lib/Target/LLVM/NVVM/Target.cpp
index 71b15a92782ed7..d5b6645631edd6 100644
--- a/mlir/lib/Target/LLVM/NVVM/Target.cpp
+++ b/mlir/lib/Target/LLVM/NVVM/Target.cpp
@@ -13,6 +13,7 @@
#include "mlir/Target/LLVM/NVVM/Target.h"
+#include "mlir/Config/mlir-config.h"
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
#include "mlir/Target/LLVM/NVVM/Utils.h"
@@ -156,7 +157,7 @@ SerializeGPUModuleBase::loadBitcodeFiles(llvm::Module &module) {
return std::move(bcFiles);
}
-#if MLIR_CUDA_CONVERSIONS_ENABLED == 1
+#if MLIR_ENABLE_CUDA_CONVERSIONS
namespace {
class NVPTXSerializer : public SerializeGPUModuleBase {
public:
@@ -562,7 +563,7 @@ NVPTXSerializer::moduleToObject(llvm::Module &llvmModule) {
return compileToBinary(*serializedISA);
#endif // MLIR_NVPTXCOMPILER_ENABLED == 1
}
-#endif // MLIR_CUDA_CONVERSIONS_ENABLED == 1
+#endif // MLIR_ENABLE_CUDA_CONVERSIONS
std::optional<SmallVector<char, 0>>
NVVMTargetAttrImpl::serializeToObject(Attribute attribute, Operation *module,
@@ -574,7 +575,7 @@ NVVMTargetAttrImpl::serializeToObject(Attribute attribute, Operation *module,
module->emitError("Module must be a GPU module.");
return std::nullopt;
}
-#if MLIR_CUDA_CONVERSIONS_ENABLED == 1
+#if MLIR_ENABLE_CUDA_CONVERSIONS
NVPTXSerializer serializer(*module, cast<NVVMTargetAttr>(attribute), options);
serializer.init();
return serializer.run();
@@ -582,7 +583,7 @@ NVVMTargetAttrImpl::serializeToObject(Attribute attribute, Operation *module,
module->emitError(
"The `NVPTX` target was not built. Please enable it when building LLVM.");
return std::nullopt;
-#endif // MLIR_CUDA_CONVERSIONS_ENABLED == 1
+#endif // MLIR_ENABLE_CUDA_CONVERSIONS
}
Attribute
diff --git a/mlir/unittests/Target/LLVM/SerializeNVVMTarget.cpp b/mlir/unittests/Target/LLVM/SerializeNVVMTarget.cpp
index 26bfbd5c11e81e..06ed53152ba1fe 100644
--- a/mlir/unittests/Target/LLVM/SerializeNVVMTarget.cpp
+++ b/mlir/unittests/Target/LLVM/SerializeNVVMTarget.cpp
@@ -6,6 +6,7 @@
//
//===----------------------------------------------------------------------===//
+#include "mlir/Config/mlir-config.h"
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
#include "mlir/IR/MLIRContext.h"
@@ -29,7 +30,7 @@
using namespace mlir;
// Skip the test if the NVPTX target was not built.
-#if MLIR_CUDA_CONVERSIONS_ENABLED == 0
+#if MLIR_ENABLE_CUDA_CONVERSIONS == 0
#define SKIP_WITHOUT_NVPTX(x) DISABLED_##x
#else
#define SKIP_WITHOUT_NVPTX(x) x
diff --git a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
index 59ee03d9a3213f..737c8e7765f6a6 100644
--- a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
+++ b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
@@ -6,7 +6,6 @@
# The MLIR "Multi-Level Intermediate Representation" Compiler Infrastructure
load("@bazel_skylib//rules:expand_template.bzl", "expand_template")
-load("@bazel_skylib//rules:write_file.bzl", "write_file")
load(
":build_defs.bzl",
"cc_headers_only",
@@ -36,7 +35,10 @@ expand_template(
"#cmakedefine01 MLIR_ENABLE_EXPENSIVE_PATTERN_API_CHECKS": "#define MLIR_ENABLE_EXPENSIVE_PATTERN_API_CHECKS 0",
"#cmakedefine MLIR_GREEDY_REWRITE_RANDOMIZER_SEED ${MLIR_GREEDY_REWRITE_RANDOMIZER_SEED}": "/* #undef MLIR_GREEDY_REWRITE_RANDOMIZER_SEED */",
"#cmakedefine01 MLIR_ENABLE_PDL_IN_PATTERNMATCH": "#define MLIR_ENABLE_PDL_IN_PATTERNMATCH 1",
- },
+ } | if_cuda_available(
+ {"#cmakedefine01 MLIR_ENABLE_CUDA_CONVERSIONS": "#define MLIR_ENABLE_CUDA_CONVERSIONS 1"},
+ {"#cmakedefine01 MLIR_ENABLE_CUDA_CONVERSIONS": "#define MLIR_ENABLE_CUDA_CONVERSIONS 0"},
+ ),
template = "include/mlir/Config/mlir-config.h.cmake",
)
@@ -5468,7 +5470,6 @@ cc_library(
srcs = ["lib/Dialect/GPU/Pipelines/GPUToNVVMPipeline.cpp"],
hdrs = ["include/mlir/Dialect/GPU/Pipelines/Passes.h"],
includes = ["include"],
- local_defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"],
deps = [
":AffineToStandard",
":ArithToLLVM",
@@ -5492,6 +5493,7 @@ cc_library(
":Transforms",
":VectorToLLVM",
":VectorToSCF",
+ ":config",
],
)
@@ -5541,6 +5543,7 @@ cc_library(
":Transforms",
":VCIXToLLVMIRTranslation",
":VectorDialect",
+ ":config",
"//llvm:Core",
"//llvm:MC",
"//llvm:Support",
@@ -6176,6 +6179,7 @@ cc_library(
":NVVMToLLVMIRTranslation",
":TargetLLVM",
":ToLLVMIRTranslation",
+ ":config",
"//llvm:NVPTXCodeGen",
"//llvm:Support",
],
diff --git a/utils/bazel/llvm-project-overlay/mlir/test/BUILD.bazel b/utils/bazel/llvm-project-overlay/mlir/test/BUILD.bazel
index 68d9b23fd5643e..367c5a63bbf107 100644
--- a/utils/bazel/llvm-project-overlay/mlir/test/BUILD.bazel
+++ b/utils/bazel/llvm-project-overlay/mlir/test/BUILD.bazel
@@ -552,7 +552,7 @@ cc_library(
cc_library(
name = "TestTransforms",
srcs = glob(["lib/Transforms/*.cpp"]),
- defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"],
+ defines = ["MLIR_ENABLE_CUDA_CONVERSIONS"],
includes = ["lib/Dialect/Test"],
deps = [
":TestDialect",
@@ -579,7 +579,7 @@ cc_library(
cc_library(
name = "TestFuncToLLVM",
srcs = glob(["lib/Conversion/FuncToLLVM/*.cpp"]),
- defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"],
+ defines = ["MLIR_ENABLE_CUDA_CONVERSIONS"],
includes = ["lib/Dialect/Test"],
deps = [
":TestDialect",
@@ -594,7 +594,7 @@ cc_library(
cc_library(
name = "TestOneToNTypeConversion",
srcs = glob(["lib/Conversion/OneToNTypeConversion/*.cpp"]),
- defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"],
+ defines = ["MLIR_ENABLE_CUDA_CONVERSIONS"],
includes = ["lib/Dialect/Test"],
deps = [
":TestDialect",
@@ -653,7 +653,7 @@ cc_library(
cc_library(
name = "TestDLTI",
srcs = glob(["lib/Dialect/DLTI/*.cpp"]),
- defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"],
+ defines = ["MLIR_ENABLE_CUDA_CONVERSIONS"],
includes = ["lib/Dialect/Test"],
deps = [
":TestDialect",
@@ -667,7 +667,7 @@ cc_library(
cc_library(
name = "TestGPU",
srcs = glob(["lib/Dialect/GPU/*.cpp"]),
- defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"] + if_cuda_available([
+ defines = ["MLIR_ENABLE_CUDA_CONVERSIONS"] + if_cuda_available([
"MLIR_GPU_TO_CUBIN_PASS_ENABLE",
]),
includes = ["lib/Dialect/Test"],
@@ -714,7 +714,7 @@ cc_library(
cc_library(
name = "TestLinalg",
srcs = glob(["lib/Dialect/Linalg/*.cpp"]),
- defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"],
+ defines = ["MLIR_ENABLE_CUDA_CONVERSIONS"],
includes = ["lib/Dialect/Test"],
deps = [
"//llvm:Support",
@@ -748,7 +748,7 @@ cc_library(
cc_library(
name = "TestLLVM",
srcs = glob(["lib/Dialect/LLVM/*.cpp"]),
- defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"],
+ defines = ["MLIR_ENABLE_CUDA_CONVERSIONS"],
includes = ["lib/Dialect/Test"],
deps = [
"//mlir:AffineToStandard",
@@ -773,7 +773,7 @@ cc_library(
cc_library(
name = "TestMath",
srcs = glob(["lib/Dialect/Math/*.cpp"]),
- defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"],
+ defines = ["MLIR_ENABLE_CUDA_CONVERSIONS"],
includes = ["lib/Dialect/Test"],
deps = [
"//mlir:ArithDialect",
@@ -790,7 +790,7 @@ cc_library(
cc_library(
name = "TestMathToVCIX",
srcs = glob(["lib/Conversion/MathToVCIX/*.cpp"]),
- defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"],
+ defines = ["MLIR_ENABLE_CUDA_CONVERSIONS"],
includes = ["lib/Dialect/Test"],
deps = [
"//mlir:ArithDialect",
@@ -807,7 +807,7 @@ cc_library(
cc_library(
name = "TestMemRef",
srcs = glob(["lib/Dialect/MemRef/*.cpp"]),
- defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"],
+ defines = ["MLIR_ENABLE_CUDA_CONVERSIONS"],
includes = ["lib/Dialect/Test"],
deps = [
":TestDialect",
@@ -847,7 +847,7 @@ cc_library(
cc_library(
name = "TestNVGPU",
srcs = glob(["lib/Dialect/NVGPU/*.cpp"]),
- defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"],
+ defines = ["MLIR_ENABLE_CUDA_CONVERSIONS"],
includes = ["lib/Dialect/Test"],
deps = [
"//mlir:AffineDialect",
@@ -871,7 +871,7 @@ cc_library(
cc_library(
name = "TestSCF",
srcs = glob(["lib/Dialect/SCF/*.cpp"]),
- defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"],
+ defines = ["MLIR_ENABLE_CUDA_CONVERSIONS"],
includes = ["lib/Dialect/Test"],
deps = [
"//llvm:Support",
@@ -891,7 +891,7 @@ cc_library(
cc_library(
name = "TestArith",
srcs = glob(["lib/Dialect/Arith/*.cpp"]),
- defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"],
+ defines = ["MLIR_ENABLE_CUDA_CONVERSIONS"],
includes = ["lib/Dialect/Test"],
deps = [
"//mlir:ArithDialect",
@@ -908,7 +908,7 @@ cc_library(
cc_library(
name = "TestArmSME",
srcs = glob(["lib/Dialect/ArmSME/*.cpp"]),
- defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"],
+ defines = ["MLIR_ENABLE_CUDA_CONVERSIONS"],
includes = ["lib/Dialect/Test"],
deps = [
"//mlir:ArithToArmSME",
@@ -927,7 +927,7 @@ cc_library(
cc_library(
name = "TestBufferization",
srcs = glob(["lib/Dialect/Bufferization/*.cpp"]),
- defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"],
+ defines = ["MLIR_ENABLE_CUDA_CONVERSIONS"],
includes = ["lib/Dialect/Test"],
deps = [
"//mlir:BufferizationDialect",
@@ -989,7 +989,7 @@ cc_library(
cc_library(
name = "TestFunc",
srcs = glob(["lib/Dialect/Func/*.cpp"]),
- defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"],
+ defines = ["MLIR_ENABLE_CUDA_CONVERSIONS"],
includes = ["lib/Dialect/Test"],
deps = [
":TestDialect",
@@ -1005,7 +1005,7 @@ cc_library(
cc_library(
name = "TestTensor",
srcs = glob(["lib/Dialect/Tensor/*.cpp"]),
- defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"],
+ defines = ["MLIR_ENABLE_CUDA_CONVERSIONS"],
includes = ["lib/Dialect/Test"],
deps = [
"//mlir:ArithDialect",
@@ -1022,7 +1022,7 @@ cc_library(
cc_library(
name = "TestVector",
srcs = glob(["lib/Dialect/Vector/*.cpp"]),
- defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"],
+ defines = ["MLIR_ENABLE_CUDA_CONVERSIONS"],
includes = ["lib/Dialect/Test"],
deps = [
"//mlir:AffineDialect",
More information about the llvm-commits
mailing list