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 * @{ */