[Mlir-commits] [mlir] 1e77536 - Revert "[mlir][ROCDL] Adds the ROCDL target attribute."

Fabian Mora llvmlistbot at llvm.org
Fri Aug 11 12:51:19 PDT 2023


Author: Fabian Mora
Date: 2023-08-11T19:50:05Z
New Revision: 1e77536e1d14c17075f3e484c189a8284c1c45e3

URL: https://github.com/llvm/llvm-project/commit/1e77536e1d14c17075f3e484c189a8284c1c45e3
DIFF: https://github.com/llvm/llvm-project/commit/1e77536e1d14c17075f3e484c189a8284c1c45e3.diff

LOG: Revert "[mlir][ROCDL] Adds the ROCDL target attribute."

This reverts commit 6a0feb1503e21432e63d93b44357bad43f8026d1.

Added: 
    

Modified: 
    mlir/include/mlir/Dialect/LLVMIR/CMakeLists.txt
    mlir/include/mlir/Dialect/LLVMIR/ROCDLDialect.h
    mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
    mlir/include/mlir/InitAllExtensions.h
    mlir/lib/Dialect/LLVMIR/IR/ROCDLDialect.cpp
    mlir/lib/Target/LLVM/CMakeLists.txt
    mlir/test/Dialect/GPU/ops.mlir
    mlir/test/Dialect/LLVMIR/rocdl.mlir
    mlir/unittests/Target/LLVM/CMakeLists.txt

Removed: 
    mlir/include/mlir/Target/LLVM/ROCDL/Target.h
    mlir/include/mlir/Target/LLVM/ROCDL/Utils.h
    mlir/lib/Target/LLVM/ROCDL/Target.cpp
    mlir/unittests/Target/LLVM/SerializeROCDLTarget.cpp


################################################################################
diff  --git a/mlir/include/mlir/Dialect/LLVMIR/CMakeLists.txt b/mlir/include/mlir/Dialect/LLVMIR/CMakeLists.txt
index aaf64e63321f20..2f65c1a3d6bcde 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/CMakeLists.txt
+++ b/mlir/include/mlir/Dialect/LLVMIR/CMakeLists.txt
@@ -62,6 +62,4 @@ add_mlir_dialect(ROCDLOps rocdl)
 add_mlir_doc(ROCDLOps ROCDLDialect Dialects/ -gen-dialect-doc -dialect=rocdl)
 set(LLVM_TARGET_DEFINITIONS ROCDLOps.td)
 mlir_tablegen(ROCDLConversions.inc -gen-llvmir-conversions)
-mlir_tablegen(ROCDLOpsAttributes.h.inc -gen-attrdef-decls -attrdefs-dialect=rocdl)
-mlir_tablegen(ROCDLOpsAttributes.cpp.inc -gen-attrdef-defs -attrdefs-dialect=rocdl)
 add_public_tablegen_target(MLIRROCDLConversionsIncGen)

diff  --git a/mlir/include/mlir/Dialect/LLVMIR/ROCDLDialect.h b/mlir/include/mlir/Dialect/LLVMIR/ROCDLDialect.h
index d4208a4e251631..6e1072423a0f5b 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/ROCDLDialect.h
+++ b/mlir/include/mlir/Dialect/LLVMIR/ROCDLDialect.h
@@ -23,16 +23,12 @@
 #define MLIR_DIALECT_LLVMIR_ROCDLDIALECT_H_
 
 #include "mlir/Bytecode/BytecodeOpInterface.h"
-#include "mlir/Dialect/GPU/IR/CompilationInterfaces.h"
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/IR/Dialect.h"
 #include "mlir/IR/OpDefinition.h"
 #include "mlir/Interfaces/SideEffectInterfaces.h"
 
 ///// Ops /////
-#define GET_ATTRDEF_CLASSES
-#include "mlir/Dialect/LLVMIR/ROCDLOpsAttributes.h.inc"
-
 #define GET_OP_CLASSES
 #include "mlir/Dialect/LLVMIR/ROCDLOps.h.inc"
 

diff  --git a/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td b/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
index 0738f09ac4164a..e9c349577dfbca 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
@@ -13,7 +13,6 @@
 #ifndef ROCDLIR_OPS
 #define ROCDLIR_OPS
 
-include "mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td"
 include "mlir/Dialect/LLVMIR/LLVMOpBase.td"
 include "mlir/Interfaces/SideEffectInterfaces.td"
 
@@ -45,20 +44,8 @@ def ROCDL_Dialect : Dialect {
     /// The address space value that represents private memory.
     static constexpr unsigned kPrivateMemoryAddressSpace = 5;
   }];
-
-  let useDefaultAttributePrinterParser = 1;
 }
 
-//===----------------------------------------------------------------------===//
-// ROCDL attribute definitions
-//===----------------------------------------------------------------------===//
-
-class ROCDL_Attr<string attrName, string attrMnemonic, list<Trait> traits = []>
-    : AttrDef<ROCDL_Dialect, attrName, traits> {
-  let mnemonic = attrMnemonic;
-}
-
-
 //===----------------------------------------------------------------------===//
 // ROCDL op definitions
 //===----------------------------------------------------------------------===//
@@ -399,92 +386,4 @@ def ROCDL_RawBufferAtomicUMinOp :
   let hasCustomAssemblyFormat = 1;
 }
 
-//===----------------------------------------------------------------------===//
-// ROCDL target attribute.
-//===----------------------------------------------------------------------===//
-
-def ROCDL_TargettAttr :
-    ROCDL_Attr<"ROCDLTarget", "target"> {
-  let description = [{
-    ROCDL target attribute for controlling compilation of AMDGPU targets. All
-    parameters decay into default values if not present.
-
-    Examples:
-
-    1. Target with default values.
-    ```
-      gpu.module @mymodule [#rocdl.target] attributes {...} {
-        ...
-      }
-    ```
-
-    2. Target with `gfx90a` chip and fast math.
-    ```
-      gpu.module @mymodule [#rocdl.target<chip = "gfx90a", flags = {fast, no_wave64}>] {
-        ...
-      }
-    ```
-  }];
-  let parameters = (ins
-    DefaultValuedParameter<"int", "2", "Optimization level to apply.">:$O,
-    StringRefParameter<"Target triple.", "\"amdgcn-amd-amdhsa\"">:$triple,
-    StringRefParameter<"Target chip.", "\"gfx900\"">:$chip,
-    StringRefParameter<"Target chip features.", "\"\"">:$features,
-    StringRefParameter<"ABI version.", "\"400\"">:$abi,
-    OptionalParameter<"DictionaryAttr", "Target specific flags.">:$flags,
-    OptionalParameter<"ArrayAttr", "Files to link to the LLVM module.">:$link
-  );
-  let assemblyFormat = [{
-    (`<` struct($O, $triple, $chip, $features, $abi, $flags, $link)^ `>`)?
-  }];
-  let builders = [
-    AttrBuilder<(ins CArg<"int", "2">:$optLevel,
-                     CArg<"StringRef", "\"amdgcn-amd-amdhsa\"">:$triple,
-                     CArg<"StringRef", "\"gfx900\"">:$chip,
-                     CArg<"StringRef", "\"\"">:$features,
-                     CArg<"StringRef", "\"400\"">:$abiVersion,
-                     CArg<"DictionaryAttr", "nullptr">:$targetFlags,
-                     CArg<"ArrayAttr", "nullptr">:$linkFiles), [{
-      return Base::get($_ctxt, optLevel, triple, chip, features, abiVersion,
-                       targetFlags, linkFiles);
-    }]>
-  ];
-  let skipDefaultBuilders = 1;
-  let genVerifyDecl = 1;
-  let extraClassDeclaration = [{
-    bool hasFlag(StringRef flag) const;
-    bool hasWave64() const;
-    bool hasFastMath() const;
-    bool hasDaz() const;
-    bool hasFiniteOnly() const;
-    bool hasUnsafeMath() const;
-    bool hasCorrectSqrt() const;
-  }];
-  let extraClassDefinition = [{
-    bool $cppClass::hasFlag(StringRef flag) const {
-      if (DictionaryAttr flags = getFlags())
-        return flags.get(flag) != nullptr;
-      return false;
-    }
-    bool $cppClass::hasWave64() const {
-      return hasFlag("wave64") || !hasFlag("no_wave64");
-    }
-    bool $cppClass::hasFastMath() const {
-      return hasFlag("fast");
-    }
-    bool $cppClass::hasDaz() const {
-      return hasFlag("daz");
-    }
-    bool $cppClass::hasFiniteOnly() const {
-      return hasFlag("finite_only");
-    }
-    bool $cppClass::hasUnsafeMath() const {
-      return hasFlag("unsafe_math");
-    }
-    bool $cppClass::hasCorrectSqrt() const {
-      return !hasFlag("unsafe_sqrt");
-    }
-  }];
-}
-
 #endif // ROCDLIR_OPS

diff  --git a/mlir/include/mlir/InitAllExtensions.h b/mlir/include/mlir/InitAllExtensions.h
index fed2e15c4efb92..ddcb27302ff657 100644
--- a/mlir/include/mlir/InitAllExtensions.h
+++ b/mlir/include/mlir/InitAllExtensions.h
@@ -25,7 +25,6 @@
 #include "mlir/Conversion/UBToLLVM/UBToLLVM.h"
 #include "mlir/Dialect/Func/Extensions/AllExtensions.h"
 #include "mlir/Target/LLVM/NVVM/Target.h"
-#include "mlir/Target/LLVM/ROCDL/Target.h"
 
 #include <cstdlib>
 
@@ -48,7 +47,6 @@ inline void registerAllExtensions(DialectRegistry &registry) {
   registerConvertNVVMToLLVMInterface(registry);
   ub::registerConvertUBToLLVMInterface(registry);
   registerNVVMTarget(registry);
-  registerROCDLTarget(registry);
 }
 
 } // namespace mlir

diff  --git a/mlir/include/mlir/Target/LLVM/ROCDL/Target.h b/mlir/include/mlir/Target/LLVM/ROCDL/Target.h
deleted file mode 100644
index f6a548e550868c..00000000000000
--- a/mlir/include/mlir/Target/LLVM/ROCDL/Target.h
+++ /dev/null
@@ -1,28 +0,0 @@
-//===- Target.h - MLIR ROCDL target registration ----------------*- C++ -*-===//
-//
-// 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 provides registration calls for attaching the ROCDL target interface.
-//
-//===----------------------------------------------------------------------===//
-
-#ifndef MLIR_TARGET_LLVM_ROCDL_TARGET_H
-#define MLIR_TARGET_LLVM_ROCDL_TARGET_H
-
-namespace mlir {
-class DialectRegistry;
-class MLIRContext;
-/// Registers the `TargetAttrInterface` for the `#rocdl.target` attribute in the
-/// given registry.
-void registerROCDLTarget(DialectRegistry &registry);
-
-/// Registers the `TargetAttrInterface` for the `#rocdl.target` attribute in the
-/// registry associated with the given context.
-void registerROCDLTarget(MLIRContext &context);
-} // namespace mlir
-
-#endif // MLIR_TARGET_LLVM_ROCDL_TARGET_H

diff  --git a/mlir/include/mlir/Target/LLVM/ROCDL/Utils.h b/mlir/include/mlir/Target/LLVM/ROCDL/Utils.h
deleted file mode 100644
index c14fa80056a879..00000000000000
--- a/mlir/include/mlir/Target/LLVM/ROCDL/Utils.h
+++ /dev/null
@@ -1,94 +0,0 @@
-//===- Utils.h - MLIR ROCDL target utils ------------------------*- C++ -*-===//
-//
-// 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 files declares ROCDL target related utility classes and functions.
-//
-//===----------------------------------------------------------------------===//
-
-#ifndef MLIR_TARGET_LLVM_ROCDL_UTILS_H
-#define MLIR_TARGET_LLVM_ROCDL_UTILS_H
-
-#include "mlir/Dialect/GPU/IR/CompilationInterfaces.h"
-#include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
-#include "mlir/Target/LLVM/ModuleToObject.h"
-
-namespace mlir {
-namespace ROCDL {
-/// Searches & returns the path ROCM toolkit path, the search order is:
-/// 1. The `ROCM_PATH` environment variable.
-/// 2. The `ROCM_ROOT` environment variable.
-/// 3. The `ROCM_HOME` environment variable.
-/// 4. The ROCM path detected by CMake.
-/// 5. Returns an empty string.
-StringRef getROCMPath();
-
-/// Base class for all ROCDL serializations from GPU modules into binary
-/// strings. By default this class serializes into LLVM bitcode.
-class SerializeGPUModuleBase : public LLVM::ModuleToObject {
-public:
-  /// Initializes the `toolkitPath` with the path in `targetOptions` or if empty
-  /// with the path in `getROCMPath`.
-  SerializeGPUModuleBase(Operation &module, ROCDLTargetAttr target,
-                         const gpu::TargetOptions &targetOptions = {});
-
-  /// Initializes the LLVM AMDGPU target by safely calling
-  /// `LLVMInitializeAMDGPU*` methods if available.
-  static void init();
-
-  /// Returns the target attribute.
-  ROCDLTargetAttr getTarget() const;
-
-  /// Returns the ROCM toolkit path.
-  StringRef getToolkitPath() const;
-
-  /// Returns the bitcode files to be loaded.
-  ArrayRef<std::string> getFileList() const;
-
-  /// Appends standard ROCm device libraries like `ocml.bc`, `ockl.bc`, etc.
-  LogicalResult appendStandardLibs();
-
-  /// Loads the bitcode files in `fileList`.
-  virtual std::optional<SmallVector<std::unique_ptr<llvm::Module>>>
-  loadBitcodeFiles(llvm::Module &module,
-                   llvm::TargetMachine &targetMachine) override;
-
-  /// Adds `oclc` control variables to the LLVM module.
-  void handleModulePreLink(llvm::Module &module,
-                           llvm::TargetMachine &targetMachine) override;
-
-  /// Removes unnecessary metadata from the loaded bitcode files.
-  LogicalResult handleBitcodeFile(llvm::Module &module,
-                                  llvm::TargetMachine &targetMachine) override;
-
-protected:
-  /// Appends the paths of common ROCm device libraries to `libs`.
-  LogicalResult getCommonBitcodeLibs(llvm::SmallVector<std::string> &libs,
-                                     SmallVector<char, 256> &libPath,
-                                     StringRef isaVersion);
-
-  /// Adds `oclc` control variables to the LLVM module.
-  void addControlVariables(llvm::Module &module, bool wave64, bool daz,
-                           bool finiteOnly, bool unsafeMath, bool fastMath,
-                           bool correctSqrt, StringRef abiVer);
-
-  /// Returns the assembled ISA.
-  std::optional<SmallVector<char, 0>> assembleIsa(StringRef isa);
-
-  /// ROCDL target attribute.
-  ROCDLTargetAttr target;
-
-  /// ROCM toolkit path.
-  std::string toolkitPath;
-
-  /// List of LLVM bitcode files to link to.
-  SmallVector<std::string> fileList;
-};
-} // namespace ROCDL
-} // namespace mlir
-
-#endif // MLIR_TARGET_LLVM_ROCDL_UTILS_H

diff  --git a/mlir/lib/Dialect/LLVMIR/IR/ROCDLDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/ROCDLDialect.cpp
index a6ebdf409768b5..1a607d92ac7419 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/ROCDLDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/ROCDLDialect.cpp
@@ -16,14 +16,11 @@
 
 #include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
 
-#include "mlir/Dialect/GPU/IR/GPUDialect.h"
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/BuiltinTypes.h"
-#include "mlir/IR/DialectImplementation.h"
 #include "mlir/IR/MLIRContext.h"
 #include "mlir/IR/Operation.h"
-#include "llvm/ADT/TypeSwitch.h"
 #include "llvm/AsmParser/Parser.h"
 #include "llvm/IR/Attributes.h"
 #include "llvm/IR/Function.h"
@@ -240,14 +237,8 @@ void ROCDLDialect::initialize() {
 #include "mlir/Dialect/LLVMIR/ROCDLOps.cpp.inc"
       >();
 
-  addAttributes<
-#define GET_ATTRDEF_LIST
-#include "mlir/Dialect/LLVMIR/ROCDLOpsAttributes.cpp.inc"
-      >();
-
   // Support unknown operations because not all ROCDL operations are registered.
   allowUnknownOperations();
-  declarePromisedInterface<gpu::TargetAttrInterface>();
 }
 
 LogicalResult ROCDLDialect::verifyOperationAttribute(Operation *op,
@@ -262,41 +253,5 @@ LogicalResult ROCDLDialect::verifyOperationAttribute(Operation *op,
   return success();
 }
 
-//===----------------------------------------------------------------------===//
-// ROCDL target attribute.
-//===----------------------------------------------------------------------===//
-LogicalResult
-ROCDLTargetAttr::verify(function_ref<InFlightDiagnostic()> emitError,
-                        int optLevel, StringRef triple, StringRef chip,
-                        StringRef features, StringRef abiVersion,
-                        DictionaryAttr flags, ArrayAttr files) {
-  if (optLevel < 0 || optLevel > 3) {
-    emitError() << "The optimization level must be a number between 0 and 3.";
-    return failure();
-  }
-  if (triple.empty()) {
-    emitError() << "The target triple cannot be empty.";
-    return failure();
-  }
-  if (chip.empty()) {
-    emitError() << "The target chip cannot be empty.";
-    return failure();
-  }
-  if (abiVersion != "400" && abiVersion != "500") {
-    emitError() << "Invalid ABI version, it must be either `400` or `500`.";
-    return failure();
-  }
-  if (files && !llvm::all_of(files, [](::mlir::Attribute attr) {
-        return attr && mlir::isa<StringAttr>(attr);
-      })) {
-    emitError() << "All the elements in the `link` array must be strings.";
-    return failure();
-  }
-  return success();
-}
-
 #define GET_OP_CLASSES
 #include "mlir/Dialect/LLVMIR/ROCDLOps.cpp.inc"
-
-#define GET_ATTRDEF_CLASSES
-#include "mlir/Dialect/LLVMIR/ROCDLOpsAttributes.cpp.inc"

diff  --git a/mlir/lib/Target/LLVM/CMakeLists.txt b/mlir/lib/Target/LLVM/CMakeLists.txt
index 842adefb7156b8..5f311d084e43f1 100644
--- a/mlir/lib/Target/LLVM/CMakeLists.txt
+++ b/mlir/lib/Target/LLVM/CMakeLists.txt
@@ -97,53 +97,3 @@ if(MLIR_ENABLE_CUDA_CONVERSIONS)
     __DEFAULT_CUDATOOLKIT_PATH__="${MLIR_CUDAToolkit_ROOT}"
   )
 endif()
-
-if (MLIR_ENABLE_ROCM_CONVERSIONS)
-  set(AMDGPU_LIBS
-    IRReader
-    IPO
-    linker
-    MCParser
-    AMDGPUAsmParser
-    AMDGPUCodeGen
-    AMDGPUDesc
-    AMDGPUInfo
-    target
-  )
-endif()
-
-add_mlir_dialect_library(MLIRROCDLTarget
-  ROCDL/Target.cpp
-
-  LINK_COMPONENTS
-  ${AMDGPU_LIBS}
-
-  LINK_LIBS PUBLIC
-  MLIRIR
-  MLIRExecutionEngineUtils
-  MLIRSupport
-  MLIRGPUDialect
-  MLIRTargetLLVM
-  MLIRROCDLToLLVMIRTranslation
-  )
-
-if(MLIR_ENABLE_ROCM_CONVERSIONS)
-  if (NOT ("AMDGPU" IN_LIST LLVM_TARGETS_TO_BUILD))
-    message(SEND_ERROR
-      "Building mlir with ROCm support requires the AMDGPU backend")
-  endif()
-
-  if (DEFINED ROCM_PATH)
-    set(DEFAULT_ROCM_PATH "${ROCM_PATH}" CACHE PATH "Fallback path to search for ROCm installs")
-  elseif(DEFINED ENV{ROCM_PATH})
-    set(DEFAULT_ROCM_PATH "$ENV{ROCM_PATH}" CACHE PATH "Fallback path to search for ROCm installs")
-  else()
-    set(DEFAULT_ROCM_PATH "/opt/rocm" CACHE PATH "Fallback path to search for ROCm installs")
-  endif()
-  message(VERBOSE "MLIR Default ROCM toolkit path: ${DEFAULT_ROCM_PATH}")
-
-  target_compile_definitions(obj.MLIRROCDLTarget
-    PRIVATE
-    __DEFAULT_ROCM_PATH__="${DEFAULT_ROCM_PATH}"
-  )
-endif()

diff  --git a/mlir/lib/Target/LLVM/ROCDL/Target.cpp b/mlir/lib/Target/LLVM/ROCDL/Target.cpp
deleted file mode 100644
index fbbf22eedef2d5..00000000000000
--- a/mlir/lib/Target/LLVM/ROCDL/Target.cpp
+++ /dev/null
@@ -1,464 +0,0 @@
-//===- Target.cpp - MLIR LLVM ROCDL target compilation ----------*- C++ -*-===//
-//
-// 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 files defines ROCDL target related functions including registration
-// calls for the `#rocdl.target` compilation attribute.
-//
-//===----------------------------------------------------------------------===//
-
-#include "mlir/Target/LLVM/ROCDL/Target.h"
-
-#include "mlir/Dialect/GPU/IR/GPUDialect.h"
-#include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
-#include "mlir/Support/FileUtilities.h"
-#include "mlir/Target/LLVM/ROCDL/Utils.h"
-#include "mlir/Target/LLVMIR/Dialect/GPU/GPUToLLVMIRTranslation.h"
-#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
-#include "mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h"
-#include "mlir/Target/LLVMIR/Export.h"
-
-#include "llvm/IR/Constants.h"
-#include "llvm/MC/MCAsmBackend.h"
-#include "llvm/MC/MCAsmInfo.h"
-#include "llvm/MC/MCCodeEmitter.h"
-#include "llvm/MC/MCContext.h"
-#include "llvm/MC/MCInstrInfo.h"
-#include "llvm/MC/MCObjectFileInfo.h"
-#include "llvm/MC/MCObjectWriter.h"
-#include "llvm/MC/MCParser/MCTargetAsmParser.h"
-#include "llvm/MC/MCRegisterInfo.h"
-#include "llvm/MC/MCStreamer.h"
-#include "llvm/MC/MCSubtargetInfo.h"
-#include "llvm/MC/TargetRegistry.h"
-#include "llvm/Support/FileSystem.h"
-#include "llvm/Support/FileUtilities.h"
-#include "llvm/Support/Path.h"
-#include "llvm/Support/Program.h"
-#include "llvm/Support/SourceMgr.h"
-#include "llvm/Support/TargetSelect.h"
-#include "llvm/TargetParser/TargetParser.h"
-
-#include <cstdlib>
-
-using namespace mlir;
-using namespace mlir::ROCDL;
-
-#ifndef __DEFAULT_ROCM_PATH__
-#define __DEFAULT_ROCM_PATH__ ""
-#endif
-
-namespace {
-// Implementation of the `TargetAttrInterface` model.
-class ROCDLTargetAttrImpl
-    : public gpu::TargetAttrInterface::FallbackModel<ROCDLTargetAttrImpl> {
-public:
-  std::optional<SmallVector<char, 0>>
-  serializeToObject(Attribute attribute, Operation *module,
-                    const gpu::TargetOptions &options) const;
-};
-} // namespace
-
-// Register the ROCDL dialect, the ROCDL translation and the target interface.
-void mlir::registerROCDLTarget(DialectRegistry &registry) {
-  registerROCDLDialectTranslation(registry);
-  registry.addExtension(+[](MLIRContext *ctx, ROCDL::ROCDLDialect *dialect) {
-    ROCDLTargetAttr::attachInterface<ROCDLTargetAttrImpl>(*ctx);
-  });
-}
-
-void mlir::registerROCDLTarget(MLIRContext &context) {
-  DialectRegistry registry;
-  registerROCDLTarget(registry);
-  context.appendDialectRegistry(registry);
-}
-
-// Search for the ROCM path.
-StringRef mlir::ROCDL::getROCMPath() {
-  if (const char *var = std::getenv("ROCM_PATH"))
-    return var;
-  if (const char *var = std::getenv("ROCM_ROOT"))
-    return var;
-  if (const char *var = std::getenv("ROCM_HOME"))
-    return var;
-  return __DEFAULT_ROCM_PATH__;
-}
-
-SerializeGPUModuleBase::SerializeGPUModuleBase(
-    Operation &module, ROCDLTargetAttr target,
-    const gpu::TargetOptions &targetOptions)
-    : ModuleToObject(module, target.getTriple(), target.getChip(),
-                     target.getFeatures(), target.getO()),
-      target(target), toolkitPath(targetOptions.getToolkitPath()),
-      fileList(targetOptions.getLinkFiles()) {
-
-  // If `targetOptions` has an empty toolkitPath use `getROCMPath`
-  if (toolkitPath.empty())
-    toolkitPath = getROCMPath();
-
-  // Append the files in the target attribute.
-  if (ArrayAttr files = target.getLink())
-    for (Attribute attr : files.getValue())
-      if (auto file = dyn_cast<StringAttr>(attr))
-        fileList.push_back(file.str());
-
-  // Append standard ROCm device bitcode libraries to the files to be loaded.
-  (void)appendStandardLibs();
-}
-
-void SerializeGPUModuleBase::init() {
-  static llvm::once_flag initializeBackendOnce;
-  llvm::call_once(initializeBackendOnce, []() {
-  // If the `AMDGPU` LLVM target was built, initialize it.
-#if MLIR_ROCM_CONVERSIONS_ENABLED == 1
-    LLVMInitializeAMDGPUTarget();
-    LLVMInitializeAMDGPUTargetInfo();
-    LLVMInitializeAMDGPUTargetMC();
-    LLVMInitializeAMDGPUAsmParser();
-    LLVMInitializeAMDGPUAsmPrinter();
-#endif
-  });
-}
-
-ROCDLTargetAttr SerializeGPUModuleBase::getTarget() const { return target; }
-
-StringRef SerializeGPUModuleBase::getToolkitPath() const { return toolkitPath; }
-
-ArrayRef<std::string> SerializeGPUModuleBase::getFileList() const {
-  return fileList;
-}
-
-LogicalResult SerializeGPUModuleBase::appendStandardLibs() {
-  StringRef pathRef = getToolkitPath();
-  if (pathRef.size()) {
-    SmallVector<char, 256> path;
-    path.insert(path.begin(), pathRef.begin(), pathRef.end());
-    llvm::sys::path::append(path, "amdgcn", "bitcode");
-    pathRef = StringRef(path.data(), path.size());
-    if (!llvm::sys::fs::is_directory(pathRef)) {
-      getOperation().emitRemark() << "ROCm amdgcn bitcode path: " << pathRef
-                                  << " does not exist or is not a directory.";
-      return failure();
-    }
-    StringRef isaVersion =
-        llvm::AMDGPU::getArchNameAMDGCN(llvm::AMDGPU::parseArchAMDGCN(chip));
-    isaVersion.consume_front("gfx");
-    return getCommonBitcodeLibs(fileList, path, isaVersion);
-  }
-  return success();
-}
-
-std::optional<SmallVector<std::unique_ptr<llvm::Module>>>
-SerializeGPUModuleBase::loadBitcodeFiles(llvm::Module &module,
-                                         llvm::TargetMachine &targetMachine) {
-  SmallVector<std::unique_ptr<llvm::Module>> bcFiles;
-  if (failed(loadBitcodeFilesFromList(module.getContext(), targetMachine,
-                                      fileList, bcFiles, true)))
-    return std::nullopt;
-  return bcFiles;
-}
-
-LogicalResult
-SerializeGPUModuleBase::handleBitcodeFile(llvm::Module &module,
-                                          llvm::TargetMachine &targetMachine) {
-  // Some ROCM builds don't strip this like they should
-  if (auto *openclVersion = module.getNamedMetadata("opencl.ocl.version"))
-    module.eraseNamedMetadata(openclVersion);
-  // Stop spamming us with clang version numbers
-  if (auto *ident = module.getNamedMetadata("llvm.ident"))
-    module.eraseNamedMetadata(ident);
-  return success();
-}
-
-void SerializeGPUModuleBase::handleModulePreLink(
-    llvm::Module &module, llvm::TargetMachine &targetMachine) {
-  addControlVariables(module, target.hasWave64(), target.hasDaz(),
-                      target.hasFiniteOnly(), target.hasUnsafeMath(),
-                      target.hasFastMath(), target.hasCorrectSqrt(),
-                      target.getAbi());
-}
-
-// Get the paths of ROCm device libraries.
-LogicalResult SerializeGPUModuleBase::getCommonBitcodeLibs(
-    llvm::SmallVector<std::string> &libs, SmallVector<char, 256> &libPath,
-    StringRef isaVersion) {
-  auto addLib = [&](StringRef path) -> bool {
-    if (!llvm::sys::fs::is_regular_file(path)) {
-      getOperation().emitRemark() << "Bitcode library path: " << path
-                                  << " does not exist or is not a file.\n";
-      return true;
-    }
-    libs.push_back(path.str());
-    return false;
-  };
-  auto getLibPath = [&libPath](Twine lib) {
-    auto baseSize = libPath.size();
-    llvm::sys::path::append(libPath, lib + ".bc");
-    std::string path(StringRef(libPath.data(), libPath.size()).str());
-    libPath.truncate(baseSize);
-    return path;
-  };
-
-  // Add ROCm device libraries. Fail if any of the libraries is not found.
-  if (addLib(getLibPath("ocml")) || addLib(getLibPath("ockl")) ||
-      addLib(getLibPath("hip")) || addLib(getLibPath("opencl")) ||
-      addLib(getLibPath("oclc_isa_version_" + isaVersion)))
-    return failure();
-  return success();
-}
-
-void SerializeGPUModuleBase::addControlVariables(
-    llvm::Module &module, bool wave64, bool daz, bool finiteOnly,
-    bool unsafeMath, bool fastMath, bool correctSqrt, StringRef abiVer) {
-  llvm::Type *i8Ty = llvm::Type::getInt8Ty(module.getContext());
-  auto addControlVariable = [i8Ty, &module](StringRef name, bool enable) {
-    llvm::GlobalVariable *controlVariable = new llvm::GlobalVariable(
-        module, i8Ty, true, llvm::GlobalValue::LinkageTypes::LinkOnceODRLinkage,
-        llvm::ConstantInt::get(i8Ty, enable), name, nullptr,
-        llvm::GlobalValue::ThreadLocalMode::NotThreadLocal, 4);
-    controlVariable->setVisibility(
-        llvm::GlobalValue::VisibilityTypes::ProtectedVisibility);
-    controlVariable->setAlignment(llvm::MaybeAlign(1));
-    controlVariable->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Local);
-  };
-  addControlVariable("__oclc_finite_only_opt", finiteOnly || fastMath);
-  addControlVariable("__oclc_unsafe_math_opt", unsafeMath || fastMath);
-  addControlVariable("__oclc_daz_opt", daz || fastMath);
-  addControlVariable("__oclc_correctly_rounded_sqrt32",
-                     correctSqrt && !fastMath);
-  addControlVariable("__oclc_wavefrontsize64", wave64);
-
-  llvm::Type *i32Ty = llvm::Type::getInt32Ty(module.getContext());
-  int abi = 400;
-  abiVer.getAsInteger(0, abi);
-  llvm::GlobalVariable *abiVersion = new llvm::GlobalVariable(
-      module, i32Ty, true, llvm::GlobalValue::LinkageTypes::LinkOnceODRLinkage,
-      llvm::ConstantInt::get(i32Ty, abi), "__oclc_ABI_version", nullptr,
-      llvm::GlobalValue::ThreadLocalMode::NotThreadLocal, 4);
-  abiVersion->setVisibility(
-      llvm::GlobalValue::VisibilityTypes::ProtectedVisibility);
-  abiVersion->setAlignment(llvm::MaybeAlign(4));
-  abiVersion->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Local);
-}
-
-std::optional<SmallVector<char, 0>>
-SerializeGPUModuleBase::assembleIsa(StringRef isa) {
-  auto loc = getOperation().getLoc();
-
-  StringRef targetTriple = this->triple;
-
-  SmallVector<char, 0> result;
-  llvm::raw_svector_ostream os(result);
-
-  llvm::Triple triple(llvm::Triple::normalize(targetTriple));
-  std::string error;
-  const llvm::Target *target =
-      llvm::TargetRegistry::lookupTarget(triple.normalize(), error);
-  if (!target) {
-    emitError(loc, Twine("failed to lookup target: ") + error);
-    return std::nullopt;
-  }
-
-  llvm::SourceMgr srcMgr;
-  srcMgr.AddNewSourceBuffer(llvm::MemoryBuffer::getMemBuffer(isa), SMLoc());
-
-  const llvm::MCTargetOptions mcOptions;
-  std::unique_ptr<llvm::MCRegisterInfo> mri(
-      target->createMCRegInfo(targetTriple));
-  std::unique_ptr<llvm::MCAsmInfo> mai(
-      target->createMCAsmInfo(*mri, targetTriple, mcOptions));
-  mai->setRelaxELFRelocations(true);
-  std::unique_ptr<llvm::MCSubtargetInfo> sti(
-      target->createMCSubtargetInfo(targetTriple, chip, features));
-
-  llvm::MCContext ctx(triple, mai.get(), mri.get(), sti.get(), &srcMgr,
-                      &mcOptions);
-  std::unique_ptr<llvm::MCObjectFileInfo> mofi(target->createMCObjectFileInfo(
-      ctx, /*PIC=*/false, /*LargeCodeModel=*/false));
-  ctx.setObjectFileInfo(mofi.get());
-
-  SmallString<128> cwd;
-  if (!llvm::sys::fs::current_path(cwd))
-    ctx.setCompilationDir(cwd);
-
-  std::unique_ptr<llvm::MCStreamer> mcStreamer;
-  std::unique_ptr<llvm::MCInstrInfo> mcii(target->createMCInstrInfo());
-
-  llvm::MCCodeEmitter *ce = target->createMCCodeEmitter(*mcii, ctx);
-  llvm::MCAsmBackend *mab = target->createMCAsmBackend(*sti, *mri, mcOptions);
-  mcStreamer.reset(target->createMCObjectStreamer(
-      triple, ctx, std::unique_ptr<llvm::MCAsmBackend>(mab),
-      mab->createObjectWriter(os), std::unique_ptr<llvm::MCCodeEmitter>(ce),
-      *sti, mcOptions.MCRelaxAll, mcOptions.MCIncrementalLinkerCompatible,
-      /*DWARFMustBeAtTheEnd*/ false));
-  mcStreamer->setUseAssemblerInfoForParsing(true);
-
-  std::unique_ptr<llvm::MCAsmParser> parser(
-      createMCAsmParser(srcMgr, ctx, *mcStreamer, *mai));
-  std::unique_ptr<llvm::MCTargetAsmParser> tap(
-      target->createMCAsmParser(*sti, *parser, *mcii, mcOptions));
-
-  if (!tap) {
-    emitError(loc, "assembler initialization error");
-    return {};
-  }
-
-  parser->setTargetParser(*tap);
-  parser->Run(false);
-
-  return result;
-}
-
-#ifdef MLIR_ROCM_CONVERSIONS_ENABLED
-namespace {
-class AMDGPUSerializer : public SerializeGPUModuleBase {
-public:
-  AMDGPUSerializer(Operation &module, ROCDLTargetAttr target,
-                   const gpu::TargetOptions &targetOptions);
-
-  gpu::GPUModuleOp getOperation();
-
-  // Compile to HSA.
-  std::optional<SmallVector<char, 0>>
-  compileToBinary(const std::string &serializedISA);
-
-  std::optional<SmallVector<char, 0>>
-  moduleToObject(llvm::Module &llvmModule,
-                 llvm::TargetMachine &targetMachine) override;
-
-private:
-  // Target options.
-  gpu::TargetOptions targetOptions;
-};
-} // namespace
-
-AMDGPUSerializer::AMDGPUSerializer(Operation &module, ROCDLTargetAttr target,
-                                   const gpu::TargetOptions &targetOptions)
-    : SerializeGPUModuleBase(module, target, targetOptions),
-      targetOptions(targetOptions) {}
-
-gpu::GPUModuleOp AMDGPUSerializer::getOperation() {
-  return dyn_cast<gpu::GPUModuleOp>(&SerializeGPUModuleBase::getOperation());
-}
-
-std::optional<SmallVector<char, 0>>
-AMDGPUSerializer::compileToBinary(const std::string &serializedISA) {
-  // Assemble the ISA.
-  std::optional<SmallVector<char, 0>> isaBinary = assembleIsa(serializedISA);
-
-  if (!isaBinary) {
-    getOperation().emitError() << "Failed during ISA assembling.";
-    return std::nullopt;
-  }
-
-  // Save the ISA binary to a temp file.
-  int tempIsaBinaryFd = -1;
-  SmallString<128> tempIsaBinaryFilename;
-  if (llvm::sys::fs::createTemporaryFile("kernel%%", "o", tempIsaBinaryFd,
-                                         tempIsaBinaryFilename)) {
-    getOperation().emitError()
-        << "Failed to create a temporary file for dumping the ISA binary.";
-    return std::nullopt;
-  }
-  llvm::FileRemover cleanupIsaBinary(tempIsaBinaryFilename);
-  {
-    llvm::raw_fd_ostream tempIsaBinaryOs(tempIsaBinaryFd, true);
-    tempIsaBinaryOs << StringRef(isaBinary->data(), isaBinary->size());
-    tempIsaBinaryOs.flush();
-  }
-
-  // Create a temp file for HSA code object.
-  int tempHsacoFD = -1;
-  SmallString<128> tempHsacoFilename;
-  if (llvm::sys::fs::createTemporaryFile("kernel", "hsaco", tempHsacoFD,
-                                         tempHsacoFilename)) {
-    getOperation().emitError()
-        << "Failed to create a temporary file for the HSA code object.";
-    return std::nullopt;
-  }
-  llvm::FileRemover cleanupHsaco(tempHsacoFilename);
-
-  llvm::SmallString<128> lldPath(toolkitPath);
-  llvm::sys::path::append(lldPath, "llvm", "bin", "ld.lld");
-  int lldResult = llvm::sys::ExecuteAndWait(
-      lldPath,
-      {"ld.lld", "-shared", tempIsaBinaryFilename, "-o", tempHsacoFilename});
-  if (lldResult != 0) {
-    getOperation().emitError() << "lld invocation failed.";
-    return std::nullopt;
-  }
-
-  // Load the HSA code object.
-  auto hsacoFile = openInputFile(tempHsacoFilename);
-  if (!hsacoFile) {
-    getOperation().emitError()
-        << "Failed to read the HSA code object from the temp file.";
-    return std::nullopt;
-  }
-
-  StringRef buffer = hsacoFile->getBuffer();
-
-  return SmallVector<char, 0>(buffer.begin(), buffer.end());
-}
-
-std::optional<SmallVector<char, 0>>
-AMDGPUSerializer::moduleToObject(llvm::Module &llvmModule,
-                                 llvm::TargetMachine &targetMachine) {
-  // Return LLVM IR if the compilation target is offload.
-#define DEBUG_TYPE "serialize-to-llvm"
-  LLVM_DEBUG({
-    llvm::dbgs() << "LLVM IR for module: " << getOperation().getNameAttr()
-                 << "\n"
-                 << llvmModule << "\n";
-  });
-#undef DEBUG_TYPE
-  if (targetOptions.getCompilationTarget() == gpu::TargetOptions::offload)
-    return SerializeGPUModuleBase::moduleToObject(llvmModule, targetMachine);
-
-  // Translate the Module to ISA.
-  std::optional<std::string> serializedISA =
-      translateToISA(llvmModule, targetMachine);
-  if (!serializedISA) {
-    getOperation().emitError() << "Failed translating the module to ISA.";
-    return std::nullopt;
-  }
-#define DEBUG_TYPE "serialize-to-isa"
-  LLVM_DEBUG({
-    llvm::dbgs() << "ISA for module: " << getOperation().getNameAttr() << "\n"
-                 << *serializedISA << "\n";
-  });
-#undef DEBUG_TYPE
-  // Return ISA assembly code if the compilation target is assembly.
-  if (targetOptions.getCompilationTarget() == gpu::TargetOptions::assembly)
-    return SmallVector<char, 0>(serializedISA->begin(), serializedISA->end());
-
-  // Compile to binary.
-  return compileToBinary(*serializedISA);
-}
-#endif // MLIR_ROCM_CONVERSIONS_ENABLED
-
-std::optional<SmallVector<char, 0>> ROCDLTargetAttrImpl::serializeToObject(
-    Attribute attribute, Operation *module,
-    const gpu::TargetOptions &options) const {
-  assert(module && "The module must be non null.");
-  if (!module)
-    return std::nullopt;
-  if (!mlir::isa<gpu::GPUModuleOp>(module)) {
-    module->emitError("Module must be a GPU module.");
-    return std::nullopt;
-  }
-#if MLIR_ROCM_CONVERSIONS_ENABLED == 1
-  AMDGPUSerializer serializer(*module, cast<ROCDLTargetAttr>(attribute),
-                              options);
-  serializer.init();
-  return serializer.run();
-#else
-  module->emitError("The `AMDGPU` target was not built. Please enable it when "
-                    "building LLVM.");
-  return std::nullopt;
-#endif // MLIR_ROCM_CONVERSIONS_ENABLED == 1
-}

diff  --git a/mlir/test/Dialect/GPU/ops.mlir b/mlir/test/Dialect/GPU/ops.mlir
index 7e1f64f69119f5..4ae39816693602 100644
--- a/mlir/test/Dialect/GPU/ops.mlir
+++ b/mlir/test/Dialect/GPU/ops.mlir
@@ -371,9 +371,3 @@ gpu.module @module_with_one_target [#nvvm.target] {
     gpu.return
   }
 }
-
-gpu.module @module_with_two_target [#nvvm.target, #rocdl.target<chip = "gfx90a">] {
-  gpu.func @kernel(%arg0 : f32) kernel {
-    gpu.return
-  }
-}

diff  --git a/mlir/test/Dialect/LLVMIR/rocdl.mlir b/mlir/test/Dialect/LLVMIR/rocdl.mlir
index e2a35e96fbc558..11605afc7d16a6 100644
--- a/mlir/test/Dialect/LLVMIR/rocdl.mlir
+++ b/mlir/test/Dialect/LLVMIR/rocdl.mlir
@@ -274,12 +274,3 @@ llvm.func @rocdl.raw.buffer.i32(%rsrc : vector<4xi32>,
 
 // expected-error at below {{attribute attached to unexpected op}}
 func.func private @expected_llvm_func() attributes { rocdl.kernel }
-
-// -----
-
-// Just check these don't emit errors.
-gpu.module @module_1 [#rocdl.target<O = 1, chip = "gfx900", abi = "500", link = ["my_device_lib.bc"], flags = {fast, daz, unsafe_math}>] {
-}
-
-gpu.module @module_2 [#rocdl.target<chip = "gfx900">, #rocdl.target<chip = "gfx90a">] {
-}

diff  --git a/mlir/unittests/Target/LLVM/CMakeLists.txt b/mlir/unittests/Target/LLVM/CMakeLists.txt
index 6d612548a94c0f..6b6e117df58f3d 100644
--- a/mlir/unittests/Target/LLVM/CMakeLists.txt
+++ b/mlir/unittests/Target/LLVM/CMakeLists.txt
@@ -1,6 +1,5 @@
 add_mlir_unittest(MLIRTargetLLVMTests
   SerializeNVVMTarget.cpp
-  SerializeROCDLTarget.cpp
   SerializeToLLVMBitcode.cpp
 )
 
@@ -10,14 +9,12 @@ target_link_libraries(MLIRTargetLLVMTests
   PRIVATE
   MLIRTargetLLVM
   MLIRNVVMTarget
-  MLIRROCDLTarget
   MLIRGPUDialect
   MLIRNVVMDialect
   MLIRLLVMDialect
   MLIRLLVMToLLVMIRTranslation
   MLIRBuiltinToLLVMIRTranslation
   MLIRNVVMToLLVMIRTranslation
-  MLIRROCDLToLLVMIRTranslation
   MLIRGPUToLLVMIRTranslation
   ${llvm_libs}
 )

diff  --git a/mlir/unittests/Target/LLVM/SerializeROCDLTarget.cpp b/mlir/unittests/Target/LLVM/SerializeROCDLTarget.cpp
deleted file mode 100644
index 9353c22bfa31f1..00000000000000
--- a/mlir/unittests/Target/LLVM/SerializeROCDLTarget.cpp
+++ /dev/null
@@ -1,158 +0,0 @@
-//===- SerializeROCDLTarget.cpp ---------------------------------*- C++ -*-===//
-//
-// This file is licensed 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
-//
-//===----------------------------------------------------------------------===//
-
-#include "mlir/Dialect/GPU/IR/GPUDialect.h"
-#include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
-#include "mlir/IR/MLIRContext.h"
-#include "mlir/InitAllDialects.h"
-#include "mlir/Parser/Parser.h"
-#include "mlir/Target/LLVM/ROCDL/Target.h"
-#include "mlir/Target/LLVM/ROCDL/Utils.h"
-#include "mlir/Target/LLVMIR/Dialect/Builtin/BuiltinToLLVMIRTranslation.h"
-#include "mlir/Target/LLVMIR/Dialect/GPU/GPUToLLVMIRTranslation.h"
-#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
-
-#include "llvm/IRReader/IRReader.h"
-#include "llvm/Support/FileSystem.h"
-#include "llvm/Support/MemoryBufferRef.h"
-#include "llvm/Support/Path.h"
-#include "llvm/Support/TargetSelect.h"
-#include "llvm/Support/raw_ostream.h"
-#include "llvm/TargetParser/Host.h"
-
-#include "gmock/gmock.h"
-
-using namespace mlir;
-
-// Skip the test if the AMDGPU target was not built.
-#if MLIR_ROCM_CONVERSIONS_ENABLED == 0
-#define SKIP_WITHOUT_AMDGPU(x) DISABLED_##x
-#else
-#define SKIP_WITHOUT_AMDGPU(x) x
-#endif
-
-class MLIRTargetLLVMROCDL : public ::testing::Test {
-protected:
-  virtual void SetUp() {
-    registerBuiltinDialectTranslation(registry);
-    registerLLVMDialectTranslation(registry);
-    registerGPUDialectTranslation(registry);
-    registerROCDLTarget(registry);
-  }
-
-  // Checks if a ROCm installation is available.
-  bool hasROCMTools() {
-    StringRef rocmPath = ROCDL::getROCMPath();
-    if (rocmPath.empty())
-      return false;
-    llvm::SmallString<128> lldPath(rocmPath);
-    llvm::sys::path::append(lldPath, "llvm", "bin", "ld.lld");
-    return llvm::sys::fs::can_execute(lldPath);
-  }
-
-  // Dialect registry.
-  DialectRegistry registry;
-
-  // MLIR module used for the tests.
-  const std::string moduleStr = R"mlir(
-      gpu.module @rocdl_test {
-        llvm.func @rocdl_kernel(%arg0: f32) attributes {gpu.kernel, rocdl.kernel} {
-        llvm.return
-      }
-    })mlir";
-};
-
-// Test ROCDL serialization to LLVM.
-TEST_F(MLIRTargetLLVMROCDL, SKIP_WITHOUT_AMDGPU(SerializeROCDLMToLLVM)) {
-  MLIRContext context(registry);
-
-  OwningOpRef<ModuleOp> module =
-      parseSourceString<ModuleOp>(moduleStr, &context);
-  ASSERT_TRUE(!!module);
-
-  // Create a ROCDL target.
-  ROCDL::ROCDLTargetAttr target = ROCDL::ROCDLTargetAttr::get(&context);
-
-  // Serialize the module.
-  auto serializer = dyn_cast<gpu::TargetAttrInterface>(target);
-  ASSERT_TRUE(!!serializer);
-  gpu::TargetOptions options("", {}, "", gpu::TargetOptions::offload);
-  for (auto gpuModule : (*module).getBody()->getOps<gpu::GPUModuleOp>()) {
-    std::optional<SmallVector<char, 0>> object =
-        serializer.serializeToObject(gpuModule, options);
-    // Check that the serializer was successful.
-    ASSERT_TRUE(object != std::nullopt);
-    ASSERT_TRUE(object->size() > 0);
-
-    // Read the serialized module.
-    llvm::MemoryBufferRef buffer(StringRef(object->data(), object->size()),
-                                 "module");
-    llvm::LLVMContext llvmContext;
-    llvm::Expected<std::unique_ptr<llvm::Module>> llvmModule =
-        llvm::getLazyBitcodeModule(buffer, llvmContext);
-    ASSERT_TRUE(!!llvmModule);
-    ASSERT_TRUE(!!*llvmModule);
-
-    // Check that it has a function named `foo`.
-    ASSERT_TRUE((*llvmModule)->getFunction("rocdl_kernel") != nullptr);
-  }
-}
-
-// Test ROCDL serialization to PTX.
-TEST_F(MLIRTargetLLVMROCDL, SKIP_WITHOUT_AMDGPU(SerializeROCDLToPTX)) {
-  MLIRContext context(registry);
-
-  OwningOpRef<ModuleOp> module =
-      parseSourceString<ModuleOp>(moduleStr, &context);
-  ASSERT_TRUE(!!module);
-
-  // Create a ROCDL target.
-  ROCDL::ROCDLTargetAttr target = ROCDL::ROCDLTargetAttr::get(&context);
-
-  // Serialize the module.
-  auto serializer = dyn_cast<gpu::TargetAttrInterface>(target);
-  ASSERT_TRUE(!!serializer);
-  gpu::TargetOptions options("", {}, "", gpu::TargetOptions::assembly);
-  for (auto gpuModule : (*module).getBody()->getOps<gpu::GPUModuleOp>()) {
-    std::optional<SmallVector<char, 0>> object =
-        serializer.serializeToObject(gpuModule, options);
-    // Check that the serializer was successful.
-    ASSERT_TRUE(object != std::nullopt);
-    ASSERT_TRUE(object->size() > 0);
-
-    ASSERT_TRUE(
-        StringRef(object->data(), object->size()).contains("rocdl_kernel"));
-  }
-}
-
-// Test ROCDL serialization to Binary.
-TEST_F(MLIRTargetLLVMROCDL, SKIP_WITHOUT_AMDGPU(SerializeROCDLToBinary)) {
-  if (!hasROCMTools())
-    GTEST_SKIP() << "ROCm installation not found, skipping test.";
-
-  MLIRContext context(registry);
-
-  OwningOpRef<ModuleOp> module =
-      parseSourceString<ModuleOp>(moduleStr, &context);
-  ASSERT_TRUE(!!module);
-
-  // Create a ROCDL target.
-  ROCDL::ROCDLTargetAttr target = ROCDL::ROCDLTargetAttr::get(&context);
-
-  // Serialize the module.
-  auto serializer = dyn_cast<gpu::TargetAttrInterface>(target);
-  ASSERT_TRUE(!!serializer);
-  gpu::TargetOptions options("", {}, "", gpu::TargetOptions::binary);
-  for (auto gpuModule : (*module).getBody()->getOps<gpu::GPUModuleOp>()) {
-    std::optional<SmallVector<char, 0>> object =
-        serializer.serializeToObject(gpuModule, options);
-    // Check that the serializer was successful.
-    ASSERT_TRUE(object != std::nullopt);
-    ASSERT_FALSE(object->empty());
-  }
-}


        


More information about the Mlir-commits mailing list