[mlir] Replace MLIR_ENABLE_ROCM_CONVERSIONS with LLVM_HAS_AMDGPU_TARGET (#182652)

`LLVM_HAS_NVPTX_TARGET` is already defined in `llvm/Config/Targets.h`
and used to gate NVPTX-related code in MLIR. The same macro exists for
AMDGPU as `LLVM_HAS_AMDGPU_TARGET`, but MLIR defined its own
`MLIR_ENABLE_ROCM_CONVERSIONS` variable for this purpose. This PR
removes `MLIR_ENABLE_ROCM_CONVERSIONS` and replaces it with
`LLVM_HAS_AMDGPU_TARGET`, bringing parity with the NVPTX target.

---------

Co-authored-by: William Moses <gh@wsmoses.com>

mlir/CMakeLists.txt

@@ -132,14 +132,6 @@ option(MLIR_ENABLE_EXECUTION_ENGINE
        "Enable building the MLIR Execution Engine."
        ${MLIR_ENABLE_EXECUTION_ENGINE_default})
 
-# Build the ROCm conversions and run according tests if the AMDGPU backend
-# is available.
-if ("AMDGPU" IN_LIST LLVM_TARGETS_TO_BUILD)
-  set(MLIR_ENABLE_ROCM_CONVERSIONS 1)
-else()
-  set(MLIR_ENABLE_ROCM_CONVERSIONS 0)
-endif()
-
 # Build the XeVM conversions and run according tests if the SPIRV backend
 # is available.
 if ("SPIRV" IN_LIST LLVM_TARGETS_TO_BUILD)

mlir/include/mlir/Config/mlir-config.h.cmake

@@ -39,8 +39,4 @@
 /* If set, enables features that depend on the NVIDIA's PTX compiler. */
 #cmakedefine01 MLIR_ENABLE_NVPTXCOMPILER
 
-/* If set, enables ROCm-related features in ROCM-related transforms, pipelines,
-   and targets. */
-#cmakedefine01 MLIR_ENABLE_ROCM_CONVERSIONS
-
 #endif

mlir/lib/Conversion/GPUCommon/CMakeLists.txt

@@ -6,7 +6,7 @@ if ("NVPTX" IN_LIST LLVM_TARGETS_TO_BUILD)
   )
 endif()
 
-if (MLIR_ENABLE_ROCM_CONVERSIONS)
+if ("AMDGPU" IN_LIST LLVM_TARGETS_TO_BUILD)
   set(AMDGPU_LIBS
     AMDGPUCodeGen
     AMDGPUDesc

mlir/lib/Target/LLVM/CMakeLists.txt

@@ -159,7 +159,7 @@ else()
 endif()
 
 
-if (MLIR_ENABLE_ROCM_CONVERSIONS)
+if ("AMDGPU" IN_LIST LLVM_TARGETS_TO_BUILD)
   set(AMDGPU_LIBS
     AMDGPUAsmParser
     AMDGPUCodeGen
@@ -188,7 +188,7 @@ add_mlir_dialect_library(MLIRROCDLTarget
   MLIRROCDLToLLVMIRTranslation
   )
 
-if(MLIR_ENABLE_ROCM_CONVERSIONS)
+if("AMDGPU" IN_LIST LLVM_TARGETS_TO_BUILD)
   if (DEFINED ROCM_PATH)
     set(DEFAULT_ROCM_PATH "${ROCM_PATH}" CACHE PATH "Fallback path to search for ROCm installs")
   elseif(DEFINED ENV{ROCM_PATH})

mlir/lib/Target/LLVM/ROCDL/Target.cpp

@@ -19,6 +19,7 @@
 #include "mlir/Target/LLVM/ROCDL/Utils.h"
 #include "mlir/Target/LLVMIR/Export.h"
 
+#include "llvm/Config/Targets.h"
 #include "llvm/IR/Constants.h"
 #include "llvm/MC/MCAsmBackend.h"
 #include "llvm/MC/MCAsmInfo.h"
@@ -112,7 +113,7 @@ void SerializeGPUModuleBase::init() {
   static llvm::once_flag initializeBackendOnce;
   llvm::call_once(initializeBackendOnce, []() {
   // If the `AMDGPU` LLVM target was built, initialize it.
-#if MLIR_ENABLE_ROCM_CONVERSIONS
+#if LLVM_HAS_AMDGPU_TARGET
     LLVMInitializeAMDGPUTarget();
     LLVMInitializeAMDGPUTargetInfo();
     LLVMInitializeAMDGPUTargetMC();
@@ -446,7 +447,7 @@ FailureOr<SmallVector<char, 0>> SerializeGPUModuleBase::moduleToObjectImpl(
   return compileToBinary(*serializedISA);
 }
 
-#if MLIR_ENABLE_ROCM_CONVERSIONS
+#if LLVM_HAS_AMDGPU_TARGET
 namespace {
 class AMDGPUSerializer : public SerializeGPUModuleBase {
 public:
@@ -471,7 +472,7 @@ FailureOr<SmallVector<char, 0>>
 AMDGPUSerializer::moduleToObject(llvm::Module &llvmModule) {
   return moduleToObjectImpl(targetOptions, llvmModule);
 }
-#endif // MLIR_ENABLE_ROCM_CONVERSIONS
+#endif // LLVM_HAS_AMDGPU_TARGET
 
 std::optional<mlir::gpu::SerializedObject>
 ROCDLTargetAttrImpl::serializeToObject(
@@ -484,7 +485,7 @@ ROCDLTargetAttrImpl::serializeToObject(
     module->emitError("module must be a GPU module");
     return std::nullopt;
   }
-#if MLIR_ENABLE_ROCM_CONVERSIONS
+#if LLVM_HAS_AMDGPU_TARGET
   AMDGPUSerializer serializer(*module, cast<ROCDLTargetAttr>(attribute),
                               options);
   serializer.init();
@@ -496,7 +497,7 @@ ROCDLTargetAttrImpl::serializeToObject(
   module->emitError("the `AMDGPU` target was not built. Please enable it when "
                     "building LLVM");
   return std::nullopt;
-#endif // MLIR_ENABLE_ROCM_CONVERSIONS
+#endif // LLVM_HAS_AMDGPU_TARGET
 }
 
 Attribute

mlir/test/CMakeLists.txt

@@ -69,10 +69,10 @@ llvm_canonicalize_cmake_booleans(
   LLVM_BUILD_EXAMPLES
   LLVM_HAS_NVPTX_TARGET
   LLVM_INCLUDE_SPIRV_TOOLS_TESTS
+  LLVM_HAS_AMDGPU_TARGET
   MLIR_ENABLE_BINDINGS_PYTHON
   MLIR_ENABLE_PYTHON_STABLE_ABI
   MLIR_ENABLE_CUDA_RUNNER
-  MLIR_ENABLE_ROCM_CONVERSIONS
   MLIR_ENABLE_ROCM_RUNNER
   MLIR_ENABLE_SPIRV_CPU_RUNNER
   MLIR_ENABLE_VULKAN_RUNNER

mlir/test/lit.site.cfg.py.in

@@ -34,7 +34,7 @@ config.mlir_lib_dir = "@MLIR_LIB_DIR@"
 config.build_examples = @LLVM_BUILD_EXAMPLES@
 config.run_nvptx_tests = @LLVM_HAS_NVPTX_TARGET@
 config.enable_cuda_runner = @MLIR_ENABLE_CUDA_RUNNER@
-config.run_rocm_tests = @MLIR_ENABLE_ROCM_CONVERSIONS@
+config.run_rocm_tests = @LLVM_HAS_AMDGPU_TARGET@
 config.enable_rocm_runner = @MLIR_ENABLE_ROCM_RUNNER@
 config.gpu_compilation_format = "@MLIR_GPU_COMPILATION_TEST_FORMAT@"
 config.rocm_test_chipset = "@ROCM_TEST_CHIPSET@"

mlir/unittests/Target/LLVM/SerializeROCDLTarget.cpp

@@ -18,6 +18,7 @@
 #include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h"
 
+#include "llvm/Config/Targets.h"
 #include "llvm/IRReader/IRReader.h"
 #include "llvm/Support/FileSystem.h"
 #include "llvm/Support/MemoryBufferRef.h"
@@ -31,7 +32,7 @@
 using namespace mlir;
 
 // Skip the test if the AMDGPU target was not built.
-#if MLIR_ENABLE_ROCM_CONVERSIONS
+#if LLVM_HAS_AMDGPU_TARGET
 #define SKIP_WITHOUT_AMDGPU(x) x
 #else
 #define SKIP_WITHOUT_AMDGPU(x) DISABLED_##x