From cbeb372e4624e5db4f332aaaec1bc015a68860b0 Mon Sep 17 00:00:00 2001 From: Vladislav Sytchenko Date: Mon, 12 Apr 2021 14:55:06 -0400 Subject: [PATCH] SWDEV-280473 - Remove HSAIL support from the ROCm backend In adition to removing the HSAIL logic from the ROCm backend, guard all of the HSAIL includes in the common layer behind the WITH_COMPILER_LIB define. This is to avoid including HSAIL headers when building with no support for it. In common logic replace the use of the aclType enum with the new Program::file_type_t enum. This is essentially a local copy of the HSAIL enum to avoid including any HSAIL headers. Change-Id: Ica0651d1b29dfccc255cc584eb82a5cb35e1b520 --- rocclr/device/device.cpp | 12 ++ rocclr/device/device.hpp | 8 ++ rocclr/device/devkernel.cpp | 6 + rocclr/device/devkernel.hpp | 4 + rocclr/device/devprogram.cpp | 93 +++++++------ rocclr/device/devprogram.hpp | 40 +++++- rocclr/device/pal/paldevice.cpp | 2 + rocclr/device/pal/paldevice.hpp | 4 + rocclr/device/pal/palprogram.hpp | 2 + rocclr/device/rocm/CMakeLists.txt | 1 - rocclr/device/rocm/rocdevice.cpp | 26 ---- rocclr/device/rocm/rocdevice.hpp | 4 - rocclr/device/rocm/rockernel.cpp | 113 ---------------- rocclr/device/rocm/rockernel.hpp | 1 - rocclr/device/rocm/rocprogram.cpp | 210 ----------------------------- rocclr/device/rocm/rocprogram.hpp | 3 - rocclr/device/rocm/rocsettings.hpp | 2 - 17 files changed, 126 insertions(+), 405 deletions(-) diff --git a/rocclr/device/device.cpp b/rocclr/device/device.cpp index 5b640cc47a..d997fd4264 100644 --- a/rocclr/device/device.cpp +++ b/rocclr/device/device.cpp @@ -56,9 +56,11 @@ extern void DeviceUnload(); #include "blowfish/oclcrypt.hpp" #endif +#if defined(WITH_COMPILER_LIB) #include "utils/bif_section_labels.hpp" #include "utils/libUtils.h" #include "spirv/spirvUtils.h" +#endif #include #include @@ -757,6 +759,7 @@ bool ClBinary::setElfTarget() { return true; } +#if defined(WITH_COMPILER_LIB) std::string ClBinary::getBIFSymbol(unsigned int symbolID) const { size_t nSymbols = 0; // Due to PRE & POST defines in bif_section_labels.hpp conflict with @@ -791,6 +794,7 @@ std::string ClBinary::getBIFSymbol(unsigned int symbolID) const { } return ""; } +#endif void ClBinary::init(amd::option::Options* optionsObj) { // option has higher priority than environment variable. @@ -1089,6 +1093,7 @@ bool ClBinary::loadCompileOptions(std::string& compileOptions) const { char* options = nullptr; size_t sz; compileOptions.clear(); +#if defined(WITH_COMPILER_LIB) if (elfIn_->getSymbol(amd::Elf::COMMENT, getBIFSymbol(symOpenclCompilerOptions).c_str(), &options, &sz)) { if (sz > 0) { @@ -1096,6 +1101,7 @@ bool ClBinary::loadCompileOptions(std::string& compileOptions) const { } return true; } +#endif return false; } @@ -1103,6 +1109,7 @@ bool ClBinary::loadLinkOptions(std::string& linkOptions) const { char* options = nullptr; size_t sz; linkOptions.clear(); +#if defined(WITH_COMPILER_LIB) if (elfIn_->getSymbol(amd::Elf::COMMENT, getBIFSymbol(symOpenclLinkerOptions).c_str(), &options, &sz)) { if (sz > 0) { @@ -1110,17 +1117,22 @@ bool ClBinary::loadLinkOptions(std::string& linkOptions) const { } return true; } +#endif return false; } void ClBinary::storeCompileOptions(const std::string& compileOptions) { +#if defined(WITH_COMPILER_LIB) elfOut()->addSymbol(amd::Elf::COMMENT, getBIFSymbol(symOpenclCompilerOptions).c_str(), compileOptions.c_str(), compileOptions.length()); +#endif } void ClBinary::storeLinkOptions(const std::string& linkOptions) { +#if defined(WITH_COMPILER_LIB) elfOut()->addSymbol(amd::Elf::COMMENT, getBIFSymbol(symOpenclLinkerOptions).c_str(), linkOptions.c_str(), linkOptions.length()); +#endif } bool ClBinary::isSPIR() const { diff --git a/rocclr/device/device.hpp b/rocclr/device/device.hpp index ce3beb8bd5..08ba24d693 100644 --- a/rocclr/device/device.hpp +++ b/rocclr/device/device.hpp @@ -34,7 +34,9 @@ #include "devprogram.hpp" #include "devkernel.hpp" #include "amdocl/cl_profile_amd.h" +#if defined(WITH_COMPILER_LIB) #include "acl.h" +#endif #include "hwdebug.hpp" #include "devsignal.hpp" @@ -1030,9 +1032,11 @@ class ClBinary : public amd::HeapObject { //! Returns TRUE if binary file was allocated bool isBinaryAllocated() const { return (flags_ & BinaryAllocated) ? true : false; } +#if defined(WITH_COMPILER_LIB) //! Returns BIF symbol name by symbolID, //! returns empty string if not found or if BIF version is unsupported std::string getBIFSymbol(unsigned int symbolID) const; +#endif protected: const amd::Device& dev_; //!< Device object @@ -1445,7 +1449,9 @@ class Isa { */ class Device : public RuntimeObject { protected: +#if defined(WITH_COMPILER_LIB) typedef aclCompiler Compiler; +#endif public: // The structures below for MGPU launch match the device library format @@ -1498,8 +1504,10 @@ class Device : public RuntimeObject { ); }; +#if defined(WITH_COMPILER_LIB) virtual Compiler* compiler() const = 0; virtual Compiler* binCompiler() const { return compiler(); } +#endif Device(); virtual ~Device(); diff --git a/rocclr/device/devkernel.cpp b/rocclr/device/devkernel.cpp index 47f94c6246..965c268801 100644 --- a/rocclr/device/devkernel.cpp +++ b/rocclr/device/devkernel.cpp @@ -24,15 +24,19 @@ #include "devkernel.hpp" #include "utils/macros.hpp" #include "utils/options.hpp" +#if defined(WITH_COMPILER_LIB) #include "utils/bif_section_labels.hpp" #include "utils/libUtils.h" +#endif #include "comgrctx.hpp" #include #include #include +#if defined(WITH_COMPILER_LIB) #include "acl.h" +#endif namespace device { @@ -660,11 +664,13 @@ bool Kernel::createSignature( Kernel::~Kernel() { delete signature_; } // ================================================================================================ +#if defined(WITH_COMPILER_LIB) std::string Kernel::openclMangledName(const std::string& name) { const oclBIFSymbolStruct* bifSym = findBIF30SymStruct(symOpenclKernel); assert(bifSym && "symbol not found"); return std::string("&") + bifSym->str[bif::PRE] + name + bifSym->str[bif::POST]; } +#endif // ================================================================================================ void Kernel::FindLocalWorkSize(size_t workDim, const amd::NDRange& gblWorkSize, diff --git a/rocclr/device/devkernel.hpp b/rocclr/device/devkernel.hpp index ad97d9a2f7..155b219a69 100644 --- a/rocclr/device/devkernel.hpp +++ b/rocclr/device/devkernel.hpp @@ -20,7 +20,9 @@ #pragma once +#if defined(WITH_COMPILER_LIB) #include "aclTypes.h" +#endif #include "platform/context.hpp" #include "platform/object.hpp" #include "platform/memory.hpp" @@ -421,7 +423,9 @@ class Kernel : public amd::HeapObject { //! Return the build log const std::string& buildLog() const { return buildLog_; } +#if defined(WITH_COMPILER_LIB) static std::string openclMangledName(const std::string& name); +#endif const std::unordered_map& patch() const { return patchReferences_; } diff --git a/rocclr/device/devprogram.cpp b/rocclr/device/devprogram.cpp index 865b0a1718..617e594312 100644 --- a/rocclr/device/devprogram.cpp +++ b/rocclr/device/devprogram.cpp @@ -25,8 +25,10 @@ #include "devkernel.hpp" #include "utils/macros.hpp" #include "utils/options.hpp" +#if defined(WITH_COMPILER_LIB) #include "utils/bif_section_labels.hpp" #include "utils/libUtils.h" +#endif #include "comgrctx.hpp" #include @@ -43,8 +45,10 @@ #include #endif // defined(ATI_OS_LINUX) +#if defined(WITH_COMPILER_LIB) #include "spirv/spirvUtils.h" #include "acl.h" +#endif #ifdef EARLY_INLINE #define AMDGPU_EARLY_INLINE_ALL_OPTION " -mllvm -amdgpu-early-inline-all" @@ -78,19 +82,23 @@ Program::Program(amd::Device& device, amd::Program& owner) elfSectionType_(amd::Elf::LLVMIR), compileOptions_(), linkOptions_(), +#if defined(WITH_COMPILER_LIB) binaryElf_(nullptr), +#endif lastBuildOptionsArg_(), buildStatus_(CL_BUILD_NONE), buildError_(CL_SUCCESS), globalVariableTotalSize_(0), programOptions_(nullptr) { +#if defined(WITH_COMPILER_LIB) memset(&binOpts_, 0, sizeof(binOpts_)); binOpts_.struct_size = sizeof(binOpts_); binOpts_.elfclass = LP64_SWITCH(ELFCLASS32, ELFCLASS64); binOpts_.bitness = ELFDATA2LSB; binOpts_.alloc = &::malloc; binOpts_.dealloc = &::free; +#endif } // ================================================================================================ @@ -1103,7 +1111,7 @@ static void dumpCodeObject(const std::string& image) { // ================================================================================================ bool Program::linkImplLC(amd::option::Options* options) { #if defined(USE_COMGR_LIBRARY) - aclType continueCompileFrom = ACL_TYPE_LLVMIR_BINARY; + file_type_t continueCompileFrom = FILE_TYPE_LLVMIR_BINARY; internal_ = (compileOptions_.find("-cl-internal-kernel") != std::string::npos) ? true : false; @@ -1120,11 +1128,11 @@ bool Program::linkImplLC(amd::option::Options* options) { } switch (continueCompileFrom) { - case ACL_TYPE_CG: - case ACL_TYPE_LLVMIR_BINARY: { + case FILE_TYPE_CG: + case FILE_TYPE_LLVMIR_BINARY: { break; } - case ACL_TYPE_ASM_TEXT: { + case FILE_TYPE_ASM_TEXT: { char* section; size_t sz; clBinary()->elfOut()->getSection(amd::Elf::SOURCE, §ion, &sz); @@ -1139,7 +1147,7 @@ bool Program::linkImplLC(amd::option::Options* options) { bLinkLLVMBitcode = false; break; } - case ACL_TYPE_ISA: { + case FILE_TYPE_ISA: { amd::Comgr::destroy_data_set(inputs); binary_t isaBinary = binary(); finfo_t isaFdesc = BinaryFd(); @@ -1262,7 +1270,7 @@ bool Program::linkImplLC(amd::option::Options* options) { amd::Comgr::destroy_data_set(inputs); if (!ret) { - if (continueCompileFrom == ACL_TYPE_ASM_TEXT) { + if (continueCompileFrom == FILE_TYPE_ASM_TEXT) { buildLog_ += "Error: Creating the executable from ISA assembly text failed.\n"; } else { buildLog_ += "Error: Creating the executable from LLVM IRs failed.\n"; @@ -2100,9 +2108,10 @@ bool Program::setBinary(const char* binaryIn, size_t size, const device::Program } // ================================================================================================ -aclType Program::getCompilationStagesFromBinary(std::vector& completeStages, +Program::file_type_t Program::getCompilationStagesFromBinary( + std::vector& completeStages, bool& needOptionsCheck) { - aclType from = ACL_TYPE_DEFAULT; + Program::file_type_t from = FILE_TYPE_DEFAULT; if (isLC()) { #if defined(USE_COMGR_LIBRARY) completeStages.clear(); @@ -2115,30 +2124,30 @@ aclType Program::getCompilationStagesFromBinary(std::vector& completeSt if (containsLlvmirText && containsOpts) { completeStages.push_back(from); - from = ACL_TYPE_LLVMIR_BINARY; + from = FILE_TYPE_LLVMIR_BINARY; } if (containsShaderIsa) { completeStages.push_back(from); - from = ACL_TYPE_ISA; + from = FILE_TYPE_ISA; } std::string sCurOptions = compileOptions_ + linkOptions_; amd::option::Options curOptions; if (!amd::option::parseAllOptions(sCurOptions, curOptions, false, isLC())) { buildLog_ += curOptions.optionsLog(); LogError("Parsing compile options failed."); - return ACL_TYPE_DEFAULT; + return FILE_TYPE_DEFAULT; } switch (from) { - case ACL_TYPE_CG: - case ACL_TYPE_ISA: + case FILE_TYPE_CG: + case FILE_TYPE_ISA: // do not check options, if LLVMIR is absent or might be absent or options are absent if (!curOptions.oVariables->BinLLVMIR || !containsLlvmirText || !containsOpts) { needOptionsCheck = false; } break; // recompilation might be needed - case ACL_TYPE_LLVMIR_BINARY: - case ACL_TYPE_DEFAULT: + case FILE_TYPE_LLVMIR_BINARY: + case FILE_TYPE_DEFAULT: default: break; } @@ -2159,7 +2168,7 @@ aclType Program::getCompilationStagesFromBinary(std::vector& completeSt } if (containsSpirv) { completeStages.push_back(from); - from = ACL_TYPE_SPIRV_BINARY; + from = FILE_TYPE_SPIRV_BINARY; } bool containsSpirText = true; errorCode = aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_SPIR, nullptr, @@ -2169,7 +2178,7 @@ aclType Program::getCompilationStagesFromBinary(std::vector& completeSt } if (containsSpirText) { completeStages.push_back(from); - from = ACL_TYPE_SPIR_BINARY; + from = FILE_TYPE_SPIR_BINARY; } bool containsLlvmirText = true; errorCode = aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_LLVMIR, nullptr, @@ -2186,7 +2195,7 @@ aclType Program::getCompilationStagesFromBinary(std::vector& completeSt } if (containsLlvmirText && containsOpts) { completeStages.push_back(from); - from = ACL_TYPE_LLVMIR_BINARY; + from = FILE_TYPE_LLVMIR_BINARY; } // Checking HSAIL in .cg section bool containsHsailText = true; @@ -2204,11 +2213,11 @@ aclType Program::getCompilationStagesFromBinary(std::vector& completeSt } if (containsBrig) { completeStages.push_back(from); - from = ACL_TYPE_HSAIL_BINARY; + from = FILE_TYPE_HSAIL_BINARY; } else if (containsHsailText) { completeStages.push_back(from); - from = ACL_TYPE_HSAIL_TEXT; + from = FILE_TYPE_HSAIL_TEXT; } // Checking Loader Map symbol from CG section bool containsLoaderMap = true; @@ -2219,7 +2228,7 @@ aclType Program::getCompilationStagesFromBinary(std::vector& completeSt } if (containsLoaderMap) { completeStages.push_back(from); - from = ACL_TYPE_CG; + from = FILE_TYPE_CG; } // Checking ISA in .text section bool containsShaderIsa = true; @@ -2230,28 +2239,28 @@ aclType Program::getCompilationStagesFromBinary(std::vector& completeSt } if (containsShaderIsa) { completeStages.push_back(from); - from = ACL_TYPE_ISA; + from = FILE_TYPE_ISA; } std::string sCurOptions = compileOptions_ + linkOptions_; amd::option::Options curOptions; if (!amd::option::parseAllOptions(sCurOptions, curOptions, false, isLC())) { buildLog_ += curOptions.optionsLog(); LogError("Parsing compile options failed."); - return ACL_TYPE_DEFAULT; + return FILE_TYPE_DEFAULT; } switch (from) { // compile from HSAIL text, no matter prev. stages and options - case ACL_TYPE_HSAIL_TEXT: + case FILE_TYPE_HSAIL_TEXT: needOptionsCheck = false; break; - case ACL_TYPE_HSAIL_BINARY: + case FILE_TYPE_HSAIL_BINARY: // do not check options, if LLVMIR is absent or might be absent or options are absent if (!curOptions.oVariables->BinLLVMIR || !containsLlvmirText || !containsOpts) { needOptionsCheck = false; } break; - case ACL_TYPE_CG: - case ACL_TYPE_ISA: + case FILE_TYPE_CG: + case FILE_TYPE_ISA: // do not check options, if LLVMIR is absent or might be absent or options are absent if (!curOptions.oVariables->BinLLVMIR || !containsLlvmirText || !containsOpts) { needOptionsCheck = false; @@ -2262,8 +2271,8 @@ aclType Program::getCompilationStagesFromBinary(std::vector& completeSt } break; // recompilation might be needed - case ACL_TYPE_LLVMIR_BINARY: - case ACL_TYPE_DEFAULT: + case FILE_TYPE_LLVMIR_BINARY: + case FILE_TYPE_DEFAULT: default: break; } @@ -2273,8 +2282,8 @@ aclType Program::getCompilationStagesFromBinary(std::vector& completeSt } // ================================================================================================ -aclType Program::getNextCompilationStageFromBinary(amd::option::Options* options) { - aclType continueCompileFrom = ACL_TYPE_DEFAULT; +Program::file_type_t Program::getNextCompilationStageFromBinary(amd::option::Options* options) { + Program::file_type_t continueCompileFrom = FILE_TYPE_DEFAULT; binary_t binary = this->binary(); finfo_t finfo = this->BinaryFd(); std::string uri = this->BinaryURI(); @@ -2303,7 +2312,7 @@ aclType Program::getNextCompilationStageFromBinary(amd::option::Options* options // Calculate the next stage to compile from, based on sections in binaryElf_; // No any validity checks here - std::vector completeStages; + std::vector completeStages; bool needOptionsCheck = true; continueCompileFrom = getCompilationStagesFromBinary(completeStages, needOptionsCheck); if (!options || !needOptionsCheck) { @@ -2312,9 +2321,9 @@ aclType Program::getNextCompilationStageFromBinary(amd::option::Options* options bool recompile = false; //! @todo Should we also check for ACL_TYPE_OPENCL & ACL_TYPE_LLVMIR_TEXT? switch (continueCompileFrom) { - case ACL_TYPE_HSAIL_BINARY: - case ACL_TYPE_CG: - case ACL_TYPE_ISA: { + case FILE_TYPE_HSAIL_BINARY: + case FILE_TYPE_CG: + case FILE_TYPE_ISA: { // Compare options loaded from binary with current ones, recompile if differ; // If compile options are absent in binary, do not compare and recompile if (compileOptions_.empty()) break; @@ -2350,12 +2359,12 @@ aclType Program::getNextCompilationStageFromBinary(amd::option::Options* options if (!amd::option::parseAllOptions(sBinOptions, binOptions, false, isLC())) { buildLog_ += binOptions.optionsLog(); LogError("Parsing compile options from binary failed."); - return ACL_TYPE_DEFAULT; + return FILE_TYPE_DEFAULT; } if (!amd::option::parseAllOptions(sCurOptions, curOptions, false, isLC())) { buildLog_ += curOptions.optionsLog(); LogError("Parsing compile options failed."); - return ACL_TYPE_DEFAULT; + return FILE_TYPE_DEFAULT; } if (!curOptions.equals(binOptions)) { recompile = true; @@ -2368,10 +2377,10 @@ aclType Program::getNextCompilationStageFromBinary(amd::option::Options* options if (recompile) { while (!completeStages.empty()) { continueCompileFrom = completeStages.back(); - if (continueCompileFrom == ACL_TYPE_SPIRV_BINARY || - continueCompileFrom == ACL_TYPE_LLVMIR_BINARY || - continueCompileFrom == ACL_TYPE_SPIR_BINARY || - continueCompileFrom == ACL_TYPE_DEFAULT) { + if (continueCompileFrom == FILE_TYPE_SPIRV_BINARY || + continueCompileFrom == FILE_TYPE_LLVMIR_BINARY || + continueCompileFrom == FILE_TYPE_SPIR_BINARY || + continueCompileFrom == FILE_TYPE_DEFAULT) { break; } completeStages.pop_back(); @@ -2381,7 +2390,7 @@ aclType Program::getNextCompilationStageFromBinary(amd::option::Options* options else { const char* xLang = options->oVariables->XLang; if (xLang != nullptr && strcmp(xLang, "asm") == 0) { - continueCompileFrom = ACL_TYPE_ASM_TEXT; + continueCompileFrom = FILE_TYPE_ASM_TEXT; } } return continueCompileFrom; diff --git a/rocclr/device/devprogram.hpp b/rocclr/device/devprogram.hpp index 610810cba1..69ca9aa7cf 100644 --- a/rocclr/device/devprogram.hpp +++ b/rocclr/device/devprogram.hpp @@ -20,7 +20,9 @@ #pragma once +#if defined(WITH_COMPILER_LIB) #include "aclTypes.h" +#endif #include "platform/context.hpp" #include "platform/object.hpp" #include "platform/memory.hpp" @@ -82,6 +84,30 @@ class Program : public amd::HeapObject { TYPE_INTERMEDIATE // intermediate } type_t; + //! type of the input file + typedef enum { + FILE_TYPE_DEFAULT = 0, + FILE_TYPE_OPENCL = 1, + FILE_TYPE_LLVMIR_TEXT = 2, + FILE_TYPE_LLVMIR_BINARY = 3, + FILE_TYPE_SPIR_TEXT = 4, + FILE_TYPE_SPIR_BINARY = 5, + FILE_TYPE_AMDIL_TEXT = 6, + FILE_TYPE_AMDIL_BINARY = 7, + FILE_TYPE_HSAIL_TEXT = 8, + FILE_TYPE_HSAIL_BINARY = 9, + FILE_TYPE_X86_TEXT = 10, + FILE_TYPE_X86_BINARY = 11, + FILE_TYPE_CG = 12, + FILE_TYPE_SOURCE = 13, + FILE_TYPE_ISA = 14, + FILE_TYPE_HEADER = 15, + FILE_TYPE_RSLLVMIR_BINARY = 16, + FILE_TYPE_SPIRV_BINARY = 17, + FILE_TYPE_ASM_TEXT = 18, + FILE_TYPE_LAST = 19 + } file_type_t; + private: //! The device target for this binary. amd::SharedReference device_; @@ -109,15 +135,19 @@ class Program : public amd::HeapObject { std::string linkOptions_; //!< link options. //!< the option arg passed in to clCompileProgram(), clLinkProgram(), //! or clBuildProgram(), whichever is called last +#if defined(WITH_COMPILER_LIB) aclBinaryOptions binOpts_; //!< Binary options to create aclBinary aclBinary* binaryElf_; //!< Binary for the new compiler library +#endif std::string lastBuildOptionsArg_; mutable std::string buildLog_; //!< build log. int32_t buildStatus_; //!< build status. int32_t buildError_; //!< build error +#if defined(WITH_COMPILER_LIB) aclTargetInfo info_; //!< The info target for this binary. +#endif size_t globalVariableTotalSize_; amd::option::Options* programOptions_; @@ -201,8 +231,10 @@ class Program : public amd::HeapObject { size_t globalVariableTotalSize() const { return globalVariableTotalSize_; } +#if defined(WITH_COMPILER_LIB) //! Returns the aclBinary associated with the program aclBinary* binaryElf() const { return static_cast(binaryElf_); } +#endif //! Returns TRUE if the program just compiled bool isNull() const { return isNull_; } @@ -288,8 +320,10 @@ class Program : public amd::HeapObject { //! Release the Binary void releaseClBinary(); +#if defined(WITH_COMPILER_LIB) //! return target info virtual const aclTargetInfo& info() = 0; +#endif virtual bool setKernels( amd::option::Options* options, void* binary, size_t binSize, @@ -314,13 +348,13 @@ class Program : public amd::HeapObject { * also returns completeStages in a vector, which contains at least ACL_TYPE_DEFAULT, * sets needOptionsCheck to true if options check is needed to decide whether or not to recompile */ - aclType getCompilationStagesFromBinary( - std::vector& completeStages, + file_type_t getCompilationStagesFromBinary( + std::vector& completeStages, bool& needOptionsCheck); /* \brief Returns the next stage to compile from, based on sections and options in binary */ - aclType getNextCompilationStageFromBinary(amd::option::Options* options); + file_type_t getNextCompilationStageFromBinary(amd::option::Options* options); //! Finds the total size of all global variables in the program bool FindGlobalVarSize(void* binary, size_t binSize); diff --git a/rocclr/device/pal/paldevice.cpp b/rocclr/device/pal/paldevice.cpp index d72c996516..7a6b266d0a 100644 --- a/rocclr/device/pal/paldevice.cpp +++ b/rocclr/device/pal/paldevice.cpp @@ -154,7 +154,9 @@ Util::GenericAllocator NullDevice::allocator_; char* Device::platformObj_; Pal::IPlatform* Device::platform_; +#if defined(WITH_COMPILER_LIB) NullDevice::Compiler* NullDevice::compiler_; +#endif AppProfile Device::appProfile_; NullDevice::NullDevice() : amd::Device(), ipLevel_(Pal::GfxIpLevel::None), palName_(nullptr) {} diff --git a/rocclr/device/pal/paldevice.hpp b/rocclr/device/pal/paldevice.hpp index 778f9f9f86..51cc25331b 100644 --- a/rocclr/device/pal/paldevice.hpp +++ b/rocclr/device/pal/paldevice.hpp @@ -53,10 +53,14 @@ namespace pal { //! A nil device object class NullDevice : public amd::Device { protected: +#if defined(WITH_COMPILER_LIB) static Compiler* compiler_; +#endif public: +#if defined(WITH_COMPILER_LIB) Compiler* compiler() const { return compiler_; } +#endif public: static bool init(void); diff --git a/rocclr/device/pal/palprogram.hpp b/rocclr/device/pal/palprogram.hpp index 14ae14e510..f2e26297c3 100644 --- a/rocclr/device/pal/palprogram.hpp +++ b/rocclr/device/pal/palprogram.hpp @@ -202,7 +202,9 @@ class HSAILProgram : public device::Program { virtual bool createBinary(amd::option::Options* options); +#if defined(WITH_COMPILER_LIB) virtual const aclTargetInfo& info(); +#endif virtual bool setKernels(amd::option::Options* options, void* binary, size_t binSize, amd::Os::FileDesc fdesc = amd::Os::FDescInit(), size_t foffset = 0, diff --git a/rocclr/device/rocm/CMakeLists.txt b/rocclr/device/rocm/CMakeLists.txt index 72cbcc1fd1..79367e553b 100644 --- a/rocclr/device/rocm/CMakeLists.txt +++ b/rocclr/device/rocm/CMakeLists.txt @@ -46,7 +46,6 @@ target_include_directories(oclrocm PRIVATE ${OPENCL_DIR} ${PROJECT_SOURCE_DIR}/compiler/lib - ${PROJECT_SOURCE_DIR}/compiler/lib/include ${PROJECT_SOURCE_DIR}/compiler/lib/backends/common ${PROJECT_SOURCE_DIR}/elf ${CMAKE_CURRENT_BINARY_DIR} diff --git a/rocclr/device/rocm/rocdevice.cpp b/rocclr/device/rocm/rocdevice.cpp index ca2f976efa..ac634d4d23 100644 --- a/rocclr/device/rocm/rocdevice.cpp +++ b/rocclr/device/rocm/rocdevice.cpp @@ -94,7 +94,6 @@ extern const char* BlitSourceCode; } // namespace device namespace roc { -amd::Device::Compiler* NullDevice::compilerHandle_; bool roc::Device::isHsaInitialized_ = false; std::vector roc::Device::gpu_agents_; std::vector roc::Device::cpu_agents_; @@ -264,35 +263,10 @@ Device::~Device() { } bool NullDevice::initCompiler(bool isOffline) { -#if defined(WITH_COMPILER_LIB) - // Initialize the compiler handle if has already not been initialized - // This is destroyed in Device::teardown - acl_error error; - if (!compilerHandle_) { - aclCompilerOptions opts = { - sizeof(aclCompilerOptions_0_8), "libamdoclcl64.so", - NULL, NULL, NULL, NULL, NULL, NULL - }; - compilerHandle_ = aclCompilerInit(&opts, &error); - if (!GPU_ENABLE_LC && error != ACL_SUCCESS) { - LogError("Error initializing the compiler handle"); - return false; - } - } -#endif // defined(WITH_COMPILER_LIB) return true; } bool NullDevice::destroyCompiler() { -#if defined(WITH_COMPILER_LIB) - if (compilerHandle_ != nullptr) { - acl_error error = aclCompilerFini(compilerHandle_); - if (error != ACL_SUCCESS) { - LogError("Error closing the compiler"); - return false; - } - } -#endif // defined(WITH_COMPILER_LIB) return true; } diff --git a/rocclr/device/rocm/rocdevice.hpp b/rocclr/device/rocm/rocdevice.hpp index 82d33ccbb4..b5e323378e 100644 --- a/rocclr/device/rocm/rocdevice.hpp +++ b/rocclr/device/rocm/rocdevice.hpp @@ -120,8 +120,6 @@ class NullDevice : public amd::Device { //! Destructor for the Null device virtual ~NullDevice(); - Compiler* compiler() const { return compilerHandle_; } - const Settings& settings() const { return static_cast(*settings_); } //! Construct an HSAIL program object from the ELF assuming it is valid @@ -234,8 +232,6 @@ class NullDevice : public amd::Device { static bool initCompiler(bool isOffline); //! destroy compiler instance and handle static bool destroyCompiler(); - //! Handle to the the compiler - static Compiler* compilerHandle_; private: static constexpr bool offlineDevice_ = true; diff --git a/rocclr/device/rocm/rockernel.cpp b/rocclr/device/rocm/rockernel.cpp index 273dee2e31..58ae759f4e 100644 --- a/rocclr/device/rocm/rockernel.cpp +++ b/rocclr/device/rocm/rockernel.cpp @@ -174,118 +174,5 @@ bool LightningKernel::init() { } #endif // defined(USE_COMGR_LIBRARY) -#if defined(WITH_COMPILER_LIB) -bool HSAILKernel::init() { - acl_error errorCode; - // Pull out metadata from the ELF - size_t sizeOfArgList; - aclCompiler* compileHandle = program()->rocDevice().compiler(); - std::string openClKernelName("&__OpenCL_" + name() + "_kernel"); - errorCode = aclQueryInfo(compileHandle, program()->binaryElf(), RT_ARGUMENT_ARRAY, - openClKernelName.c_str(), nullptr, &sizeOfArgList); - if (errorCode != ACL_SUCCESS) { - LogPrintfError("Query Info failed with error code: %d \n", errorCode); - return false; - } - std::unique_ptr argList(new char[sizeOfArgList]); - errorCode = aclQueryInfo(compileHandle, program()->binaryElf(), RT_ARGUMENT_ARRAY, - openClKernelName.c_str(), argList.get(), &sizeOfArgList); - if (errorCode != ACL_SUCCESS) { - LogPrintfError("Query Info failed with error code: %d \n", errorCode); - return false; - } - - // Set the argList - InitParameters((const aclArgData*)argList.get(), KernargSegmentByteSize()); - - // Set the workgroup information for the kernel - memset(&workGroupInfo_, 0, sizeof(workGroupInfo_)); - workGroupInfo_.availableLDSSize_ = program()->rocDevice().info().localMemSizePerCU_; - assert(workGroupInfo_.availableLDSSize_ > 0); - workGroupInfo_.availableSGPRs_ = 104; - workGroupInfo_.availableVGPRs_ = 256; - size_t sizeOfWorkGroupSize; - errorCode = aclQueryInfo(compileHandle, program()->binaryElf(), RT_WORK_GROUP_SIZE, - openClKernelName.c_str(), nullptr, &sizeOfWorkGroupSize); - if (errorCode != ACL_SUCCESS) { - LogPrintfError("Query Info failed with error code: %d \n", errorCode); - return false; - } - errorCode = aclQueryInfo(compileHandle, program()->binaryElf(), RT_WORK_GROUP_SIZE, - openClKernelName.c_str(), workGroupInfo_.compileSize_, - &sizeOfWorkGroupSize); - if (errorCode != ACL_SUCCESS) { - LogPrintfError("Query Info failed with error code: %d \n ", errorCode); - return false; - } - - uint32_t wavefront_size = 0; - hsa_status_t hsaStatus = hsa_agent_get_info(program()->rocDevice().getBackendDevice(), - HSA_AGENT_INFO_WAVEFRONT_SIZE, &wavefront_size); - if (HSA_STATUS_SUCCESS != hsaStatus) { - DevLogPrintfError("Could not get Wave Info Size: %d, failed with hsa_status: %d \n", - errorCode, hsaStatus); - return false; - } - assert(wavefront_size > 0); - - // Setting it the same as used LDS. - workGroupInfo_.localMemSize_ = workgroupGroupSegmentByteSize_; - workGroupInfo_.privateMemSize_ = workitemPrivateSegmentByteSize_; - workGroupInfo_.usedLDSSize_ = workgroupGroupSegmentByteSize_; - workGroupInfo_.preferredSizeMultiple_ = wavefront_size; - - // Query kernel header object to initialize the number of - // SGPR's and VGPR's used by the kernel - const void* kernelHostPtr = nullptr; - if (Device::loaderQueryHostAddress(reinterpret_cast(kernelCodeHandle_), - &kernelHostPtr) == HSA_STATUS_SUCCESS) { - auto akc = reinterpret_cast(kernelHostPtr); - workGroupInfo_.usedSGPRs_ = akc->wavefront_sgpr_count; - workGroupInfo_.usedVGPRs_ = akc->workitem_vgpr_count; - } else { - workGroupInfo_.usedSGPRs_ = 0; - workGroupInfo_.usedVGPRs_ = 0; - } - - workGroupInfo_.usedStackSize_ = 0; - workGroupInfo_.wavefrontPerSIMD_ = program()->rocDevice().info().maxWorkItemSizes_[0] / wavefront_size; - workGroupInfo_.wavefrontSize_ = wavefront_size; - if (workGroupInfo_.compileSize_[0] != 0) { - workGroupInfo_.size_ = workGroupInfo_.compileSize_[0] * workGroupInfo_.compileSize_[1] * - workGroupInfo_.compileSize_[2]; - } else { - workGroupInfo_.size_ = program()->rocDevice().info().preferredWorkGroupSize_; - } - - // Pull out printf metadata from the ELF - size_t sizeOfPrintfList; - errorCode = aclQueryInfo(compileHandle, program()->binaryElf(), RT_GPU_PRINTF_ARRAY, - openClKernelName.c_str(), nullptr, &sizeOfPrintfList); - if (errorCode != ACL_SUCCESS) { - LogPrintfError("Query Info failed with error code: %d \n", errorCode); - return false; - } - - // Make sure kernel has any printf info - if (0 != sizeOfPrintfList) { - std::unique_ptr aclPrintfList(new char[sizeOfPrintfList]); - if (!aclPrintfList) { - return false; - } - errorCode = aclQueryInfo(compileHandle, program()->binaryElf(), - RT_GPU_PRINTF_ARRAY, openClKernelName.c_str(), - aclPrintfList.get(), &sizeOfPrintfList); - if (errorCode != ACL_SUCCESS) { - return false; - } - - // Set the Printf List - InitPrintf(reinterpret_cast(aclPrintfList.get())); - } - return true; -} -#endif // defined(WITH_COMPILER_LIB) - } // namespace roc #endif // WITHOUT_HSA_BACKEND diff --git a/rocclr/device/rocm/rockernel.hpp b/rocclr/device/rocm/rockernel.hpp index 5cdf4141b0..e4719c82ee 100644 --- a/rocclr/device/rocm/rockernel.hpp +++ b/rocclr/device/rocm/rockernel.hpp @@ -21,7 +21,6 @@ #pragma once #include -#include "acl.h" #include "rocprogram.hpp" #include "top.hpp" #include "rocprintf.hpp" diff --git a/rocclr/device/rocm/rocprogram.cpp b/rocclr/device/rocm/rocprogram.cpp index 2d547f953c..755e094250 100644 --- a/rocclr/device/rocm/rocprogram.cpp +++ b/rocclr/device/rocm/rocprogram.cpp @@ -25,7 +25,6 @@ #include "utils/options.hpp" #include "rockernel.hpp" -#include "utils/bif_section_labels.hpp" #include "amd_hsa_kernel_code.h" #include @@ -38,30 +37,6 @@ namespace roc { -#if defined(WITH_COMPILER_LIB) -static hsa_status_t GetKernelNamesCallback(hsa_executable_t exec, hsa_agent_t agent, - hsa_executable_symbol_t symbol, void* data) { - std::vector* symNameList = reinterpret_cast*>(data); - - hsa_symbol_kind_t sym_type; - hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &sym_type); - - if (sym_type == HSA_SYMBOL_KIND_KERNEL) { - uint32_t len; - hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH, &len); - - char* symName = (char*)alloca(len + 1); - hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME, symName); - symName[len] = '\0'; - - std::string kernelName(symName); - symNameList->push_back(kernelName); - } - - return HSA_STATUS_SUCCESS; -} -#endif - static inline const char* hsa_strerror(hsa_status_t status) { const char* str = nullptr; if (hsa_status_string(status, &str) == HSA_STATUS_SUCCESS) { @@ -238,199 +213,14 @@ HSAILProgram::HSAILProgram(roc::NullDevice& device, amd::Program& owner) : roc::Program(device, owner) {} HSAILProgram::~HSAILProgram() { -#if defined(WITH_COMPILER_LIB) - acl_error error; - // Free the elf binary - if (binaryElf_ != nullptr) { - error = aclBinaryFini(binaryElf_); - if (error != ACL_SUCCESS) { - LogWarning("Error while destroying the acl binary \n"); - } - } -#endif // defined(WITH_COMPILER_LIB) } bool HSAILProgram::saveBinaryAndSetType(type_t type) { -#if defined(WITH_COMPILER_LIB) - void* rawBinary; - size_t size; - - // Write binary to memory - if (aclWriteToMem(binaryElf_, &rawBinary, &size) != ACL_SUCCESS) { - buildLog_ += "Failed to write binary to memory \n"; - return false; - } - clBinary()->saveBIFBinary((char*)rawBinary, size); - // Set the type of binary - setType(type); - -// Free memory containing rawBinary - binaryElf_->binOpts.dealloc(rawBinary); -#endif // defined(WITH_COMPILER_LIB) return true; } bool HSAILProgram::setKernels(amd::option::Options* options, void* binary, size_t binSize, amd::Os::FileDesc fdesc, size_t foffset, std::string uri) { -#if defined(WITH_COMPILER_LIB) - // Stop compilation if it is an offline device - HSA runtime does not - // support ISA compiled offline - if (!device().isOnline()) { - return true; - } - - size_t secSize = binSize; - void* data = binary; - - // Create an executable. - hsa_status_t status = hsa_executable_create_alt( - HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, nullptr, &hsaExecutable_); - if (status != HSA_STATUS_SUCCESS) { - buildLog_ += "Error: Failed to create executable: "; - buildLog_ += hsa_strerror(status); - buildLog_ += "\n"; - return false; - } - - // Load the code object. - status = hsa_code_object_reader_create_from_memory(data, secSize, &hsaCodeObjectReader_); - if (status != HSA_STATUS_SUCCESS) { - buildLog_ += "Error: AMD HSA Code Object Reader create failed: "; - buildLog_ += hsa_strerror(status); - buildLog_ += "\n"; - return false; - } - - hsa_agent_t hsaDevice = rocDevice().getBackendDevice(); - status = hsa_executable_load_agent_code_object(hsaExecutable_, hsaDevice, hsaCodeObjectReader_, - nullptr, nullptr); - if (status != HSA_STATUS_SUCCESS) { - buildLog_ += "Error: AMD HSA Code Object loading failed: "; - buildLog_ += hsa_strerror(status); - buildLog_ += "\n"; - return false; - } - - // Freeze the executable. - status = hsa_executable_freeze(hsaExecutable_, nullptr); - if (status != HSA_STATUS_SUCCESS) { - buildLog_ += "Error: Failed to freeze executable: "; - buildLog_ += hsa_strerror(status); - buildLog_ += "\n"; - return false; - } - - // Get the list of kernels - std::vector kernelNameList; - status = hsa_executable_iterate_agent_symbols(hsaExecutable_, hsaDevice, GetKernelNamesCallback, - (void*)&kernelNameList); - if (status != HSA_STATUS_SUCCESS) { - buildLog_ += "Error: Failed to get kernel names: "; - buildLog_ += hsa_strerror(status); - buildLog_ += "\n"; - return false; - } - - for (auto& kernelName : kernelNameList) { - // Query symbol handle for this symbol. - hsa_executable_symbol_t kernelSymbol; - status = hsa_executable_get_symbol_by_name(hsaExecutable_, kernelName.c_str(), &hsaDevice, - &kernelSymbol); - if (status != HSA_STATUS_SUCCESS) { - buildLog_ += "Error: Failed to get executable symbol: "; - buildLog_ += hsa_strerror(status); - buildLog_ += "\n"; - return false; - } - - // Query code handle for this symbol. - uint64_t kernelCodeHandle; - status = hsa_executable_symbol_get_info(kernelSymbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, - &kernelCodeHandle); - if (status != HSA_STATUS_SUCCESS) { - buildLog_ += "Error: Failed to get executable symbol info: "; - buildLog_ += hsa_strerror(status); - buildLog_ += "\n"; - return false; - } - - std::string openclKernelName = kernelName; - // Strip the opencl and kernel name - kernelName = kernelName.substr(strlen("&__OpenCL_"), kernelName.size()); - kernelName = kernelName.substr(0, kernelName.size() - strlen("_kernel")); - aclMetadata md; - md.numHiddenKernelArgs = 0; - - size_t sizeOfnumHiddenKernelArgs = sizeof(md.numHiddenKernelArgs); - acl_error errorCode = aclQueryInfo(device().compiler(), binaryElf_, - RT_NUM_KERNEL_HIDDEN_ARGS, openclKernelName.c_str(), - &md.numHiddenKernelArgs, &sizeOfnumHiddenKernelArgs); - if (errorCode != ACL_SUCCESS) { - buildLog_ += - "Error while Finalization phase: Kernel extra arguments count querying from the ELF " - "failed\n"; - return false; - } - - uint32_t workgroupGroupSegmentByteSize; - status = hsa_executable_symbol_get_info(kernelSymbol, - HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE, - &workgroupGroupSegmentByteSize); - if (status != HSA_STATUS_SUCCESS) { - buildLog_ += "Error: Failed to get group segment size info: "; - buildLog_ += hsa_strerror(status); - buildLog_ += "\n"; - return false; - } - - uint32_t workitemPrivateSegmentByteSize; - status = hsa_executable_symbol_get_info(kernelSymbol, - HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE, - &workitemPrivateSegmentByteSize); - if (status != HSA_STATUS_SUCCESS) { - buildLog_ += "Error: Failed to get private segment size info: "; - buildLog_ += hsa_strerror(status); - buildLog_ += "\n"; - return false; - } - - uint32_t kernargSegmentByteSize; - status = hsa_executable_symbol_get_info(kernelSymbol, - HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE, - &kernargSegmentByteSize); - if (status != HSA_STATUS_SUCCESS) { - buildLog_ += "Error: Failed to get kernarg segment size info: "; - buildLog_ += hsa_strerror(status); - buildLog_ += "\n"; - return false; - } - - uint32_t kernargSegmentAlignment; - status = hsa_executable_symbol_get_info( - kernelSymbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_ALIGNMENT, - &kernargSegmentAlignment); - if (status != HSA_STATUS_SUCCESS) { - buildLog_ += "Error: Failed to get kernarg segment alignment info: "; - buildLog_ += hsa_strerror(status); - buildLog_ += "\n"; - return false; - } - - Kernel* aKernel = new roc::HSAILKernel(kernelName, this, kernelCodeHandle, - workgroupGroupSegmentByteSize, - workitemPrivateSegmentByteSize, - kernargSegmentByteSize, kernargSegmentAlignment); - if (!aKernel->init()) { - buildLog_ += "Error: Kernel Init Failed "; - buildLog_ += "\n"; - return false; - } - aKernel->setUniformWorkGroupSize(options->oVariables->UniformWorkGroupSize); - aKernel->setInternalKernelFlag(compileOptions_.find("-cl-internal-kernel") != - std::string::npos); - kernels()[kernelName] = aKernel; - } -#endif // defined(WITH_COMPILER_LIB) return true; } diff --git a/rocclr/device/rocm/rocprogram.hpp b/rocclr/device/rocm/rocprogram.hpp index 36b5aa0bfb..6ed31dab0f 100644 --- a/rocclr/device/rocm/rocprogram.hpp +++ b/rocclr/device/rocm/rocprogram.hpp @@ -22,7 +22,6 @@ #ifndef WITHOUT_HSA_BACKEND -#include "acl.h" #include #include #include @@ -74,8 +73,6 @@ class Program : public device::Program { ); virtual bool createBinary(amd::option::Options* options) = 0; - virtual const aclTargetInfo& info() { return info_; } - protected: //! Disable default copy constructor Program(const Program&) = delete; diff --git a/rocclr/device/rocm/rocsettings.hpp b/rocclr/device/rocm/rocsettings.hpp index 96a80b7e8c..34baa1ec4c 100644 --- a/rocclr/device/rocm/rocsettings.hpp +++ b/rocclr/device/rocm/rocsettings.hpp @@ -22,8 +22,6 @@ #ifndef WITHOUT_HSA_BACKEND -#include "library.hpp" - /*! \addtogroup HSA OCL Stub Implementation * @{ */