diff --git a/mlir/CMakeLists.txt b/mlir/CMakeLists.txt index 16c898bdeb6e0..070609c94a3b3 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 e152a36c0ce0c..4a7d75e226682 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 e28921619fe58..5d90c197a6cce 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 935f0deaf9c8a..db1974ddb3773 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 0527073da85b6..f379ea8193923 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(); registry.insert(); -#if MLIR_CUDA_CONVERSIONS_ENABLED == 1 +#if MLIR_ENABLE_CUDA_CONVERSIONS registry.insert(); #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 71b15a92782ed..d5b6645631edd 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> 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(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 26bfbd5c11e81..cea49356538f0 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,10 +30,10 @@ using namespace mlir; // Skip the test if the NVPTX target was not built. -#if MLIR_CUDA_CONVERSIONS_ENABLED == 0 -#define SKIP_WITHOUT_NVPTX(x) DISABLED_##x -#else +#if MLIR_ENABLE_CUDA_CONVERSIONS #define SKIP_WITHOUT_NVPTX(x) x +#else +#define SKIP_WITHOUT_NVPTX(x) DISABLED_##x #endif class MLIRTargetLLVMNVVM : public ::testing::Test { diff --git a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel index 59ee03d9a3213..f58808476cbbc 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", ], @@ -9131,6 +9135,7 @@ cc_library( ":VectorTransforms", ":X86VectorDialect", ":X86VectorTransforms", + ":config", ], ) diff --git a/utils/bazel/llvm-project-overlay/mlir/test/BUILD.bazel b/utils/bazel/llvm-project-overlay/mlir/test/BUILD.bazel index 68d9b23fd5643..583411aa60e55 100644 --- a/utils/bazel/llvm-project-overlay/mlir/test/BUILD.bazel +++ b/utils/bazel/llvm-project-overlay/mlir/test/BUILD.bazel @@ -552,7 +552,6 @@ cc_library( cc_library( name = "TestTransforms", srcs = glob(["lib/Transforms/*.cpp"]), - defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"], includes = ["lib/Dialect/Test"], deps = [ ":TestDialect", @@ -579,7 +578,6 @@ cc_library( cc_library( name = "TestFuncToLLVM", srcs = glob(["lib/Conversion/FuncToLLVM/*.cpp"]), - defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"], includes = ["lib/Dialect/Test"], deps = [ ":TestDialect", @@ -594,7 +592,6 @@ cc_library( cc_library( name = "TestOneToNTypeConversion", srcs = glob(["lib/Conversion/OneToNTypeConversion/*.cpp"]), - defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"], includes = ["lib/Dialect/Test"], deps = [ ":TestDialect", @@ -653,7 +650,6 @@ cc_library( cc_library( name = "TestDLTI", srcs = glob(["lib/Dialect/DLTI/*.cpp"]), - defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"], includes = ["lib/Dialect/Test"], deps = [ ":TestDialect", @@ -667,7 +663,7 @@ cc_library( cc_library( name = "TestGPU", srcs = glob(["lib/Dialect/GPU/*.cpp"]), - defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"] + if_cuda_available([ + defines = if_cuda_available([ "MLIR_GPU_TO_CUBIN_PASS_ENABLE", ]), includes = ["lib/Dialect/Test"], @@ -714,7 +710,6 @@ cc_library( cc_library( name = "TestLinalg", srcs = glob(["lib/Dialect/Linalg/*.cpp"]), - defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"], includes = ["lib/Dialect/Test"], deps = [ "//llvm:Support", @@ -748,7 +743,6 @@ cc_library( cc_library( name = "TestLLVM", srcs = glob(["lib/Dialect/LLVM/*.cpp"]), - defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"], includes = ["lib/Dialect/Test"], deps = [ "//mlir:AffineToStandard", @@ -773,7 +767,6 @@ cc_library( cc_library( name = "TestMath", srcs = glob(["lib/Dialect/Math/*.cpp"]), - defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"], includes = ["lib/Dialect/Test"], deps = [ "//mlir:ArithDialect", @@ -790,7 +783,6 @@ cc_library( cc_library( name = "TestMathToVCIX", srcs = glob(["lib/Conversion/MathToVCIX/*.cpp"]), - defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"], includes = ["lib/Dialect/Test"], deps = [ "//mlir:ArithDialect", @@ -807,7 +799,6 @@ cc_library( cc_library( name = "TestMemRef", srcs = glob(["lib/Dialect/MemRef/*.cpp"]), - defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"], includes = ["lib/Dialect/Test"], deps = [ ":TestDialect", @@ -847,7 +838,6 @@ cc_library( cc_library( name = "TestNVGPU", srcs = glob(["lib/Dialect/NVGPU/*.cpp"]), - defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"], includes = ["lib/Dialect/Test"], deps = [ "//mlir:AffineDialect", @@ -871,7 +861,6 @@ cc_library( cc_library( name = "TestSCF", srcs = glob(["lib/Dialect/SCF/*.cpp"]), - defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"], includes = ["lib/Dialect/Test"], deps = [ "//llvm:Support", @@ -891,7 +880,6 @@ cc_library( cc_library( name = "TestArith", srcs = glob(["lib/Dialect/Arith/*.cpp"]), - defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"], includes = ["lib/Dialect/Test"], deps = [ "//mlir:ArithDialect", @@ -908,7 +896,6 @@ cc_library( cc_library( name = "TestArmSME", srcs = glob(["lib/Dialect/ArmSME/*.cpp"]), - defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"], includes = ["lib/Dialect/Test"], deps = [ "//mlir:ArithToArmSME", @@ -927,7 +914,6 @@ cc_library( cc_library( name = "TestBufferization", srcs = glob(["lib/Dialect/Bufferization/*.cpp"]), - defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"], includes = ["lib/Dialect/Test"], deps = [ "//mlir:BufferizationDialect", @@ -989,7 +975,6 @@ cc_library( cc_library( name = "TestFunc", srcs = glob(["lib/Dialect/Func/*.cpp"]), - defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"], includes = ["lib/Dialect/Test"], deps = [ ":TestDialect", @@ -1005,7 +990,6 @@ cc_library( cc_library( name = "TestTensor", srcs = glob(["lib/Dialect/Tensor/*.cpp"]), - defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"], includes = ["lib/Dialect/Test"], deps = [ "//mlir:ArithDialect", @@ -1022,7 +1006,6 @@ cc_library( cc_library( name = "TestVector", srcs = glob(["lib/Dialect/Vector/*.cpp"]), - defines = ["MLIR_CUDA_CONVERSIONS_ENABLED"], includes = ["lib/Dialect/Test"], deps = [ "//mlir:AffineDialect",