From cdfc1b06beca65091f3a4b6218c8f15d2c927832 Mon Sep 17 00:00:00 2001 From: foreman Date: Fri, 10 Nov 2017 16:21:49 -0500 Subject: [PATCH] P4 to Git Change 1481396 by lmoriche@lmoriche_opencl_dev2 on 2017/11/10 16:12:29 SWDEV-118564 - [OCL-LC-ROCm] Refactor the Lightning Compiler program manager to allow the compiler library API and the ROCm-OpenCL-Driver to coexist in the same platform. - Default compiler is the Lightning Compiler - Fall back to the HSAIL compiler if the amdoclcl compiler library is in the PATH and the -legacy option is specified (or app-detect) Affected files ... ... //depot/stg/opencl/drivers/opencl/Makefile#59 edit ... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/build/Makefile.api#168 edit ... //depot/stg/opencl/drivers/opencl/compiler/lib/api/v0_8/acl.cpp#44 edit ... //depot/stg/opencl/drivers/opencl/compiler/lib/api/v0_8/aclLoaders.cpp#14 edit ... //depot/stg/opencl/drivers/opencl/compiler/lib/api/v0_8/aclValidation.cpp#7 edit ... //depot/stg/opencl/drivers/opencl/compiler/lib/build/Makefile.complib#98 edit ... //depot/stg/opencl/drivers/opencl/compiler/lib/loaders/elf/elf.hpp#26 edit ... //depot/stg/opencl/drivers/opencl/compiler/lib/utils/OPTIONS.def#137 edit ... //depot/stg/opencl/drivers/opencl/compiler/lib/utils/build/Makefile.utils#19 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/cpu/cpuprogram.cpp#71 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/device.cpp#213 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/device.hpp#292 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpuprogram.cpp#235 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palprogram.cpp#49 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/build/Makefile.oclrocm#20 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/roccompiler.cpp#38 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/roccompilerlib.cpp#8 delete ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/roccompilerlib.hpp#6 delete ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdevice.cpp#75 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rockernel.cpp#30 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rockernel.hpp#17 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocprogram.cpp#75 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocprogram.hpp#30 edit ... //depot/stg/opencl/drivers/opencl/runtime/platform/program.cpp#88 edit [ROCm/clr commit: 88ba77a1b49173920af5ab063e29618139c942ce] --- .../rocclr/compiler/lib/loaders/elf/elf.hpp | 3 + .../clr/rocclr/compiler/lib/utils/OPTIONS.def | 2 +- .../rocclr/runtime/device/cpu/cpuprogram.cpp | 4 +- projects/clr/rocclr/runtime/device/device.cpp | 63 +- projects/clr/rocclr/runtime/device/device.hpp | 28 +- .../rocclr/runtime/device/gpu/gpuprogram.cpp | 4 +- .../rocclr/runtime/device/pal/palprogram.cpp | 2 +- .../runtime/device/rocm/roccompiler.cpp | 248 ++-- .../runtime/device/rocm/roccompilerlib.cpp | 63 - .../runtime/device/rocm/roccompilerlib.hpp | 94 -- .../rocclr/runtime/device/rocm/rocdevice.cpp | 44 +- .../rocclr/runtime/device/rocm/rockernel.cpp | 37 +- .../rocclr/runtime/device/rocm/rockernel.hpp | 72 +- .../rocclr/runtime/device/rocm/rocprogram.cpp | 1321 +++++++++-------- .../rocclr/runtime/device/rocm/rocprogram.hpp | 125 +- .../clr/rocclr/runtime/platform/program.cpp | 34 +- 16 files changed, 1010 insertions(+), 1134 deletions(-) delete mode 100644 projects/clr/rocclr/runtime/device/rocm/roccompilerlib.cpp delete mode 100644 projects/clr/rocclr/runtime/device/rocm/roccompilerlib.hpp diff --git a/projects/clr/rocclr/compiler/lib/loaders/elf/elf.hpp b/projects/clr/rocclr/compiler/lib/loaders/elf/elf.hpp index 552ca360f6..3964aa30bb 100644 --- a/projects/clr/rocclr/compiler/lib/loaders/elf/elf.hpp +++ b/projects/clr/rocclr/compiler/lib/loaders/elf/elf.hpp @@ -34,6 +34,9 @@ #ifndef EM_ATI_CALIMAGE_BINARY #define EM_ATI_CALIMAGE_BINARY 125 #endif +#ifndef EM_AMDGPU +#define EM_AMDGPU 224 +#endif #ifndef ELFOSABI_AMD_OPENCL #define ELFOSABI_AMD_OPENCL 201 #endif diff --git a/projects/clr/rocclr/compiler/lib/utils/OPTIONS.def b/projects/clr/rocclr/compiler/lib/utils/OPTIONS.def index b30fbb6657..196cb35852 100644 --- a/projects/clr/rocclr/compiler/lib/utils/OPTIONS.def +++ b/projects/clr/rocclr/compiler/lib/utils/OPTIONS.def @@ -499,7 +499,7 @@ OPTION(OT_BOOL, \ OA_RUNTIME|OVIS_INTERNAL|OVA_DISALLOWED|OFA_PREFIX_F, \ "bin-gpu64", NULL, \ EnableGpuElf64, \ - false, 0, 0, NULL, \ + IS_LIGHTNING ? true : false, 0, 0, NULL, \ "Generate 64-bit ELF binary for GPU (default: 32-bit)") // -fdebug-linker -fno-debug-linker diff --git a/projects/clr/rocclr/runtime/device/cpu/cpuprogram.cpp b/projects/clr/rocclr/runtime/device/cpu/cpuprogram.cpp index 179c483083..436b6ed956 100644 --- a/projects/clr/rocclr/runtime/device/cpu/cpuprogram.cpp +++ b/projects/clr/rocclr/runtime/device/cpu/cpuprogram.cpp @@ -919,7 +919,7 @@ bool Program::linkImpl(amd::option::Options* options) { // Load ISA // For elf format, setup elfIn() and this elfIn() will be released // at the end of build by finiBuild(). - if (!clBinary()->setElfIn(LP64_SWITCH(ELFCLASS32, ELFCLASS64))) { + if (!clBinary()->setElfIn()) { buildLog_ += "Internal error: Setting up input OpenCL binary failed!\n"; LogError("Setting up input binary failed"); return false; @@ -1044,7 +1044,7 @@ bool Program::linkImpl(const std::vector& inputPrograms, // Load ISA // For elf format, setup elfIn() and this elfIn() will be released // at the end of build by finiBuild(). - if (!program->clBinary()->setElfIn(LP64_SWITCH(ELFCLASS32, ELFCLASS64))) { + if (!program->clBinary()->setElfIn()) { buildLog_ += "Internal error: Setting up input OpenCL binary" " failed!\n"; diff --git a/projects/clr/rocclr/runtime/device/device.cpp b/projects/clr/rocclr/runtime/device/device.cpp index 64d49d35dd..b8c1d362d3 100644 --- a/projects/clr/rocclr/runtime/device/device.cpp +++ b/projects/clr/rocclr/runtime/device/device.cpp @@ -614,7 +614,7 @@ bool CacheCompilation::compileAndLinkExecutable(amd::opencl_driver::Compiler* C, return true; } -#endif +#endif // defined(WITH_LIGHTNING_COMPILER) } // namespace amd @@ -1088,7 +1088,7 @@ bool Program::getCompileOptionsAtLinking(const std::vector& inputProgr return true; } -bool Program::initClBinary(char* binaryIn, size_t size) { +bool Program::initClBinary(const char* binaryIn, size_t size) { if (!initClBinary()) { return false; } @@ -1096,7 +1096,7 @@ bool Program::initClBinary(char* binaryIn, size_t size) { // Save the original binary that isn't owned by ClBinary clBinary()->saveOrigBinary(binaryIn, size); - char* bin = binaryIn; + const char* bin = binaryIn; size_t sz = size; // unencrypted @@ -1130,7 +1130,7 @@ bool Program::initClBinary(char* binaryIn, size_t size) { return false; } if (info().arch_id == aclHSAIL || info().arch_id == aclHSAIL64) { - err = aclWriteToMem(aclbin_v30, reinterpret_cast(&bin), &sz); + err = aclWriteToMem(aclbin_v30, (void**)const_cast(&bin), &sz); if (err != ACL_SUCCESS) { LogWarning("aclWriteToMem failed"); aclBinaryFini(aclbin_v30); @@ -1139,7 +1139,7 @@ bool Program::initClBinary(char* binaryIn, size_t size) { aclBinaryFini(aclbin_v30); } else { aclBinary* aclbin_v21 = aclCreateFromBinary(aclbin_v30, aclBIFVersion21); - err = aclWriteToMem(aclbin_v21, reinterpret_cast(&bin), &sz); + err = aclWriteToMem(aclbin_v21, (void**)const_cast(&bin), &sz); if (err != ACL_SUCCESS) { LogWarning("aclWriteToMem failed"); aclBinaryFini(aclbin_v30); @@ -1150,7 +1150,7 @@ bool Program::initClBinary(char* binaryIn, size_t size) { aclBinaryFini(aclbin_v21); } } else -#endif // defined(WITH_LIGHTNING_COMPILER) +#endif // !defined(WITH_LIGHTNING_COMPILER) { size_t decryptedSize; if (!clBinary()->decryptElf(binaryIn, size, &decryptedBin, &decryptedSize, &encryptCode)) { @@ -1177,16 +1177,12 @@ bool Program::initClBinary(char* binaryIn, size_t size) { } -bool Program::setBinary(char* binaryIn, size_t size) { +bool Program::setBinary(const char* binaryIn, size_t size) { if (!initClBinary(binaryIn, size)) { return false; } -#if defined(WITH_LIGHTNING_COMPILER) - if (!clBinary()->setElfIn(ELFCLASS64)) { -#else // !defined(WITH_LIGHTNING_COMPILER) - if (!clBinary()->setElfIn(ELFCLASS32)) { -#endif // !defined(WITH_LIGHTNING_COMPILER) + if (!clBinary()->setElfIn()) { LogError("Setting input OCL binary failed"); return false; } @@ -1209,7 +1205,14 @@ bool Program::setBinary(char* binaryIn, size_t size) { break; } case ET_DYN: { - setType(TYPE_LIBRARY); + char* sect = NULL; + size_t sz = 0; + // FIXME: we should look for the e_machine to detect an HSACO. + if (clBinary()->elfIn()->getSection(amd::OclElf::TEXT, §, &sz) && sect && sz > 0) { + setType(TYPE_EXECUTABLE); + } else { + setType(TYPE_LIBRARY); + } break; } case ET_EXEC: { @@ -1223,35 +1226,12 @@ bool Program::setBinary(char* binaryIn, size_t size) { clBinary()->loadCompileOptions(compileOptions_); clBinary()->loadLinkOptions(linkOptions_); -#if defined(WITH_LIGHTNING_COMPILER) - // TODO: Remove this once BIF is no longer used as we should have a machinasm in - // place to get the binary type correctly from above. - // It is a workaround for executable build from the library. The code object - // binary does not have the type information. - char* sect = NULL; - size_t sz = 0; - if (clBinary()->elfIn()->getSection(amd::OclElf::TEXT, §, &sz) && sect && sz > 0) { - setType(TYPE_EXECUTABLE); - } - - sect = NULL; - sz = 0; - if (type != ET_DYN && // binary is not a library - (clBinary()->elfIn()->getSection(amd::OclElf::LLVMIR, §, &sz) && sect && sz > 0)) { - setType(TYPE_COMPILED); - } - -#endif clBinary()->resetElfIn(); return true; } bool Program::createBIFBinary(aclBinary* bin) { -#if defined(WITH_LIGHTNING_COMPILER) - assert(!"createBIFBinary() should not be called when using LC"); - return false; -#else // defined(WITH_LIGHTNING_COMPILER) acl_error err; char* binaryIn = NULL; size_t size; @@ -1263,7 +1243,6 @@ bool Program::createBIFBinary(aclBinary* bin) { clBinary()->saveBIFBinary(binaryIn, size); aclFreeMem(bin, binaryIn); return true; -#endif // defined(WITH_LIGHTNING_COMPILER) } ClBinary::ClBinary(const amd::Device& dev, BinaryImageFormat bifVer) @@ -1399,7 +1378,7 @@ void ClBinary::release() { } } -void ClBinary::saveBIFBinary(char* binaryIn, size_t size) { +void ClBinary::saveBIFBinary(const char* binaryIn, size_t size) { char* image = new char[size]; memcpy(image, binaryIn, size); @@ -1492,7 +1471,7 @@ bool ClBinary::createElfBinary(bool doencrypt, Program::type_t type) { Program::binary_t ClBinary::data() const { return std::make_pair(binary_, size_); } -bool ClBinary::setBinary(char* theBinary, size_t theBinarySize, bool allocated) { +bool ClBinary::setBinary(const char* theBinary, size_t theBinarySize, bool allocated) { release(); size_ = theBinarySize; @@ -1514,7 +1493,7 @@ void ClBinary::setFlags(int encryptCode) { } } -bool ClBinary::decryptElf(char* binaryIn, size_t size, char** decryptBin, size_t* decryptSize, +bool ClBinary::decryptElf(const char* binaryIn, size_t size, char** decryptBin, size_t* decryptSize, int* encryptCode) { *decryptBin = NULL; #if defined(HAVE_BLOWFISH_H) @@ -1540,13 +1519,13 @@ bool ClBinary::decryptElf(char* binaryIn, size_t size, char** decryptBin, size_t return true; } -bool ClBinary::setElfIn(unsigned char eclass) { +bool ClBinary::setElfIn() { if (elfIn_) return true; if (binary_ == NULL) { return false; } - elfIn_ = new amd::OclElf(eclass, binary_, size_, NULL, ELF_C_READ); + elfIn_ = new amd::OclElf(ELFCLASSNONE, binary_, size_, NULL, ELF_C_READ); if ((elfIn_ == NULL) || elfIn_->hasError()) { if (elfIn_) { delete elfIn_; diff --git a/projects/clr/rocclr/runtime/device/device.hpp b/projects/clr/rocclr/runtime/device/device.hpp index 1916ef12e5..7451547bdc 100644 --- a/projects/clr/rocclr/runtime/device/device.hpp +++ b/projects/clr/rocclr/runtime/device/device.hpp @@ -19,7 +19,7 @@ #if defined(WITH_LIGHTNING_COMPILER) #include "caching/cache.hpp" #include "driver/AmdCompiler.h" -#endif // defined(WITH_LIGHTNING_COMPILER) +#endif // defined(WITH_LIGHTNING_COMPILER) #include "acl.h" #include "hwdebug.hpp" @@ -1084,7 +1084,7 @@ class Program : public amd::HeapObject { ClBinary* clBinary() { return clBinary_; } const ClBinary* clBinary() const { return clBinary_; } - bool setBinary(char* binaryIn, size_t size); + bool setBinary(const char* binaryIn, size_t size); type_t type() const { return type_; } @@ -1116,7 +1116,7 @@ class Program : public amd::HeapObject { virtual bool createBIFBinary(aclBinary* bin); //! Initialize Binary (used only for clCreateProgramWithBinary()). - bool initClBinary(char* binaryIn, size_t size); + bool initClBinary(const char* binaryIn, size_t size); //! Initialize Binary virtual bool initClBinary() = 0; @@ -1165,10 +1165,10 @@ class ClBinary : public amd::HeapObject { /** called only in loading image routines, never called in storing routines */ - bool setBinary(char* theBinary, size_t theBinarySize, bool allocated = false); + bool setBinary(const char* theBinary, size_t theBinarySize, bool allocated = false); //! setin elfIn_ - bool setElfIn(unsigned char eclass); + bool setElfIn(); void resetElfIn(); //! set out elf @@ -1189,9 +1189,9 @@ class ClBinary : public amd::HeapObject { bool createElfBinary(bool doencrypt, Program::type_t type); // save BIF binary image - void saveBIFBinary(char* binaryIn, size_t size); + void saveBIFBinary(const char* binaryIn, size_t size); - bool decryptElf(char* binaryIn, size_t size, char** decryptBin, size_t* decryptSize, + bool decryptElf(const char* binaryIn, size_t size, char** decryptBin, size_t* decryptSize, int* encryptCode); //! Returns the binary pair for the abstraction layer @@ -1222,7 +1222,7 @@ class ClBinary : public amd::HeapObject { //! Check if the binary is recompilable bool isRecompilable(std::string& llvmBinary, amd::OclElf::oclElfPlatform thePlatform); - void saveOrigBinary(char* origBinary, size_t origSize) { + void saveOrigBinary(const char* origBinary, size_t origSize) { origBinary_ = origBinary; origSize_ = origSize; } @@ -1306,11 +1306,11 @@ class ClBinary : public amd::HeapObject { //! Releases the binary data store void release(); - char* binary_; //!< binary data + const char* binary_; //!< binary data size_t size_; //!< binary size uint flags_; //!< CL binary object flags - char* origBinary_; //!< original binary data + const char* origBinary_; //!< original binary data size_t origSize_; //!< original binary size int encryptCode_; //!< Encryption Code for input binary (0 for not encrypted) @@ -1468,11 +1468,7 @@ class SvmManager : public AllStatic { */ class Device : public RuntimeObject { protected: -#if defined(WITH_LIGHTNING_COMPILER) - typedef amd::opencl_driver::Compiler Compiler; -#else // !defined(WITH_LIGHTNING_COMPILER) typedef aclCompiler Compiler; -#endif // !defined(WITH_LIGHTNING_COMPILER) public: typedef std::list CommandQueues; @@ -1753,7 +1749,6 @@ struct KernelParameterDescriptor { }; #if defined(WITH_LIGHTNING_COMPILER) - //! Compilation process with cache support. class CacheCompilation : public amd::HeapObject { public: @@ -1790,8 +1785,7 @@ class CacheCompilation : public amd::HeapObject { StringCache codeCache_; //! Cached codes const bool isCodeCacheEnabled_; //! Code cache enable }; - -#endif +#endif // defined(WITH_LIGHTNING_COMPILER) /*! @} * @} diff --git a/projects/clr/rocclr/runtime/device/gpu/gpuprogram.cpp b/projects/clr/rocclr/runtime/device/gpu/gpuprogram.cpp index f29ab699ae..354cd1d9af 100644 --- a/projects/clr/rocclr/runtime/device/gpu/gpuprogram.cpp +++ b/projects/clr/rocclr/runtime/device/gpu/gpuprogram.cpp @@ -176,7 +176,7 @@ bool NullProgram::linkImpl(amd::option::Options* options) { if (ilProgram_.empty()) { // Setup elfIn() and try to load ISA from binary // This elfIn() will be released at the end of build by finiBuild(). - if (!clBinary()->setElfIn(ELFCLASS32)) { + if (!clBinary()->setElfIn()) { buildLog_ += "Internal error: Setting input OCL binary failed!\n"; LogError("Setting input OCL binary failed"); return false; @@ -479,7 +479,7 @@ bool NullProgram::linkImpl(const std::vector& inputPrograms, // We are using CL binary directly. // Setup elfIn() and try to load llvmIR from binary // This elfIn() will be released at the end of build by finiBuild(). - if (!program->clBinary()->setElfIn(ELFCLASS32)) { + if (!program->clBinary()->setElfIn()) { buildLog_ += "Internal error: Setting input OCL binary failed!\n"; LogError("Setting input OCL binary failed"); return false; diff --git a/projects/clr/rocclr/runtime/device/pal/palprogram.cpp b/projects/clr/rocclr/runtime/device/pal/palprogram.cpp index 695fa871cd..aa196c61df 100644 --- a/projects/clr/rocclr/runtime/device/pal/palprogram.cpp +++ b/projects/clr/rocclr/runtime/device/pal/palprogram.cpp @@ -1280,7 +1280,7 @@ bool LightningProgram::linkImpl(const std::vector& inputPrograms, // We are using CL binary directly. // Setup elfIn() and try to load llvmIR from binary // This elfIn() will be released at the end of build by finiBuild(). - if (!program->clBinary()->setElfIn(ELFCLASS64)) { + if (!program->clBinary()->setElfIn()) { buildLog_ += "Internal error: Setting input OCL binary failed!\n"; return false; } diff --git a/projects/clr/rocclr/runtime/device/rocm/roccompiler.cpp b/projects/clr/rocclr/runtime/device/rocm/roccompiler.cpp index 1c167d3051..d66213945d 100644 --- a/projects/clr/rocclr/runtime/device/rocm/roccompiler.cpp +++ b/projects/clr/rocclr/runtime/device/rocm/roccompiler.cpp @@ -15,9 +15,7 @@ #if defined(WITH_LIGHTNING_COMPILER) #include "opencl1.2-c.amdgcn.inc" #include "opencl2.0-c.amdgcn.inc" -#else // !defined(WITH_LIGHTNING_COMPILER) -#include "roccompilerlib.hpp" -#endif // !defined(WITH_LIGHTNING_COMPILER) +#endif // defined(WITH_LIGHTNING_COMPILER) #include "utils/options.hpp" #include @@ -48,10 +46,125 @@ static void logFunction(const char* msg, size_t size) { static int programsCount = 0; +bool HSAILProgram::compileImpl(const std::string& sourceCode, + const std::vector& headers, + const char** headerIncludeNames, amd::option::Options* options) { + acl_error errorCode; + aclTargetInfo target; + + target = aclGetTargetInfo(LP64_SWITCH("hsail", "hsail64"), + dev().deviceInfo().complibTarget_, &errorCode); + + // end if asic info is ready + // We dump the source code for each program (param: headers) + // into their filenames (headerIncludeNames) into the TEMP + // folder specific to the OS and add the include path while + // compiling + + // Find the temp folder for the OS + std::string tempFolder = amd::Os::getEnvironment("TEMP"); + if (tempFolder.empty()) { + tempFolder = amd::Os::getEnvironment("TMP"); + if (tempFolder.empty()) { + tempFolder = WINDOWS_SWITCH(".", "/tmp"); + ; + } + } + // Iterate through each source code and dump it into tmp + std::fstream f; + std::vector headerFileNames(headers.size()); + std::vector newDirs; + for (size_t i = 0; i < headers.size(); ++i) { + std::string headerPath = tempFolder; + std::string headerIncludeName(headerIncludeNames[i]); + // replace / in path with current os's file separator + if (amd::Os::fileSeparator() != '/') { + for (std::string::iterator it = headerIncludeName.begin(), end = headerIncludeName.end(); + it != end; ++it) { + if (*it == '/') *it = amd::Os::fileSeparator(); + } + } + size_t pos = headerIncludeName.rfind(amd::Os::fileSeparator()); + if (pos != std::string::npos) { + headerPath += amd::Os::fileSeparator(); + headerPath += headerIncludeName.substr(0, pos); + headerIncludeName = headerIncludeName.substr(pos + 1); + } + if (!amd::Os::pathExists(headerPath)) { + bool ret = amd::Os::createPath(headerPath); + assert(ret && "failed creating path!"); + newDirs.push_back(headerPath); + } + std::string headerFullName = headerPath + amd::Os::fileSeparator() + headerIncludeName; + headerFileNames[i] = headerFullName; + f.open(headerFullName.c_str(), std::fstream::out); + // Should we allow asserts + assert(!f.fail() && "failed creating header file!"); + f.write(headers[i]->c_str(), headers[i]->length()); + f.close(); + } + + // Create Binary + binaryElf_ = aclBinaryInit(sizeof(aclBinary), &target, &binOpts_, &errorCode); + + if (errorCode != ACL_SUCCESS) { + buildLog_ += + "Error while compiling opencl source:\ + aclBinary init failure \n"; + LogWarning("aclBinaryInit failed"); + return false; + } + + // Insert opencl into binary + errorCode = aclInsertSection(device().compiler(), binaryElf_, sourceCode.c_str(), + strlen(sourceCode.c_str()), aclSOURCE); + + if (errorCode != ACL_SUCCESS) { + buildLog_ += + "Error while converting to BRIG: \ + Inserting openCl Source \n"; + } + + // Set the options for the compiler + // Set the include path for the temp folder that contains the includes + if (!headers.empty()) { + this->compileOptions_.append(" -I"); + this->compileOptions_.append(tempFolder); + } + + // Add only for CL2.0 and later + if (options->oVariables->CLStd[2] >= '2') { + std::stringstream opts; + opts << " -D" + << "CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE=" << device().info().maxGlobalVariableSize_; + compileOptions_.append(opts.str()); + } + + // Compile source to IR + this->compileOptions_.append(preprocessorOptions(options)); + this->compileOptions_.append(codegenOptions(options)); + + errorCode = aclCompile(device().compiler(), binaryElf_, + //"-Wf,--support_all_extensions", + this->compileOptions_.c_str(), ACL_TYPE_OPENCL, + ACL_TYPE_LLVMIR_BINARY, logFunction); + buildLog_ += aclGetCompilerLog(device().compiler()); + if (errorCode != ACL_SUCCESS) { + LogWarning("aclCompile failed"); + buildLog_ += + "Error while compiling \ + opencl source: Compiling CL to IR"; + return false; + } + // Save the binary in the interface class + saveBinaryAndSetType(TYPE_COMPILED); + return true; +} + #if defined(WITH_LIGHTNING_COMPILER) -bool HSAILProgram::compileImpl_LC(const std::string& sourceCode, - const std::vector& headers, - const char** headerIncludeNames, amd::option::Options* options) { +bool LightningProgram::compileImpl(const std::string& sourceCode, + const std::vector& headers, + const char** headerIncludeNames, amd::option::Options* options) { using namespace amd::opencl_driver; std::unique_ptr C(newCompilerInstance()); std::vector inputs; @@ -238,128 +351,7 @@ bool HSAILProgram::compileImpl_LC(const std::string& sourceCode, } return true; } -#endif // defined(WITH_LIGHTNING_COMPILER) -bool HSAILProgram::compileImpl(const std::string& sourceCode, - const std::vector& headers, - const char** headerIncludeNames, amd::option::Options* options) { -#if defined(WITH_LIGHTNING_COMPILER) - return compileImpl_LC(sourceCode, headers, headerIncludeNames, options); -#else // !defined(WITH_LIGHTNING_COMPILER) - acl_error errorCode; - aclTargetInfo target; - - target = g_complibApi._aclGetTargetInfo(LP64_SWITCH("hsail", "hsail64"), - dev().deviceInfo().complibTarget_, &errorCode); - - // end if asic info is ready - // We dump the source code for each program (param: headers) - // into their filenames (headerIncludeNames) into the TEMP - // folder specific to the OS and add the include path while - // compiling - - // Find the temp folder for the OS - std::string tempFolder = amd::Os::getEnvironment("TEMP"); - if (tempFolder.empty()) { - tempFolder = amd::Os::getEnvironment("TMP"); - if (tempFolder.empty()) { - tempFolder = WINDOWS_SWITCH(".", "/tmp"); - ; - } - } - // Iterate through each source code and dump it into tmp - std::fstream f; - std::vector headerFileNames(headers.size()); - std::vector newDirs; - for (size_t i = 0; i < headers.size(); ++i) { - std::string headerPath = tempFolder; - std::string headerIncludeName(headerIncludeNames[i]); - // replace / in path with current os's file separator - if (amd::Os::fileSeparator() != '/') { - for (std::string::iterator it = headerIncludeName.begin(), end = headerIncludeName.end(); - it != end; ++it) { - if (*it == '/') *it = amd::Os::fileSeparator(); - } - } - size_t pos = headerIncludeName.rfind(amd::Os::fileSeparator()); - if (pos != std::string::npos) { - headerPath += amd::Os::fileSeparator(); - headerPath += headerIncludeName.substr(0, pos); - headerIncludeName = headerIncludeName.substr(pos + 1); - } - if (!amd::Os::pathExists(headerPath)) { - bool ret = amd::Os::createPath(headerPath); - assert(ret && "failed creating path!"); - newDirs.push_back(headerPath); - } - std::string headerFullName = headerPath + amd::Os::fileSeparator() + headerIncludeName; - headerFileNames[i] = headerFullName; - f.open(headerFullName.c_str(), std::fstream::out); - // Should we allow asserts - assert(!f.fail() && "failed creating header file!"); - f.write(headers[i]->c_str(), headers[i]->length()); - f.close(); - } - - // Create Binary - binaryElf_ = g_complibApi._aclBinaryInit(sizeof(aclBinary), &target, &binOpts_, &errorCode); - - if (errorCode != ACL_SUCCESS) { - buildLog_ += - "Error while compiling opencl source:\ - aclBinary init failure \n"; - LogWarning("aclBinaryInit failed"); - return false; - } - - // Insert opencl into binary - errorCode = g_complibApi._aclInsertSection(device().compiler(), binaryElf_, sourceCode.c_str(), - strlen(sourceCode.c_str()), aclSOURCE); - - if (errorCode != ACL_SUCCESS) { - buildLog_ += - "Error while converting to BRIG: \ - Inserting openCl Source \n"; - } - - // Set the options for the compiler - // Set the include path for the temp folder that contains the includes - if (!headers.empty()) { - this->compileOptions_.append(" -I"); - this->compileOptions_.append(tempFolder); - } - - // Add only for CL2.0 and later - if (options->oVariables->CLStd[2] >= '2') { - std::stringstream opts; - opts << " -D" - << "CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE=" << device().info().maxGlobalVariableSize_; - compileOptions_.append(opts.str()); - } - - // Compile source to IR - this->compileOptions_.append(preprocessorOptions(options)); - this->compileOptions_.append(codegenOptions(options)); - - errorCode = g_complibApi._aclCompile(device().compiler(), binaryElf_, - //"-Wf,--support_all_extensions", - this->compileOptions_.c_str(), ACL_TYPE_OPENCL, - ACL_TYPE_LLVMIR_BINARY, logFunction); - buildLog_ += g_complibApi._aclGetCompilerLog(device().compiler()); - if (errorCode != ACL_SUCCESS) { - LogWarning("aclCompile failed"); - buildLog_ += - "Error while compiling \ - opencl source: Compiling CL to IR"; - return false; - } - // Save the binary in the interface class - saveBinaryAndSetType(TYPE_COMPILED); - return true; -#endif // !defined(WITH_LIGHTNING_COMPILER) -} - -#if defined(WITH_LIGHTNING_COMPILER) #if defined(ATI_OS_LINUX) static pthread_once_t once = PTHREAD_ONCE_INIT; @@ -396,7 +388,7 @@ static void checkLLVM_BIN() { } #endif // defined(ATI_OS_LINUX) -amd::opencl_driver::Compiler* HSAILProgram::newCompilerInstance() { +amd::opencl_driver::Compiler* LightningProgram::newCompilerInstance() { #if defined(ATI_OS_LINUX) pthread_once(&once, checkLLVM_BIN); #endif // defined(ATI_OS_LINUX) diff --git a/projects/clr/rocclr/runtime/device/rocm/roccompilerlib.cpp b/projects/clr/rocclr/runtime/device/rocm/roccompilerlib.cpp deleted file mode 100644 index 6afc916f66..0000000000 --- a/projects/clr/rocclr/runtime/device/rocm/roccompilerlib.cpp +++ /dev/null @@ -1,63 +0,0 @@ -#include "roccompilerlib.hpp" -#include "utils/flags.hpp" - -#include "acl.h" - -namespace roc { - -void* g_complibModule = nullptr; -struct CompLibApi g_complibApi; - -// -// g_complibModule is defined in LoadCompLib(). This macro must be used only in LoadCompLib() -// function. -// -#if defined(BUILD_HSA_TARGET) -#define LOADSYMBOL(api) \ - g_complibApi._##api = (pfn_##api)amd::Os::getSymbol(g_complibModule, #api); \ - if (g_complibApi._##api == nullptr) { \ - LogError("amd::Os::getSymbol() for exported func " #api " failed."); \ - amd::Os::unloadLibrary(g_complibModule); \ - return false; \ - } -#else // !defined(BUILD_HSA_TARGET) -#define LOADSYMBOL(api) g_complibApi._##api = (pfn_##api)&(api); -#endif // !defined(BUILD_HSA_TARGET) - -bool LoadCompLib(bool offline) { -#if defined(BUILD_HSA_TARGET) - g_complibModule = amd::Os::loadLibrary("amdhsacl" LP64_SWITCH(LINUX_SWITCH("32", ""), "64")); - if (g_complibModule == nullptr) { - if (!offline) { - LogError("amd::Os::loadLibrary() for loading of amdhsacl.dll failed."); - } - return false; - } -#endif // defined(BUILD_HSA_TARGET) - - LOADSYMBOL(aclCompilerInit) - LOADSYMBOL(aclGetTargetInfo) - LOADSYMBOL(aclBinaryInit) - LOADSYMBOL(aclInsertSection) - LOADSYMBOL(aclCompile) - LOADSYMBOL(aclCompilerFini) - LOADSYMBOL(aclBinaryFini) - LOADSYMBOL(aclWriteToMem) - LOADSYMBOL(aclQueryInfo) - LOADSYMBOL(aclExtractSymbol) - LOADSYMBOL(aclGetCompilerLog) - LOADSYMBOL(aclCreateFromBinary) - LOADSYMBOL(aclReadFromMem) - LOADSYMBOL(aclBinaryVersion) - LOADSYMBOL(aclLink) - - return true; -} - -void UnloadCompLib() { - if (g_complibModule) { - amd::Os::unloadLibrary(g_complibModule); - } -} - -} // namespace roc diff --git a/projects/clr/rocclr/runtime/device/rocm/roccompilerlib.hpp b/projects/clr/rocclr/runtime/device/rocm/roccompilerlib.hpp deleted file mode 100644 index 422713002c..0000000000 --- a/projects/clr/rocclr/runtime/device/rocm/roccompilerlib.hpp +++ /dev/null @@ -1,94 +0,0 @@ -#pragma once - -// -// This file hsa the code for explicity loading amdoclcl.dll. -// Exported functions from amdoclcl.dll can be added for usage as need-basis. -// With explicit/dynamic loading roc will not have any linkage to amdoclcl.lib. -// - -#include "thread/thread.hpp" -#include "acl.h" -#include "utils/debug.hpp" - -#if defined(WITH_LIGHTNING_COMPILER) -#error Should not include this file -#endif // defined(WITH_LIGHTNING_COMPILER) - -using namespace amd; - -namespace roc { - -// -// To use any new exported function from amdhsacl.dll please add/make that function specific changes -// in typedef below, struct CompLibApi and in hsacompilerLib.cpp::LoadCompLib() function. -// - -// -// Convention: The typedefed function name must be prefixed with pfn_ -// -typedef aclCompiler*(ACL_API_ENTRY* pfn_aclCompilerInit)(aclCompilerOptions* opts, - acl_error* error_code); -typedef aclTargetInfo(ACL_API_ENTRY* pfn_aclGetTargetInfo)(const char*, const char*, acl_error*); -typedef aclBinary*(ACL_API_ENTRY* pfn_aclBinaryInit)(size_t, const aclTargetInfo*, - const aclBinaryOptions*, acl_error*); -typedef acl_error(ACL_API_ENTRY* pfn_aclInsertSection)(aclCompiler* cl, aclBinary* binary, - const void* data, size_t data_size, - aclSections id); -typedef acl_error(ACL_API_ENTRY* pfn_aclCompile)(aclCompiler* cl, aclBinary* bin, - const char* options, aclType from, aclType to, - aclLogFunction compile_callback); -typedef acl_error(ACL_API_ENTRY* pfn_aclCompilerFini)(aclCompiler* cl); -typedef acl_error(ACL_API_ENTRY* pfn_aclBinaryFini)(aclBinary* bin); -typedef acl_error(ACL_API_ENTRY* pfn_aclWriteToMem)(aclBinary* bin, void** mem, size_t* size); -typedef acl_error(ACL_API_ENTRY* pfn_aclQueryInfo)(aclCompiler* cl, const aclBinary* binary, - aclQueryType query, const char* kernel, - void* data_ptr, size_t* ptr_size); -typedef const void*(ACL_API_ENTRY* pfn_aclExtractSymbol)(aclCompiler* cl, const aclBinary* binary, - size_t* size, aclSections id, - const char* symbol, acl_error* error_code); -typedef aclBinary*(ACL_API_ENTRY* pfn_aclReadFromMem)(void* mem, size_t size, - acl_error* error_code); -typedef char*(ACL_API_ENTRY* pfn_aclGetCompilerLog)(aclCompiler* cl); -typedef aclBinary*(ACL_API_ENTRY* pfn_aclCreateFromBinary)(const aclBinary* binary, - aclBIFVersion version); -typedef aclBIFVersion(ACL_API_ENTRY* pfn_aclBinaryVersion)(const aclBinary* binary); -typedef acl_error(ACL_API_ENTRY* pfn_aclLink)(aclCompiler* cl, aclBinary* src_bin, - unsigned int num_libs, aclBinary** libs, - aclType link_mode, const char* options, - aclLogFunction link_callback); -// -// Convention: prefix struct member variable with with underscore '_' -// would be nice if there was no underscore prfix, but on Linux the token -// pasting in the macro is srtict and his is the workaround. -// -struct CompLibApi { - pfn_aclCompilerInit _aclCompilerInit; - pfn_aclGetTargetInfo _aclGetTargetInfo; - pfn_aclBinaryInit _aclBinaryInit; - pfn_aclInsertSection _aclInsertSection; - pfn_aclCompile _aclCompile; - pfn_aclCompilerFini _aclCompilerFini; - pfn_aclBinaryFini _aclBinaryFini; - pfn_aclWriteToMem _aclWriteToMem; - pfn_aclQueryInfo _aclQueryInfo; - pfn_aclExtractSymbol _aclExtractSymbol; - pfn_aclReadFromMem _aclReadFromMem; - pfn_aclGetCompilerLog _aclGetCompilerLog; - pfn_aclCreateFromBinary _aclCreateFromBinary; - pfn_aclBinaryVersion _aclBinaryVersion; - pfn_aclLink _aclLink; -}; - - -// -// Use g_ prefix for all global variables. -// -extern void* g_complibModule; -extern CompLibApi g_complibApi; - -// Note: initializes global variable g_complibApi. -// Not sure what error values we have, for now returning false on failure. -bool LoadCompLib(bool isOfflineDevice = false); -void UnloadCompLib(); - -} // namespace roc diff --git a/projects/clr/rocclr/runtime/device/rocm/rocdevice.cpp b/projects/clr/rocclr/runtime/device/rocm/rocdevice.cpp index a2c805e375..04a011d275 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocdevice.cpp +++ b/projects/clr/rocclr/runtime/device/rocm/rocdevice.cpp @@ -9,6 +9,7 @@ #include "os/os.hpp" #include "utils/debug.hpp" #include "utils/flags.hpp" +#include "utils/options.hpp" #include "utils/versions.hpp" #include "thread/monitor.hpp" #include "CL/cl_ext.h" @@ -20,9 +21,7 @@ #include "device/rocm/rocprogram.hpp" #if defined(WITH_LIGHTNING_COMPILER) #include "driver/AmdCompiler.h" -#else // !defined(WITH_LIGHTNING_COMPILER) -#include "device/rocm/roccompilerlib.hpp" -#endif // !defined(WITH_LIGHTNING_COMPILER) +#endif // defined(WITH_LIGHTNING_COMPILER) #include "device/rocm/rocmemory.hpp" #include "device/rocm/rocglinterop.hpp" #ifdef WITH_AMDGPU_PRO @@ -191,46 +190,34 @@ Device::~Device() { } } bool NullDevice::initCompiler(bool isOffline) { -#if !defined(WITH_LIGHTNING_COMPILER) - // Initializes g_complibModule and g_complibApi if they were not initialized - if (g_complibModule == nullptr) { - if (!LoadCompLib(isOffline)) { - if (!isOffline) { - LogError("Error - could not find the compiler library"); - } - return false; - } - } // Initialize the compiler handle if has already not been initialized // This is destroyed in Device::teardown acl_error error; if (!compilerHandle_) { - compilerHandle_ = g_complibApi._aclCompilerInit(nullptr, &error); + aclCompilerOptions opts = { + sizeof(aclCompilerOptions_0_8), + IF(IS_LIGHTNING, "libamdoclcl64.so", NULL), + NULL, NULL, NULL, NULL, NULL, NULL + }; + compilerHandle_ = aclCompilerInit(&opts, &error); if (error != ACL_SUCCESS) { +#if !defined(WITH_LIGHTNING_COMPILER) LogError("Error initializing the compiler handle"); return false; +#endif // !defined(WITH_LIGHTNING_COMPILER) } } -#endif // !defined(WITH_LIGHTNING_COMPILER) return true; } bool NullDevice::destroyCompiler() { -#if defined(WITH_LIGHTNING_COMPILER) - delete compilerHandle_; - compilerHandle_ = nullptr; -#else // !defined(WITH_LIGHTNING_COMPILER) if (compilerHandle_ != nullptr) { - acl_error error = g_complibApi._aclCompilerFini(compilerHandle_); + acl_error error = aclCompilerFini(compilerHandle_); if (error != ACL_SUCCESS) { LogError("Error closing the compiler"); return false; } } - if (g_complibModule != nullptr) { - UnloadCompLib(); - } -#endif // !defined(WITH_LIGHTNING_COMPILER) return true; } @@ -682,6 +669,11 @@ device::Program* NullDevice::createProgram(amd::option::Options* options) { } device::Program* Device::createProgram(amd::option::Options* options) { +#if defined(WITH_LIGHTNING_COMPILER) + if (!compilerHandle_ || !options->oVariables->Legacy) { + return new roc::LightningProgram(*this); + } +#endif // defined(WITH_LIGHTNING_COMPILER) return new roc::HSAILProgram(*this); } @@ -965,7 +957,9 @@ bool Device::populateOCLDeviceConstants() { return false; } std::stringstream ss; - ss << AMD_BUILD_STRING " (HSA" << major << "." << minor << "," IF(IS_LIGHTNING, "LC", "HSAIL") ")"; + ss << AMD_BUILD_STRING " (HSA" << major << "." << minor << "," IF(IS_LIGHTNING, "LC", "HSAIL"); + if (IS_LIGHTNING && compilerHandle_) ss << "*,HSAIL"; + ss << ")"; strcpy(info_.driverVersion_, ss.str().c_str()); info_.version_ = "OpenCL " /*OPENCL_VERSION_STR*/"1.2" " "; diff --git a/projects/clr/rocclr/runtime/device/rocm/rockernel.cpp b/projects/clr/rocclr/runtime/device/rocm/rockernel.cpp index 87702089b9..56fcf2e651 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rockernel.cpp +++ b/projects/clr/rocclr/runtime/device/rocm/rockernel.cpp @@ -498,7 +498,7 @@ static inline cl_kernel_arg_type_qualifier GetOclTypeQual(const aclArgData* argI return rv; } -void Kernel::initArguments(const aclArgData* aclArg) { +void HSAILKernel::initArguments(const aclArgData* aclArg) { device::Kernel::parameters_t params; // Iterate through the arguments and insert into parameterList @@ -562,7 +562,7 @@ void Kernel::initArguments(const aclArgData* aclArg) { } #if defined(WITH_LIGHTNING_COMPILER) -void Kernel::initArguments_LC(const KernelMD& kernelMD) { +void LightningKernel::initArguments(const KernelMD& kernelMD) { device::Kernel::parameters_t params; size_t offset = 0; @@ -632,7 +632,7 @@ void Kernel::initArguments_LC(const KernelMD& kernelMD) { } #endif // defined(WITH_LIGHTNING_COMPILER) -Kernel::Kernel(std::string name, HSAILProgram* prog, const uint64_t& kernelCodeHandle, +Kernel::Kernel(std::string name, Program* prog, const uint64_t& kernelCodeHandle, const uint32_t workgroupGroupSegmentByteSize, const uint32_t workitemPrivateSegmentByteSize, const uint32_t kernargSegmentByteSize, const uint32_t kernargSegmentAlignment) @@ -645,7 +645,6 @@ Kernel::Kernel(std::string name, HSAILProgram* prog, const uint64_t& kernelCodeH kernargSegmentAlignment_(kernargSegmentAlignment) {} #if defined(WITH_LIGHTNING_COMPILER) - static const KernelMD* FindKernelMetadata(const CodeObjectMD* programMD, const std::string& name) { for (const KernelMD& kernelMD : programMD->mKernels) { if (kernelMD.mName == name) { @@ -655,18 +654,18 @@ static const KernelMD* FindKernelMetadata(const CodeObjectMD* programMD, const s return nullptr; } -bool Kernel::init_LC() { +bool LightningKernel::init() { hsa_agent_t hsaDevice = program_->hsaDevice(); // Pull out metadata from the ELF - const CodeObjectMD* programMD = program_->metadata(); + const CodeObjectMD* programMD = static_cast(program_)->metadata(); assert(programMD != nullptr); const KernelMD* kernelMD = FindKernelMetadata(programMD, name()); if (kernelMD == nullptr) { return false; } - initArguments_LC(*kernelMD); + initArguments(*kernelMD); // Set the workgroup information for the kernel workGroupInfo_.availableLDSSize_ = program_->dev().info().localMemSizePerCU_; @@ -721,16 +720,13 @@ bool Kernel::init_LC() { return false; } - initPrintf_LC(programMD->mPrintf); + initPrintf(programMD->mPrintf); return true; } #endif // defined(WITH_LIGHTNING_COMPILER) -bool Kernel::init() { -#if defined(WITH_LIGHTNING_COMPILER) - return init_LC(); -#else // !defined(WITH_LIGHTNING_COMPILER) +bool HSAILKernel::init() { acl_error errorCode; // compile kernel down to ISA hsa_agent_t hsaDevice = program_->hsaDevice(); @@ -738,13 +734,13 @@ bool Kernel::init() { size_t sizeOfArgList; aclCompiler* compileHandle = program_->dev().compiler(); std::string openClKernelName("&__OpenCL_" + name() + "_kernel"); - errorCode = g_complibApi._aclQueryInfo(compileHandle, program_->binaryElf(), RT_ARGUMENT_ARRAY, + errorCode = aclQueryInfo(compileHandle, program_->binaryElf(), RT_ARGUMENT_ARRAY, openClKernelName.c_str(), nullptr, &sizeOfArgList); if (errorCode != ACL_SUCCESS) { return false; } std::unique_ptr argList(new char[sizeOfArgList]); - errorCode = g_complibApi._aclQueryInfo(compileHandle, program_->binaryElf(), RT_ARGUMENT_ARRAY, + errorCode = aclQueryInfo(compileHandle, program_->binaryElf(), RT_ARGUMENT_ARRAY, openClKernelName.c_str(), argList.get(), &sizeOfArgList); if (errorCode != ACL_SUCCESS) { return false; @@ -760,12 +756,12 @@ bool Kernel::init() { workGroupInfo_.availableSGPRs_ = 104; workGroupInfo_.availableVGPRs_ = 256; size_t sizeOfWorkGroupSize; - errorCode = g_complibApi._aclQueryInfo(compileHandle, program_->binaryElf(), RT_WORK_GROUP_SIZE, + errorCode = aclQueryInfo(compileHandle, program_->binaryElf(), RT_WORK_GROUP_SIZE, openClKernelName.c_str(), nullptr, &sizeOfWorkGroupSize); if (errorCode != ACL_SUCCESS) { return false; } - errorCode = g_complibApi._aclQueryInfo(compileHandle, program_->binaryElf(), RT_WORK_GROUP_SIZE, + errorCode = aclQueryInfo(compileHandle, program_->binaryElf(), RT_WORK_GROUP_SIZE, openClKernelName.c_str(), workGroupInfo_.compileSize_, &sizeOfWorkGroupSize); if (errorCode != ACL_SUCCESS) { @@ -810,7 +806,7 @@ bool Kernel::init() { // Pull out printf metadata from the ELF size_t sizeOfPrintfList; - errorCode = g_complibApi._aclQueryInfo(compileHandle, program_->binaryElf(), RT_GPU_PRINTF_ARRAY, + errorCode = aclQueryInfo(compileHandle, program_->binaryElf(), RT_GPU_PRINTF_ARRAY, openClKernelName.c_str(), nullptr, &sizeOfPrintfList); if (errorCode != ACL_SUCCESS) { return false; @@ -822,7 +818,7 @@ bool Kernel::init() { if (!aclPrintfList) { return false; } - errorCode = g_complibApi._aclQueryInfo(compileHandle, program_->binaryElf(), + errorCode = aclQueryInfo(compileHandle, program_->binaryElf(), RT_GPU_PRINTF_ARRAY, openClKernelName.c_str(), aclPrintfList.get(), &sizeOfPrintfList); if (errorCode != ACL_SUCCESS) { @@ -833,11 +829,10 @@ bool Kernel::init() { initPrintf(reinterpret_cast(aclPrintfList.get())); } return true; -#endif // !defined(WITH_LIGHTNING_COMPILER) } #if defined(WITH_LIGHTNING_COMPILER) -void Kernel::initPrintf_LC(const std::vector& printfInfoStrings) { +void LightningKernel::initPrintf(const std::vector& printfInfoStrings) { for (auto str : printfInfoStrings) { std::vector tokens; @@ -928,7 +923,7 @@ void Kernel::initPrintf_LC(const std::vector& printfInfoStrings) { } #endif // defined(WITH_LIGHTNING_COMPILER) -void Kernel::initPrintf(const aclPrintfFmt* aclPrintf) { +void HSAILKernel::initPrintf(const aclPrintfFmt* aclPrintf) { PrintfInfo info; uint index = 0; for (; aclPrintf->struct_size != 0; aclPrintf++) { diff --git a/projects/clr/rocclr/runtime/device/rocm/rockernel.hpp b/projects/clr/rocclr/runtime/device/rocm/rockernel.hpp index 954324a6d8..39f4be559b 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rockernel.hpp +++ b/projects/clr/rocclr/runtime/device/rocm/rockernel.hpp @@ -85,7 +85,7 @@ class Kernel : public device::Kernel { ROC_ACCESS_TYPE access_; //!< Access type for the argument }; - Kernel(std::string name, HSAILProgram* prog, const uint64_t& kernelCodeHandle, + Kernel(std::string name, Program* prog, const uint64_t& kernelCodeHandle, const uint32_t workgroupGroupSegmentByteSize, const uint32_t workitemPrivateSegmentByteSize, const uint32_t kernargSegmentByteSize, const uint32_t kernargSegmentAlignment); @@ -103,13 +103,9 @@ class Kernel : public device::Kernel { ~Kernel(); //! Initializes the metadata required for this kernel - bool init(); -#if defined(WITH_LIGHTNING_COMPILER) - //! Initializes the metadata required for this kernel - bool init_LC(); -#endif // defined(WITH_LIGHTNING_COMPILER) + virtual bool init() = 0; - const HSAILProgram* program() const { return static_cast(program_); } + const Program* program() const { return static_cast(program_); } //! Returns the kernel argument list const std::vector& hsailArgs() const { return hsailArgList_; } @@ -131,7 +127,7 @@ class Kernel : public device::Kernel { //! set internal kernel flag void setInternalKernelFlag(bool flag) { flags_.internalKernel_ = flag; } - private: + protected: union Flags { struct { uint internalKernel_ : 1; //!< Is a blit kernel? @@ -140,21 +136,8 @@ class Kernel : public device::Kernel { Flags() : value_(0) {} } flags_; - //! Populates hsailArgList_ - void initArguments(const aclArgData* aclArg); -#if defined(WITH_LIGHTNING_COMPILER) - //! Initializes Hsail Argument metadata and info for LC - void initArguments_LC(const KernelMD& kernelMD); -#endif // defined(WITH_LIGHTNING_COMPILER) - //! Initializes HSAIL Printf metadata and info - void initPrintf(const aclPrintfFmt* aclPrintf); -#if defined(WITH_LIGHTNING_COMPILER) - //! Initializes HSAIL Printf metadata and info for LC - void initPrintf_LC(const std::vector& printfInfoStrings); -#endif // defined(WITH_LIGHTNING_COMPILER) - - HSAILProgram* program_; //!< The roc::HSAILProgram context + Program* program_; //!< The roc::Program context std::vector hsailArgList_; //!< Vector list of HSAIL Arguments uint64_t kernelCodeHandle_; //!< Kernel code handle (aka amd_kernel_code_t) const uint32_t workgroupGroupSegmentByteSize_; @@ -165,6 +148,51 @@ class Kernel : public device::Kernel { std::vector printf_; }; +class HSAILKernel : public roc::Kernel { + public: + HSAILKernel(std::string name, Program* prog, const uint64_t& kernelCodeHandle, + const uint32_t workgroupGroupSegmentByteSize, + const uint32_t workitemPrivateSegmentByteSize, + const uint32_t kernargSegmentByteSize, + const uint32_t kernargSegmentAlignment) + : roc::Kernel(name, prog, kernelCodeHandle, workgroupGroupSegmentByteSize, + workitemPrivateSegmentByteSize, kernargSegmentByteSize, kernargSegmentAlignment) { + } + + //! Initializes the metadata required for this kernel + virtual bool init() final; + + private: + //! Populates hsailArgList_ + void initArguments(const aclArgData* aclArg); + + //! Initializes HSAIL Printf metadata and info + void initPrintf(const aclPrintfFmt* aclPrintf); +}; + +#if defined(WITH_LIGHTNING_COMPILER) +class LightningKernel : public roc::Kernel { + public: + LightningKernel(std::string name, Program* prog, const uint64_t& kernelCodeHandle, + const uint32_t workgroupGroupSegmentByteSize, + const uint32_t workitemPrivateSegmentByteSize, + const uint32_t kernargSegmentByteSize, + const uint32_t kernargSegmentAlignment) + : roc::Kernel(name, prog, kernelCodeHandle, workgroupGroupSegmentByteSize, + workitemPrivateSegmentByteSize, kernargSegmentByteSize, kernargSegmentAlignment) { + } + //! Initializes the metadata required for this kernel + virtual bool init() final; + +private: + //! Initializes Hsail Argument metadata and info for LC + void initArguments(const KernelMD& kernelMD); + + //! Initializes HSAIL Printf metadata and info for LC + void initPrintf(const std::vector& printfInfoStrings); +}; +#endif // defined(WITH_LIGHTNING_COMPILER) + } // namespace roc #endif // WITHOUT_HSA_BACKEND diff --git a/projects/clr/rocclr/runtime/device/rocm/rocprogram.cpp b/projects/clr/rocclr/runtime/device/rocm/rocprogram.cpp index a0d60b24e1..2218e3ae8c 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocprogram.cpp +++ b/projects/clr/rocclr/runtime/device/rocm/rocprogram.cpp @@ -13,9 +13,7 @@ #include #include "driver/AmdCompiler.h" #include "libraries.amdgcn.inc" -#else // !defined(WITH_LIGHTNING_COMPILER) -#include "roccompilerlib.hpp" -#endif // !defined(WITH_LIGHTNING_COMPILER) +#endif // defined(WITH_LIGHTNING_COMPILER) #include "utils/bif_section_labels.hpp" #include "amd_hsa_kernel_code.h" @@ -26,7 +24,6 @@ #include #include #include -#include #include namespace roc { @@ -58,17 +55,15 @@ static void logFunction(const char* msg, size_t size) { std::cout << "Compiler Library log :" << msg << std::endl; } -HSAILProgram::~HSAILProgram() { -#if !defined(WITH_LIGHTNING_COMPILER) - acl_error error; - // Free the elf binary - if (binaryElf_ != nullptr) { - error = g_complibApi._aclBinaryFini(binaryElf_); - if (error != ACL_SUCCESS) { - LogWarning("Error while destroying the acl binary \n"); - } +static inline const char* hsa_strerror(hsa_status_t status) { + const char* str = nullptr; + if (hsa_status_string(status, &str) == HSA_STATUS_SUCCESS) { + return str; } -#endif // !defined(WITH_LIGHTNING_COMPILER) + return "Unknown error"; +} + +Program::~Program() { // Destroy the executable. if (hsaExecutable_.handle != 0) { hsa_executable_destroy(hsaExecutable_); @@ -77,13 +72,9 @@ HSAILProgram::~HSAILProgram() { hsa_code_object_reader_destroy(hsaCodeObjectReader_); } releaseClBinary(); - -#if defined(WITH_LIGHTNING_COMPILER) - delete metadata_; -#endif // defined(WITH_LIGHTNING_COMPILER) } -HSAILProgram::HSAILProgram(roc::NullDevice& device) : Program(device), binaryElf_(nullptr) { +Program::Program(roc::NullDevice& device) : device::Program(device), binaryElf_(nullptr) { memset(&binOpts_, 0, sizeof(binOpts_)); binOpts_.struct_size = sizeof(binOpts_); // binOpts_.elfclass = LP64_SWITCH( ELFCLASS32, ELFCLASS64 ); @@ -98,13 +89,9 @@ HSAILProgram::HSAILProgram(roc::NullDevice& device) : Program(device), binaryElf hsaCodeObjectReader_.handle = 0; hasGlobalStores_ = false; - -#if defined(WITH_LIGHTNING_COMPILER) - metadata_ = nullptr; -#endif // defined(WITH_LIGHTNING_COMPILER) } -bool HSAILProgram::initClBinary(char* binaryIn, size_t size) { +bool Program::initClBinary(char* binaryIn, size_t size) { // Save the original binary that isn't owned by ClBinary clBinary()->saveOrigBinary(binaryIn, size); @@ -139,7 +126,7 @@ bool HSAILProgram::initClBinary(char* binaryIn, size_t size) { } -bool HSAILProgram::initBuild(amd::option::Options* options) { +bool Program::initBuild(amd::option::Options* options) { compileOptions_ = options->origOptionStr; if (!device::Program::initBuild(options)) { @@ -159,11 +146,7 @@ bool HSAILProgram::initBuild(amd::option::Options* options) { outFileName = options->getDumpFileName(".bin"); } -#if defined(WITH_LIGHTNING_COMPILER) - bool useELF64 = true; -#else // !defined(WITH_LIGHTNING_COMPILER) bool useELF64 = getCompilerOptions()->oVariables->EnableGpuElf64; -#endif // !defined(WITH_LIGHTNING_COMPILER) if (!clBinary()->setElfOut(useELF64 ? ELFCLASS64 : ELFCLASS32, (outFileName.size() > 0) ? outFileName.c_str() : nullptr)) { LogError("Setup elf out for gpu failed"); @@ -173,7 +156,7 @@ bool HSAILProgram::initBuild(amd::option::Options* options) { } // ! post-compile setup for GPU -bool HSAILProgram::finiBuild(bool isBuildGood) { +bool Program::finiBuild(bool isBuildGood) { clBinary()->resetElfOut(); clBinary()->resetElfIn(); @@ -185,137 +168,36 @@ bool HSAILProgram::finiBuild(bool isBuildGood) { return device::Program::finiBuild(isBuildGood); } -aclType HSAILProgram::getCompilationStagesFromBinary(std::vector& completeStages, - bool& needOptionsCheck) { - acl_error errorCode; - size_t secSize = 0; - completeStages.clear(); - aclType from = ACL_TYPE_DEFAULT; - needOptionsCheck = true; - size_t boolSize = sizeof(bool); - //! @todo Should we also check for ACL_TYPE_OPENCL & ACL_TYPE_LLVMIR_TEXT? - // Checking llvmir in .llvmir section - bool containsHsailText = false; - bool containsBrig = false; - bool containsLoaderMap = false; - bool containsLlvmirText = (type() == TYPE_COMPILED); - bool containsShaderIsa = (type() == TYPE_EXECUTABLE); - bool containsOpts = !(compileOptions_.empty() && linkOptions_.empty()); -#if !defined(WITH_LIGHTNING_COMPILER) // !defined(WITH_LIGHTNING_COMPILER) - errorCode = g_complibApi._aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_LLVMIR, - nullptr, &containsLlvmirText, &boolSize); - if (errorCode != ACL_SUCCESS) { - containsLlvmirText = false; +bool Program::initClBinary() { + if (clBinary_ == nullptr) { + clBinary_ = new ClBinary(static_cast(device())); + if (clBinary_ == nullptr) { + return false; + } } - // Checking compile & link options in .comment section - errorCode = g_complibApi._aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_OPTIONS, - nullptr, &containsOpts, &boolSize); - if (errorCode != ACL_SUCCESS) { - containsOpts = false; - } - // Checking HSAIL in .cg section - containsHsailText = true; - errorCode = g_complibApi._aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_HSAIL, - nullptr, &containsHsailText, &boolSize); - if (errorCode != ACL_SUCCESS) { - containsHsailText = false; - } - // Checking BRIG sections - containsBrig = true; - errorCode = g_complibApi._aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_BRIG, nullptr, - &containsBrig, &boolSize); - if (errorCode != ACL_SUCCESS) { - containsBrig = false; - } - // Checking Loader Map symbol from CG section - containsLoaderMap = true; - errorCode = g_complibApi._aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_LOADER_MAP, - NULL, &containsLoaderMap, &boolSize); - if (errorCode != ACL_SUCCESS) { - containsLoaderMap = false; - } - if (containsBrig) { - completeStages.push_back(from); - from = ACL_TYPE_HSAIL_BINARY; - } else if (containsHsailText) { - completeStages.push_back(from); - from = ACL_TYPE_HSAIL_TEXT; - } - if (containsLoaderMap) { - completeStages.push_back(from); - from = ACL_TYPE_CG; - } - errorCode = g_complibApi._aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_ISA, nullptr, - &containsShaderIsa, &boolSize); - if (errorCode != ACL_SUCCESS) { - containsShaderIsa = false; - } -#endif // !defined(WITH_LIGHTNING_COMPILER) - - if (containsLlvmirText && containsOpts) { - completeStages.push_back(from); - from = ACL_TYPE_LLVMIR_BINARY; - } - if (containsShaderIsa) { - completeStages.push_back(from); - from = ACL_TYPE_ISA; - } - std::string sCurOptions = compileOptions_ + linkOptions_; - amd::option::Options curOptions; - if (!amd::option::parseAllOptions(sCurOptions, curOptions)) { - buildLog_ += curOptions.optionsLog(); - LogError("Parsing compile options failed."); - return ACL_TYPE_DEFAULT; - } - switch (from) { - // compile from HSAIL text, no matter prev. stages and options - case ACL_TYPE_HSAIL_TEXT: - needOptionsCheck = false; - break; - case ACL_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: - // do not check options, if LLVMIR is absent or might be absent or options are absent - if (!curOptions.oVariables->BinLLVMIR || !containsLlvmirText || !containsOpts) { - needOptionsCheck = false; - } -#if !defined(WITH_LIGHTNING_COMPILER) - // do not check options, if BRIG is absent or might be absent or LoaderMap is absent - if (!curOptions.oVariables->BinCG || !containsBrig || !containsLoaderMap) { - needOptionsCheck = false; - } -#endif - break; - // recompilation might be needed - case ACL_TYPE_LLVMIR_BINARY: - case ACL_TYPE_DEFAULT: - default: - break; - } - return from; + return true; } -aclType HSAILProgram::getNextCompilationStageFromBinary(amd::option::Options* options) { +void Program::releaseClBinary() { + if (clBinary_ != nullptr) { + delete clBinary_; + clBinary_ = nullptr; + } +} + +aclType Program::getNextCompilationStageFromBinary(amd::option::Options* options) { aclType continueCompileFrom = ACL_TYPE_DEFAULT; binary_t binary = this->binary(); // If the binary already exists if ((binary.first != nullptr) && (binary.second > 0)) { -#if defined(WITH_LIGHTNING_COMPILER) - void* mem = (void*)binary.first; -#else // !defined(WITH_LIGHTNING_COMPILER) - void* mem = const_cast(binary.first); - acl_error errorCode; - binaryElf_ = g_complibApi._aclReadFromMem(mem, binary.second, &errorCode); - if (errorCode != ACL_SUCCESS) { - buildLog_ += "Error while BRIG Codegen phase: aclReadFromMem failure \n"; - return continueCompileFrom; + if (aclValidateBinaryImage(binary.first, binary.second, BINARY_TYPE_ELF)) { + acl_error errorCode; + binaryElf_ = aclReadFromMem(binary.first, binary.second, &errorCode); + if (errorCode != ACL_SUCCESS) { + buildLog_ += "Error while BRIG Codegen phase: aclReadFromMem failure \n"; + return continueCompileFrom; + } } -#endif // !defined(WITH_LIGHTNING_COMPILER) // save the current options std::string sCurCompileOptions = compileOptions_; @@ -324,7 +206,7 @@ aclType HSAILProgram::getNextCompilationStageFromBinary(amd::option::Options* op // Saving binary in the interface class, // which also load compile & link options from binary - setBinary(static_cast(mem), binary.second); + setBinary(static_cast(binary.first), binary.second); // Calculate the next stage to compile from, based on sections in binaryElf_; // No any validity checks here @@ -344,23 +226,25 @@ aclType HSAILProgram::getNextCompilationStageFromBinary(amd::option::Options* op // If compile options are absent in binary, do not compare and recompile if (compileOptions_.empty()) break; -#if defined(WITH_LIGHTNING_COMPILER) - std::string sBinOptions = compileOptions_ + linkOptions_; -#else // !defined(WITH_LIGHTNING_COMPILER) - const oclBIFSymbolStruct* symbol = findBIF30SymStruct(symOpenclCompilerOptions); - assert(symbol && "symbol not found"); - std::string symName = - std::string(symbol->str[bif::PRE]) + std::string(symbol->str[bif::POST]); - size_t symSize = 0; + std::string sBinOptions; + if (binaryElf_ != nullptr) { + const oclBIFSymbolStruct* symbol = findBIF30SymStruct(symOpenclCompilerOptions); + assert(symbol && "symbol not found"); + std::string symName = + std::string(symbol->str[bif::PRE]) + std::string(symbol->str[bif::POST]); + size_t symSize = 0; + acl_error errorCode; - const void* opts = g_complibApi._aclExtractSymbol(device().compiler(), binaryElf_, &symSize, - aclCOMMENT, symName.c_str(), &errorCode); - if (errorCode != ACL_SUCCESS) { - recompile = true; - break; + const void* opts = aclExtractSymbol(device().compiler(), binaryElf_, &symSize, + aclCOMMENT, symName.c_str(), &errorCode); + if (errorCode != ACL_SUCCESS) { + recompile = true; + break; + } + sBinOptions = std::string((char*)opts, symSize); + } else { + sBinOptions = sCurOptions; } - std::string sBinOptions = std::string((char*)opts, symSize); -#endif // !defined(WITH_LIGHTNING_COMPILER) compileOptions_ = sCurCompileOptions; linkOptions_ = sCurLinkOptions; @@ -398,25 +282,564 @@ aclType HSAILProgram::getNextCompilationStageFromBinary(amd::option::Options* op return continueCompileFrom; } -static hsa_status_t allocFunc(size_t size, hsa_callback_data_t data, void** address) { - if (!address || 0 == size) { - return HSA_STATUS_ERROR_INVALID_ARGUMENT; +std::string HSAILProgram::codegenOptions(amd::option::Options* options) { + std::string optionsStr; + + if (dev().deviceInfo().gfxipVersion_ < 900 || !dev().settings().singleFpDenorm_) { + optionsStr.append(" -cl-denorms-are-zero"); } - *address = (char*)malloc(size); - if (!*address) { - return HSA_STATUS_ERROR_OUT_OF_RESOURCES; - } + // check if the host is 64 bit or 32 bit + LP64_ONLY(optionsStr.append(" -m64")); - return HSA_STATUS_SUCCESS; + return optionsStr; +} + +std::string Program::preprocessorOptions(amd::option::Options* options) { + std::string optionsStr; + + // Set options for the standard device specific options + + optionsStr.append(" -D__AMD__=1"); + + optionsStr.append(" -D__").append(device().info().name_).append("__=1"); + optionsStr.append(" -D__").append(device().info().name_).append("=1"); + + int major, minor; + ::sscanf(device().info().version_, "OpenCL %d.%d ", &major, &minor); + + std::stringstream ss; + ss << " -D__OPENCL_VERSION__=" << (major * 100 + minor * 10); + optionsStr.append(ss.str()); + + if (device().info().imageSupport_ && options->oVariables->ImageSupport) { + optionsStr.append(" -D__IMAGE_SUPPORT__=1"); + } + + // This is just for legacy compiler code + // All our devices support these options now + if (options->oVariables->FastFMA) { + optionsStr.append(" -DFP_FAST_FMA=1"); + } + if (options->oVariables->FastFMAF) { + optionsStr.append(" -DFP_FAST_FMAF=1"); + } + + uint clcStd = + (options->oVariables->CLStd[2] - '0') * 100 + (options->oVariables->CLStd[4] - '0') * 10; + + if (clcStd >= 200) { + std::stringstream opts; + // Add only for CL2.0 and later + opts << " -D" + << "CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE=" << device().info().maxGlobalVariableSize_; + optionsStr.append(opts.str()); + } + + // Tokenize the extensions string into a vector of strings + std::istringstream istrstr(device().info().extensions_); + std::istream_iterator sit(istrstr), end; + std::vector extensions(sit, end); + + if (IS_LIGHTNING && !options->oVariables->Legacy) { + // FIXME_lmoriche: opencl-c.h defines 'cl_khr_depth_images', so + // remove it from the command line. Should we fix opencl-c.h? + auto found = std::find(extensions.begin(), extensions.end(), "cl_khr_depth_images"); + if (found != extensions.end()) { + extensions.erase(found); + } + + if (!extensions.empty()) { + std::ostringstream clext; + + clext << " -Xclang -cl-ext=+"; + std::copy(extensions.begin(), extensions.end() - 1, + std::ostream_iterator(clext, ",+")); + clext << extensions.back(); + + optionsStr.append(clext.str()); + } + } else { + for (auto e : extensions) { + optionsStr.append(" -D").append(e).append("=1"); + } + } + + return optionsStr; +} + +HSAILProgram::HSAILProgram(roc::NullDevice& device) : roc::Program(device) { +} + +HSAILProgram::~HSAILProgram() { + 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"); + } + } +} + +bool HSAILProgram::createBinary(amd::option::Options* options) { + return false; +} + +bool HSAILProgram::saveBinaryAndSetType(type_t type) { + 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); + return true; +} + +aclType HSAILProgram::getCompilationStagesFromBinary(std::vector& completeStages, + bool& needOptionsCheck) { + acl_error errorCode; + size_t secSize = 0; + completeStages.clear(); + aclType from = ACL_TYPE_DEFAULT; + needOptionsCheck = true; + size_t boolSize = sizeof(bool); + //! @todo Should we also check for ACL_TYPE_OPENCL & ACL_TYPE_LLVMIR_TEXT? + // Checking llvmir in .llvmir section + bool containsHsailText = false; + bool containsBrig = false; + bool containsLoaderMap = false; + bool containsLlvmirText = (type() == TYPE_COMPILED); + bool containsShaderIsa = (type() == TYPE_EXECUTABLE); + bool containsOpts = !(compileOptions_.empty() && linkOptions_.empty()); + + errorCode = aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_LLVMIR, + nullptr, &containsLlvmirText, &boolSize); + if (errorCode != ACL_SUCCESS) { + containsLlvmirText = false; + } + // Checking compile & link options in .comment section + errorCode = aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_OPTIONS, + nullptr, &containsOpts, &boolSize); + if (errorCode != ACL_SUCCESS) { + containsOpts = false; + } + // Checking HSAIL in .cg section + containsHsailText = true; + errorCode = aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_HSAIL, + nullptr, &containsHsailText, &boolSize); + if (errorCode != ACL_SUCCESS) { + containsHsailText = false; + } + // Checking BRIG sections + containsBrig = true; + errorCode = aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_BRIG, nullptr, + &containsBrig, &boolSize); + if (errorCode != ACL_SUCCESS) { + containsBrig = false; + } + // Checking Loader Map symbol from CG section + containsLoaderMap = true; + errorCode = aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_LOADER_MAP, + NULL, &containsLoaderMap, &boolSize); + if (errorCode != ACL_SUCCESS) { + containsLoaderMap = false; + } + if (containsBrig) { + completeStages.push_back(from); + from = ACL_TYPE_HSAIL_BINARY; + } else if (containsHsailText) { + completeStages.push_back(from); + from = ACL_TYPE_HSAIL_TEXT; + } + if (containsLoaderMap) { + completeStages.push_back(from); + from = ACL_TYPE_CG; + } + errorCode = aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_ISA, nullptr, + &containsShaderIsa, &boolSize); + if (errorCode != ACL_SUCCESS) { + containsShaderIsa = false; + } + + if (containsLlvmirText && containsOpts) { + completeStages.push_back(from); + from = ACL_TYPE_LLVMIR_BINARY; + } + if (containsShaderIsa) { + completeStages.push_back(from); + from = ACL_TYPE_ISA; + } + std::string sCurOptions = compileOptions_ + linkOptions_; + amd::option::Options curOptions; + if (!amd::option::parseAllOptions(sCurOptions, curOptions)) { + buildLog_ += curOptions.optionsLog(); + LogError("Parsing compile options failed."); + return ACL_TYPE_DEFAULT; + } + switch (from) { + // compile from HSAIL text, no matter prev. stages and options + case ACL_TYPE_HSAIL_TEXT: + needOptionsCheck = false; + break; + case ACL_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: + // do not check options, if LLVMIR is absent or might be absent or options are absent + if (!curOptions.oVariables->BinLLVMIR || !containsLlvmirText || !containsOpts) { + needOptionsCheck = false; + } + // do not check options, if BRIG is absent or might be absent or LoaderMap is absent + if (!curOptions.oVariables->BinCG || !containsBrig || !containsLoaderMap) { + needOptionsCheck = false; + } + break; + // recompilation might be needed + case ACL_TYPE_LLVMIR_BINARY: + case ACL_TYPE_DEFAULT: + default: + break; + } + return from; +} + +bool HSAILProgram::linkImpl(const std::vector& inputPrograms, + amd::option::Options* options, bool createLibrary) { + std::vector::const_iterator it = inputPrograms.begin(); + std::vector::const_iterator itEnd = inputPrograms.end(); + acl_error errorCode; + + // For each program we need to extract the LLVMIR and create + // aclBinary for each + std::vector binaries_to_link; + + for (size_t i = 0; it != itEnd; ++it, ++i) { + Program* program = (Program*)*it; + // Check if the program was created with clCreateProgramWIthBinary + binary_t binary = program->binary(); + if ((binary.first != nullptr) && (binary.second > 0)) { + // Binary already exists -- we can also check if there is no + // opencl source code + // Need to check if LLVMIR exists in the binary + // If LLVMIR does not exist then is it valid + // We need to pull out all the compiled kernels + // We cannot do this at present because we need at least + // Hsail text to pull the kernels oout + void* mem = const_cast(binary.first); + binaryElf_ = aclReadFromMem(mem, binary.second, &errorCode); + + if (errorCode != ACL_SUCCESS) { + LogWarning("Error while linking : Could not read from raw binary"); + return false; + } + } + // At this stage each Program contains a valid binary_elf + // Check if LLVMIR is in the binary + size_t boolSize = sizeof(bool); + bool containsLLLVMIR = false; + errorCode = aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_LLVMIR, + nullptr, &containsLLLVMIR, &boolSize); + if (errorCode != ACL_SUCCESS || !containsLLLVMIR) { + buildLog_ += "Error while linking : Invalid binary (Missing LLVMIR section)"; + return false; + } + // Create a new aclBinary for each LLVMIR and save it in a list + aclBIFVersion ver = aclBinaryVersion(binaryElf_); + aclBinary* bin = aclCreateFromBinary(binaryElf_, ver); + binaries_to_link.push_back(bin); + } + + // At this stage each Program in the list has an aclBinary initialized + // and contains LLVMIR + // We can now go ahead and link them. + if (binaries_to_link.size() > 1) { + errorCode = aclLink(device().compiler(), binaries_to_link[0], + binaries_to_link.size() - 1, &binaries_to_link[1], + ACL_TYPE_LLVMIR_BINARY, "-create-library", nullptr); + } else { + errorCode = aclLink(device().compiler(), binaries_to_link[0], 0, nullptr, + ACL_TYPE_LLVMIR_BINARY, "-create-library", nullptr); + } + if (errorCode != ACL_SUCCESS) { + buildLog_ += "Failed to link programs"; + return false; + } + // Store the newly linked aclBinary for this program. + binaryElf_ = binaries_to_link[0]; + // Free all the other aclBinaries + for (size_t i = 1; i < binaries_to_link.size(); i++) { + aclBinaryFini(binaries_to_link[i]); + } + if (createLibrary) { + saveBinaryAndSetType(TYPE_LIBRARY); + return true; + } + + // Now call linkImpl with the new options + return linkImpl(options); +} + +bool HSAILProgram::linkImpl(amd::option::Options* options) { + acl_error errorCode; + aclType continueCompileFrom = ACL_TYPE_LLVMIR_BINARY; + bool finalize = true; + // If !binaryElf_ then program must have been created using clCreateProgramWithBinary + if (!binaryElf_) { + continueCompileFrom = getNextCompilationStageFromBinary(options); + } + switch (continueCompileFrom) { + // Compilation from ACL_TYPE_LLVMIR_BINARY to ACL_TYPE_CG in cases: + // 1. if the program is not created with binary; + // 2. if the program is created with binary and contains only .llvmir & .comment + // 3. if the program is created with binary, contains .llvmir, .comment, brig sections, + // but the binary's compile & link options differ from current ones (recompilation); + case ACL_TYPE_LLVMIR_BINARY: + // Compilation from ACL_TYPE_HSAIL_BINARY to ACL_TYPE_CG in cases: + // 1. if the program is created with binary and contains only brig sections + case ACL_TYPE_HSAIL_BINARY: + // Compilation from ACL_TYPE_HSAIL_TEXT to ACL_TYPE_CG in cases: + // 1. if the program is created with binary and contains only hsail text + case ACL_TYPE_HSAIL_TEXT: { + std::string curOptions = + options->origOptionStr + preprocessorOptions(options) + codegenOptions(options); + errorCode = aclCompile(device().compiler(), binaryElf_, curOptions.c_str(), + continueCompileFrom, ACL_TYPE_CG, logFunction); + buildLog_ += aclGetCompilerLog(device().compiler()); + if (errorCode != ACL_SUCCESS) { + buildLog_ += "Error while BRIG Codegen phase: compilation error \n"; + return false; + } + break; + } + case ACL_TYPE_CG: + break; + case ACL_TYPE_ISA: { + finalize = false; + break; + } + default: + buildLog_ += "Error while BRIG Codegen phase: the binary is incomplete \n"; + return false; + } + // Stop compilation if it is an offline device - HSA runtime does not + // support ISA compiled offline + if (!dev().isOnline()) { + return true; + } + + if (finalize) { + std::string fin_options(options->origOptionStr); + // Append an option so that we can selectively enable a SCOption on CZ + // whenever IOMMUv2 is enabled. + if (dev().isFineGrainedSystem(true)) { + fin_options.append(" -sc-xnack-iommu"); + } + errorCode = aclCompile(device().compiler(), binaryElf_, fin_options.c_str(), ACL_TYPE_CG, + ACL_TYPE_ISA, logFunction); + buildLog_ += aclGetCompilerLog(device().compiler()); + if (errorCode != ACL_SUCCESS) { + buildLog_ += "Error: BRIG finalization to ISA failed.\n"; + return false; + } + } + size_t secSize; + void* data = + (void*)aclExtractSection(device().compiler(), binaryElf_, &secSize, aclTEXT, &errorCode); + if (errorCode != ACL_SUCCESS) { + buildLog_ += "Error: cannot extract ISA from compiled binary.\n"; + return false; + } + + // 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. + hsa_code_object_reader_t codeObjectReader; + status = hsa_code_object_reader_create_from_memory(data, secSize, &codeObjectReader); + 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 = dev().getBackendDevice(); + status = hsa_executable_load_agent_code_object(hsaExecutable_, hsaDevice, codeObjectReader, + nullptr, nullptr); + if (status != HSA_STATUS_SUCCESS) { + buildLog_ += "Error: AMD HSA Code Object loading failed: "; + buildLog_ += hsa_strerror(status); + buildLog_ += "\n"; + return false; + } + + hsa_code_object_reader_destroy(codeObjectReader); + + // 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); + 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()) { + return false; + } + aKernel->setUniformWorkGroupSize(options->oVariables->UniformWorkGroupSize); + aKernel->setInternalKernelFlag(compileOptions_.find("-cl-internal-kernel") != + std::string::npos); + kernels()[kernelName] = aKernel; + } + saveBinaryAndSetType(TYPE_EXECUTABLE); + buildLog_ += aclGetCompilerLog(device().compiler()); + return true; } -bool HSAILProgram::saveBinaryAndSetType(type_t type, void* rawBinary, size_t size) { -// Write binary to memory #if defined(WITH_LIGHTNING_COMPILER) +LightningProgram::LightningProgram(roc::NullDevice& device) : roc::Program(device) { + metadata_ = nullptr; +} + +LightningProgram::~LightningProgram() { + delete metadata_; +} + +bool LightningProgram::createBinary(amd::option::Options* options) { + if (!clBinary()->createElfBinary(options->oVariables->BinEncrypt, type())) { + LogError("Failed to create ELF binary image!"); + return false; + } + return true; +} + +bool LightningProgram::saveBinaryAndSetType(type_t type, void* rawBinary, size_t size) { + // Write binary to memory if (type == TYPE_EXECUTABLE) { // handle code object binary assert(rawBinary != nullptr && size != 0 && "must pass in the binary"); - } else { // handle LLVM binary + } + else { // handle LLVM binary if (llvmBinary_.empty()) { buildLog_ += "ERROR: Tried to save emtpy LLVM binary \n"; return false; @@ -424,31 +847,63 @@ bool HSAILProgram::saveBinaryAndSetType(type_t type, void* rawBinary, size_t siz rawBinary = (void*)llvmBinary_.data(); size = llvmBinary_.size(); } -#else // !defined(WITH_LIGHTNING_COMPILER) - if (g_complibApi._aclWriteToMem(binaryElf_, &rawBinary, &size) != ACL_SUCCESS) { - buildLog_ += "Failed to write binary to memory \n"; - return false; - } -#endif // !defined(WITH_LIGHTNING_COMPILER) clBinary()->saveBIFBinary((char*)rawBinary, size); + // Set the type of binary setType(type); - -// Free memory containing rawBinary -#if !defined(WITH_LIGHTNING_COMPILER) - binaryElf_->binOpts.dealloc(rawBinary); -#endif return true; } -#if defined(WITH_LIGHTNING_COMPILER) -bool HSAILProgram::linkImpl_LC(const std::vector& inputPrograms, - amd::option::Options* options, bool createLibrary) { +aclType LightningProgram::getCompilationStagesFromBinary(std::vector& completeStages, + bool& needOptionsCheck) { + completeStages.clear(); + aclType from = ACL_TYPE_DEFAULT; + needOptionsCheck = true; + //! @todo Should we also check for ACL_TYPE_OPENCL & ACL_TYPE_LLVMIR_TEXT? + // Checking llvmir in .llvmir section + bool containsLlvmirText = (type() == TYPE_COMPILED); + bool containsShaderIsa = (type() == TYPE_EXECUTABLE); + bool containsOpts = !(compileOptions_.empty() && linkOptions_.empty()); + + if (containsLlvmirText && containsOpts) { + completeStages.push_back(from); + from = ACL_TYPE_LLVMIR_BINARY; + } + if (containsShaderIsa) { + completeStages.push_back(from); + from = ACL_TYPE_ISA; + } + std::string sCurOptions = compileOptions_ + linkOptions_; + amd::option::Options curOptions; + if (!amd::option::parseAllOptions(sCurOptions, curOptions)) { + buildLog_ += curOptions.optionsLog(); + LogError("Parsing compile options failed."); + return ACL_TYPE_DEFAULT; + } + switch (from) { + case ACL_TYPE_CG: + case ACL_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: + default: + break; + } + return from; +} + +bool LightningProgram::linkImpl(const std::vector& inputPrograms, + amd::option::Options* options, bool createLibrary) { using namespace amd::opencl_driver; std::unique_ptr C(newCompilerInstance()); std::vector inputs; - for (auto program : (const std::vector&)inputPrograms) { + for (auto program : (const std::vector&)inputPrograms) { if (program->llvmBinary_.empty()) { if (program->clBinary() == nullptr) { buildLog_ += "Internal error: Input program not compiled!\n"; @@ -458,7 +913,7 @@ bool HSAILProgram::linkImpl_LC(const std::vector& inputPrograms, // We are using CL binary directly. // Setup elfIn() and try to load llvmIR from binary // This elfIn() will be released at the end of build by finiBuild(). - if (!program->clBinary()->setElfIn(ELFCLASS64)) { + if (!program->clBinary()->setElfIn()) { buildLog_ += "Internal error: Setting input OCL binary failed!\n"; return false; } @@ -529,100 +984,37 @@ bool HSAILProgram::linkImpl_LC(const std::vector& inputPrograms, return true; } - return linkImpl_LC(options); + return linkImpl(options); } -#endif // defined(WITH_LIGHTNING_COMPILER) -bool HSAILProgram::linkImpl(const std::vector& inputPrograms, - amd::option::Options* options, bool createLibrary) { -#if defined(WITH_LIGHTNING_COMPILER) - return linkImpl_LC(inputPrograms, options, createLibrary); -#else // !defined(WITH_LIGHTNING_COMPILER) - std::vector::const_iterator it = inputPrograms.begin(); - std::vector::const_iterator itEnd = inputPrograms.end(); +bool LightningProgram::linkImpl(amd::option::Options* options) { acl_error errorCode; + aclType continueCompileFrom = ACL_TYPE_LLVMIR_BINARY; - // For each program we need to extract the LLVMIR and create - // aclBinary for each - std::vector binaries_to_link; - - for (size_t i = 0; it != itEnd; ++it, ++i) { - HSAILProgram* program = (HSAILProgram*)*it; - // Check if the program was created with clCreateProgramWIthBinary - binary_t binary = program->binary(); - if ((binary.first != nullptr) && (binary.second > 0)) { - // Binary already exists -- we can also check if there is no - // opencl source code - // Need to check if LLVMIR exists in the binary - // If LLVMIR does not exist then is it valid - // We need to pull out all the compiled kernels - // We cannot do this at present because we need at least - // Hsail text to pull the kernels oout - void* mem = const_cast(binary.first); - binaryElf_ = g_complibApi._aclReadFromMem(mem, binary.second, &errorCode); - - if (errorCode != ACL_SUCCESS) { - LogWarning("Error while linking : Could not read from raw binary"); + if (llvmBinary_.empty()) { + continueCompileFrom = getNextCompilationStageFromBinary(options); + } + switch (continueCompileFrom) { + case ACL_TYPE_CG: + case ACL_TYPE_LLVMIR_BINARY: { + break; + } + break; + case ACL_TYPE_ISA: { + binary_t isaBinary = binary(); + if ((isaBinary.first != nullptr) && (isaBinary.second > 0)) { + return setKernels(options, (void*)isaBinary.first, isaBinary.second); + } else { + buildLog_ += "Error: code object is empty \n"; return false; } + break; } - // At this stage each HSAILProgram contains a valid binary_elf - // Check if LLVMIR is in the binary - size_t boolSize = sizeof(bool); - bool containsLLLVMIR = false; - errorCode = g_complibApi._aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_LLVMIR, - nullptr, &containsLLLVMIR, &boolSize); - if (errorCode != ACL_SUCCESS || !containsLLLVMIR) { - buildLog_ += "Error while linking : Invalid binary (Missing LLVMIR section)"; + default: + buildLog_ += "Error while BRIG Codegen phase: the binary is incomplete \n"; return false; - } - // Create a new aclBinary for each LLVMIR and save it in a list - aclBIFVersion ver = g_complibApi._aclBinaryVersion(binaryElf_); - aclBinary* bin = g_complibApi._aclCreateFromBinary(binaryElf_, ver); - binaries_to_link.push_back(bin); } - // At this stage each HSAILProgram in the list has an aclBinary initialized - // and contains LLVMIR - // We can now go ahead and link them. - if (binaries_to_link.size() > 1) { - errorCode = g_complibApi._aclLink(device().compiler(), binaries_to_link[0], - binaries_to_link.size() - 1, &binaries_to_link[1], - ACL_TYPE_LLVMIR_BINARY, "-create-library", nullptr); - } else { - errorCode = g_complibApi._aclLink(device().compiler(), binaries_to_link[0], 0, nullptr, - ACL_TYPE_LLVMIR_BINARY, "-create-library", nullptr); - } - if (errorCode != ACL_SUCCESS) { - buildLog_ += "Failed to link programs"; - return false; - } - // Store the newly linked aclBinary for this program. - binaryElf_ = binaries_to_link[0]; - // Free all the other aclBinaries - for (size_t i = 1; i < binaries_to_link.size(); i++) { - g_complibApi._aclBinaryFini(binaries_to_link[i]); - } - if (createLibrary) { - saveBinaryAndSetType(TYPE_LIBRARY); - return true; - } - - // Now call linkImpl with the new options - return linkImpl(options); -#endif // !defined(WITH_LIGHTNING_COMPILER) -} - -static inline const char* hsa_strerror(hsa_status_t status) { - const char* str = nullptr; - if (hsa_status_string(status, &str) == HSA_STATUS_SUCCESS) { - return str; - } - return "Unknown error"; -} - -#if defined(WITH_LIGHTNING_COMPILER) -bool HSAILProgram::linkImpl_LC(amd::option::Options* options) { using namespace amd::opencl_driver; std::unique_ptr C(newCompilerInstance()); @@ -791,10 +1183,10 @@ bool HSAILProgram::linkImpl_LC(amd::option::Options* options) { } } - return setKernels_LC(options, out_exec->Buf().data(), out_exec->Size()); + return setKernels(options, out_exec->Buf().data(), out_exec->Size()); } -bool HSAILProgram::setKernels_LC(amd::option::Options* options, void* binary, size_t binSize) { +bool LightningProgram::setKernels(amd::option::Options* options, void* binary, size_t binSize) { size_t progvarsTotalSize = 0; size_t dynamicSize = 0; size_t progvarsWriteSize = 0; @@ -1006,7 +1398,7 @@ bool HSAILProgram::setKernels_LC(amd::option::Options* options, void* binary, si // the worse case: 128byte aligned. kernargSegmentAlignment = std::max(kernargSegmentAlignment, 128u); - Kernel* aKernel = new roc::Kernel( + Kernel* aKernel = new roc::LightningKernel( kernelName, this, kernelCodeHandle, workgroupGroupSegmentByteSize, workitemPrivateSegmentByteSize, kernargSegmentByteSize, amd::alignUp(kernargSegmentAlignment, device().info().globalMemCacheLineSize_)); @@ -1023,369 +1415,6 @@ bool HSAILProgram::setKernels_LC(amd::option::Options* options, void* binary, si } #endif // defined(WITH_LIGHTNING_COMPILER) -bool HSAILProgram::linkImpl(amd::option::Options* options) { - acl_error errorCode; - aclType continueCompileFrom = ACL_TYPE_LLVMIR_BINARY; - bool finalize = true; -#if !defined(WITH_LIGHTNING_COMPILER) - // If !binaryElf_ then program must have been created using clCreateProgramWithBinary - if (!binaryElf_) -#else // defined(WITH_LIGHTNING_COMPILER) - if (llvmBinary_.empty()) -#endif // defined(WITH_LIGHTNING_COMPILER) - { - continueCompileFrom = getNextCompilationStageFromBinary(options); - } - switch (continueCompileFrom) { - // Compilation from ACL_TYPE_LLVMIR_BINARY to ACL_TYPE_CG in cases: - // 1. if the program is not created with binary; - // 2. if the program is created with binary and contains only .llvmir & .comment - // 3. if the program is created with binary, contains .llvmir, .comment, brig sections, - // but the binary's compile & link options differ from current ones (recompilation); - case ACL_TYPE_LLVMIR_BINARY: - // Compilation from ACL_TYPE_HSAIL_BINARY to ACL_TYPE_CG in cases: - // 1. if the program is created with binary and contains only brig sections - case ACL_TYPE_HSAIL_BINARY: - // Compilation from ACL_TYPE_HSAIL_TEXT to ACL_TYPE_CG in cases: - // 1. if the program is created with binary and contains only hsail text - case ACL_TYPE_HSAIL_TEXT: { -#if defined(WITH_LIGHTNING_COMPILER) - if (!linkImpl_LC(options)) { - return false; - } -#else // !defined(WITH_LIGHTNING_COMPILER) - std::string curOptions = - options->origOptionStr + preprocessorOptions(options) + codegenOptions(options); - errorCode = g_complibApi._aclCompile(device().compiler(), binaryElf_, curOptions.c_str(), - continueCompileFrom, ACL_TYPE_CG, logFunction); - buildLog_ += g_complibApi._aclGetCompilerLog(device().compiler()); - if (errorCode != ACL_SUCCESS) { - buildLog_ += "Error while BRIG Codegen phase: compilation error \n"; - return false; - } -#endif // !defined(WITH_LIGHTNING_COMPILER) - break; - } - case ACL_TYPE_CG: - break; - case ACL_TYPE_ISA: { -#if defined(WITH_LIGHTNING_COMPILER) - binary_t isaBinary = binary(); - if ((isaBinary.first != nullptr) && (isaBinary.second > 0)) { - return setKernels_LC(options, (void*)isaBinary.first, isaBinary.second); - } else { - buildLog_ += "Error: code object is empty \n"; - return false; - } -#endif // !defined(WITH_LIGHTNING_COMPILER) - finalize = false; - break; - } - default: - buildLog_ += "Error while BRIG Codegen phase: the binary is incomplete \n"; - return false; - } - // Stop compilation if it is an offline device - HSA runtime does not - // support ISA compiled offline - if (!dev().isOnline()) { - return true; - } - -#if !defined(WITH_LIGHTNING_COMPILER) - if (finalize) { - std::string fin_options(options->origOptionStr); - // Append an option so that we can selectively enable a SCOption on CZ - // whenever IOMMUv2 is enabled. - if (dev().isFineGrainedSystem(true)) { - fin_options.append(" -sc-xnack-iommu"); - } - errorCode = aclCompile(dev().compiler(), binaryElf_, fin_options.c_str(), ACL_TYPE_CG, - ACL_TYPE_ISA, logFunction); - buildLog_ += aclGetCompilerLog(dev().compiler()); - if (errorCode != ACL_SUCCESS) { - buildLog_ += "Error: BRIG finalization to ISA failed.\n"; - return false; - } - } - size_t secSize; - void* data = - (void*)aclExtractSection(device().compiler(), binaryElf_, &secSize, aclTEXT, &errorCode); - if (errorCode != ACL_SUCCESS) { - buildLog_ += "Error: cannot extract ISA from compiled binary.\n"; - return false; - } - - // 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. - hsa_code_object_reader_t codeObjectReader; - status = hsa_code_object_reader_create_from_memory(data, secSize, &codeObjectReader); - 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 = dev().getBackendDevice(); - status = hsa_executable_load_agent_code_object(hsaExecutable_, hsaDevice, codeObjectReader, - nullptr, nullptr); - if (status != HSA_STATUS_SUCCESS) { - buildLog_ += "Error: AMD HSA Code Object loading failed: "; - buildLog_ += hsa_strerror(status); - buildLog_ += "\n"; - return false; - } - - hsa_code_object_reader_destroy(codeObjectReader); - - // 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); - errorCode = g_complibApi._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::Kernel(kernelName, this, kernelCodeHandle, - workgroupGroupSegmentByteSize, workitemPrivateSegmentByteSize, - kernargSegmentByteSize, kernargSegmentAlignment); - if (!aKernel->init()) { - return false; - } - aKernel->setUniformWorkGroupSize(options->oVariables->UniformWorkGroupSize); - aKernel->setInternalKernelFlag(compileOptions_.find("-cl-internal-kernel") != - std::string::npos); - kernels()[kernelName] = aKernel; - } - saveBinaryAndSetType(TYPE_EXECUTABLE); - buildLog_ += g_complibApi._aclGetCompilerLog(device().compiler()); -#endif // !defined(WITH_LIGHTNING_COMPILER) - return true; -} - -bool HSAILProgram::createBinary(amd::option::Options* options) { -#if defined(WITH_LIGHTNING_COMPILER) - if (!clBinary()->createElfBinary(options->oVariables->BinEncrypt, type())) { - LogError("Failed to create ELF binary image!"); - return false; - } - return true; -#else // !defined(WITH_LIGHTNING_COMPILER) - return false; -#endif // !defined(WITH_LIGHTNING_COMPILER) -} - -bool HSAILProgram::initClBinary() { - if (clBinary_ == nullptr) { - clBinary_ = new ClBinary(static_cast(device())); - if (clBinary_ == nullptr) { - return false; - } - } - return true; -} - -void HSAILProgram::releaseClBinary() { - if (clBinary_ != nullptr) { - delete clBinary_; - clBinary_ = nullptr; - } -} - -std::string HSAILProgram::codegenOptions(amd::option::Options* options) { - std::string optionsStr; - -#if !defined(WITH_LIGHTNING_COMPILER) - if (dev().deviceInfo().gfxipVersion_ < 900 || !dev().settings().singleFpDenorm_) { - optionsStr.append(" -cl-denorms-are-zero"); - } -#endif // !defined(WITH_LIGHTNING_COMPILER) - - // check if the host is 64 bit or 32 bit - LP64_ONLY(optionsStr.append(" -m64")); - - return optionsStr; -} - -std::string HSAILProgram::preprocessorOptions(amd::option::Options* options) { - std::string optionsStr; - - // Set options for the standard device specific options - - optionsStr.append(" -D__AMD__=1"); - - optionsStr.append(" -D__").append(device().info().name_).append("__=1"); - optionsStr.append(" -D__").append(device().info().name_).append("=1"); - - int major, minor; - ::sscanf(device().info().version_, "OpenCL %d.%d ", &major, &minor); - - std::stringstream ss; - ss << " -D__OPENCL_VERSION__=" << (major * 100 + minor * 10); - optionsStr.append(ss.str()); - - if (device().info().imageSupport_ && options->oVariables->ImageSupport) { - optionsStr.append(" -D__IMAGE_SUPPORT__=1"); - } - - // This is just for legacy compiler code - // All our devices support these options now - if (options->oVariables->FastFMA) { - optionsStr.append(" -DFP_FAST_FMA=1"); - } - if (options->oVariables->FastFMAF) { - optionsStr.append(" -DFP_FAST_FMAF=1"); - } - - uint clcStd = - (options->oVariables->CLStd[2] - '0') * 100 + (options->oVariables->CLStd[4] - '0') * 10; - - if (clcStd >= 200) { - std::stringstream opts; - // Add only for CL2.0 and later - opts << " -D" - << "CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE=" << device().info().maxGlobalVariableSize_; - optionsStr.append(opts.str()); - } - - // Tokenize the extensions string into a vector of strings - std::istringstream istrstr(device().info().extensions_); - std::istream_iterator sit(istrstr), end; - std::vector extensions(sit, end); - -#if defined(WITH_LIGHTNING_COMPILER) - // FIXME_lmoriche: opencl-c.h defines 'cl_khr_depth_images', so - // remove it from the command line. Should we fix opencl-c.h? - auto found = std::find(extensions.begin(), extensions.end(), "cl_khr_depth_images"); - if (found != extensions.end()) { - extensions.erase(found); - } - - if (!extensions.empty()) { - std::ostringstream clext; - - clext << " -Xclang -cl-ext=+"; - std::copy(extensions.begin(), extensions.end() - 1, - std::ostream_iterator(clext, ",+")); - clext << extensions.back(); - - optionsStr.append(clext.str()); - } -#else // !defined(WITH_LIGHTNING_COMPILER) - for (auto e : extensions) { - optionsStr.append(" -D").append(e).append("=1"); - } -#endif // !defined(WITH_LIGHTNING_COMPILER) - - return optionsStr; -} - } // namespace roc #endif // WITHOUT_HSA_BACKEND diff --git a/projects/clr/rocclr/runtime/device/rocm/rocprogram.hpp b/projects/clr/rocclr/runtime/device/rocm/rocprogram.hpp index de25b29377..e1154014c3 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocprogram.hpp +++ b/projects/clr/rocclr/runtime/device/rocm/rocprogram.hpp @@ -6,9 +6,6 @@ #ifndef WITHOUT_HSA_BACKEND #include "rocbinary.hpp" -#if !defined(WITH_LIGHTNING_COMPILER) -#include "roccompilerlib.hpp" -#endif // !defined(WITH_LIGHTNING_COMPILER) #include "acl.h" #include #include @@ -29,15 +26,18 @@ typedef llvm::AMDGPU::HSAMD::Kernel::Arg::Metadata KernelArgMD; //! \namespace roc HSA Device Implementation namespace roc { +class HSAILProgram; +class LightningProgram; + //! \class empty program -class HSAILProgram : public device::Program { +class Program : public device::Program { friend class ClBinary; public: //! Default constructor - HSAILProgram(roc::NullDevice& device); + Program(roc::NullDevice& device); //! Default destructor - ~HSAILProgram(); + ~Program(); // Initialize Binary for GPU (used only for clCreateProgramWithBinary()). virtual bool initClBinary(char* binaryIn, size_t size); @@ -45,11 +45,6 @@ class HSAILProgram : public device::Program { //! Returns the aclBinary associated with the program const aclBinary* binaryElf() const { return static_cast(binaryElf_); } -#if defined(WITH_LIGHTNING_COMPILER) - //! Returns the program metadata. - const CodeObjectMD* metadata() const { return metadata_; } -#endif // defined(WITH_LIGHTNING_COMPILER) - //! Return a typecasted GPU device const NullDevice& dev() const { return static_cast(device()); } @@ -65,23 +60,6 @@ class HSAILProgram : public device::Program { //! post-compile setup for GPU virtual bool finiBuild(bool isBuildGood); - /*! \brief Compiles GPU CL program to LLVM binary (compiler frontend) - * - * \return True if we successfully compiled a GPU program - */ - virtual bool compileImpl(const std::string& sourceCode, //!< the program's source code - const std::vector& headers, - const char** headerIncludeNames, - amd::option::Options* options //!< compile options's object - ); -#if defined(WITH_LIGHTNING_COMPILER) - virtual bool compileImpl_LC(const std::string& sourceCode, //!< the program's source code - const std::vector& headers, - const char** headerIncludeNames, - amd::option::Options* options //!< compile options's object - ); -#endif // defined(WITH_LIGHTNING_COMPILER) - /*! \brief Compiles LLVM binary to HSAIL code (compiler backend: link+opt+codegen) * * \return The build error code @@ -89,22 +67,7 @@ class HSAILProgram : public device::Program { int compileBinaryToHSAIL(amd::option::Options* options //!< options for compilation ); - - virtual bool linkImpl(amd::option::Options* options); -#if defined(WITH_LIGHTNING_COMPILER) - virtual bool linkImpl_LC(amd::option::Options* options); - bool setKernels_LC(amd::option::Options* options, void* binary, size_t binSize); -#endif // defined(WITH_LIGHTNING_COMPILER) - - //! Link the device programs. - virtual bool linkImpl(const std::vector& inputPrograms, amd::option::Options* options, - bool createLibrary); -#if defined(WITH_LIGHTNING_COMPILER) - virtual bool linkImpl_LC(const std::vector& inputPrograms, - amd::option::Options* options, bool createLibrary); -#endif // defined(WITH_LIGHTNING_COMPILER) - - virtual bool createBinary(amd::option::Options* options); + virtual bool createBinary(amd::option::Options* options) = 0; //! Initialize Binary virtual bool initClBinary(); @@ -127,28 +90,27 @@ class HSAILProgram : public device::Program { return static_cast(device::Program::clBinary()); } - private: + protected: /* \brief Returns the next stage to compile from, based on sections in binary, * 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, - bool& needOptionsCheck); + virtual aclType getCompilationStagesFromBinary(std::vector& completeStages, + bool& needOptionsCheck) = 0; /* \brief Returns the next stage to compile from, based on sections and options in binary */ aclType getNextCompilationStageFromBinary(amd::option::Options* options); - bool saveBinaryAndSetType(type_t type, void* binary = nullptr, size_t size = 0); //! Disable default copy constructor - HSAILProgram(const HSAILProgram&) = delete; + Program(const Program&) = delete; //! Disable operator= - HSAILProgram& operator=(const HSAILProgram&) = delete; + Program& operator=(const Program&) = delete; +protected: //! Returns all the options to be appended while passing to the // compiler std::string preprocessorOptions(amd::option::Options* options); - std::string codegenOptions(amd::option::Options* options); // aclBinary and aclCompiler - for the compiler library aclBinary* binaryElf_; //!< Binary for the new compiler library @@ -158,13 +120,72 @@ class HSAILProgram : public device::Program { /* HSA executable */ hsa_executable_t hsaExecutable_; //!< Handle to HSA executable hsa_code_object_reader_t hsaCodeObjectReader_; //!< Handle to HSA code reader +}; + +class HSAILProgram : public roc::Program { + public: + HSAILProgram(roc::NullDevice& device); + virtual ~HSAILProgram(); + + protected: + virtual bool compileImpl(const std::string& sourceCode, //!< the program's source code + const std::vector& headers, + const char** headerIncludeNames, + amd::option::Options* options //!< compile options's object + ) final; + + virtual bool linkImpl(amd::option::Options* options) final; + + virtual bool linkImpl(const std::vector& inputPrograms, + amd::option::Options* options, bool createLibrary) final; + + virtual bool createBinary(amd::option::Options* options) final; + +private: + std::string codegenOptions(amd::option::Options* options); + + virtual aclType getCompilationStagesFromBinary(std::vector& completeStages, + bool& needOptionsCheck) final; + + bool saveBinaryAndSetType(type_t type); +}; #if defined(WITH_LIGHTNING_COMPILER) +class LightningProgram : public roc::Program { +public: + LightningProgram(roc::NullDevice& device); + virtual ~LightningProgram(); + + //! Returns the program metadata. + const CodeObjectMD* metadata() const { return metadata_; } + +protected: + virtual bool compileImpl(const std::string& sourceCode, //!< the program's source code + const std::vector& headers, + const char** headerIncludeNames, + amd::option::Options* options //!< compile options's object + ) final; + + virtual bool linkImpl(amd::option::Options* options) final; + + virtual bool linkImpl(const std::vector& inputPrograms, + amd::option::Options* options, bool createLibrary) final; + + virtual bool createBinary(amd::option::Options* options) final; + +private: + bool saveBinaryAndSetType(type_t type, void* rawBinary, size_t size); + + virtual aclType getCompilationStagesFromBinary(std::vector& completeStages, + bool& needOptionsCheck) final; + + bool setKernels(amd::option::Options* options, void* binary, size_t binSize); + CodeObjectMD* metadata_; //!< Runtime metadata //! Return a new transient compiler instance. static amd::opencl_driver::Compiler* newCompilerInstance(); -#endif // defined(WITH_LIGHTNING_COMPILER) }; +#endif // defined(WITH_LIGHTNING_COMPILER) /*@}*/} // namespace roc diff --git a/projects/clr/rocclr/runtime/platform/program.cpp b/projects/clr/rocclr/runtime/platform/program.cpp index f9926a83bb..2c0b3c3039 100644 --- a/projects/clr/rocclr/runtime/platform/program.cpp +++ b/projects/clr/rocclr/runtime/platform/program.cpp @@ -49,18 +49,14 @@ const Symbol* Program::findSymbol(const char* kernelName) const { cl_int Program::addDeviceProgram(Device& device, const void* image, size_t length, amd::option::Options* options) { -#if defined(WITH_LIGHTNING_COMPILER) - // LC binary must be in ELF format - if (image != NULL && !amd::isElfMagic((const char*)image)) { + if (image != NULL && !amd::isElfMagic((const char*)image) +#if !defined(WITH_LIGHTNING_COMPILER) + && !aclValidateBinaryImage( + image, length, isSPIRV_ ? BINARY_TYPE_SPIRV : BINARY_TYPE_ELF | BINARY_TYPE_LLVM) +#endif // !defined(WITH_LIGHTNING_COMPILER) + ) { return CL_INVALID_BINARY; } -#else // !defined(WITH_LIGHTNING_COMPILER) - if (image != NULL && - !aclValidateBinaryImage(image, length, - isSPIRV_ ? BINARY_TYPE_SPIRV : BINARY_TYPE_ELF | BINARY_TYPE_LLVM)) { - return CL_INVALID_BINARY; - } -#endif // !defined(WITH_LIGHTNING_COMPILER) // Check if the device is already associated with this program if (deviceList_.find(&device) != deviceList_.end()) { @@ -80,7 +76,6 @@ cl_int Program::addDeviceProgram(Device& device, const void* image, size_t lengt emptyOptions = true; } -#if !defined(WITH_LIGHTNING_COMPILER) if (image != NULL && length != 0 && aclValidateBinaryImage(image, length, BINARY_TYPE_ELF)) { acl_error errorCode; aclBinary* binary = aclReadFromMem(image, length, &errorCode); @@ -102,10 +97,11 @@ cl_int Program::addDeviceProgram(Device& device, const void* image, size_t lengt return CL_INVALID_COMPILER_OPTIONS; } } - options->oVariables->Legacy = isAMDILTarget(*aclutGetTargetInfo(binary)); + options->oVariables->Legacy = !IS_LIGHTNING ? + isAMDILTarget(*aclutGetTargetInfo(binary)) : + isHSAILTarget(*aclutGetTargetInfo(binary)); aclBinaryFini(binary); } -#endif // !defined(WITH_LIGHTNING_COMPILER) options->oVariables->BinaryIsSpirv = isSPIRV_; device::Program* program = rootDev.createProgram(options); if (program == NULL) { @@ -133,7 +129,7 @@ cl_int Program::addDeviceProgram(Device& device, const void* image, size_t lengt return CL_INVALID_BINARY; } -#if defined(WITH_LIGHTNING_COMPILER) +#if 0 && defined(WITH_LIGHTNING_COMPILER) // load the compiler options from the binary if it is not provided std::string sBinOptions = program->compileOptions(); if (!sBinOptions.empty() && emptyOptions) { @@ -143,7 +139,7 @@ cl_int Program::addDeviceProgram(Device& device, const void* image, size_t lengt return CL_INVALID_COMPILER_OPTIONS; } } -#endif +#endif // defined(WITH_LIGHTNING_COMPILER) } devicePrograms_[&rootDev] = program; @@ -300,8 +296,8 @@ cl_int Program::link(const std::vector& devices, size_t numInputs, // Check the binary's target for the first found device program. // TODO: Revise these binary's target checks // and possibly remove them after switching to HSAIL by default. -#if !defined(WITH_LIGHTNING_COMPILER) - if (!found && binary.first != NULL && binary.second > 0) { + if (!found && binary.first != NULL && binary.second > 0 && + aclValidateBinaryImage(binary.first, binary.second, BINARY_TYPE_ELF)) { acl_error errorCode = ACL_SUCCESS; void* mem = const_cast(binary.first); aclBinary* aclBin = aclReadFromMem(mem, binary.second, &errorCode); @@ -311,12 +307,14 @@ cl_int Program::link(const std::vector& devices, size_t numInputs, } if (isHSAILTarget(*aclutGetTargetInfo(aclBin))) { parsedOptions.oVariables->Frontend = "clang"; +#if defined(WITH_LIGHTNING_COMPILER) + parsedOptions.oVariables->Legacy = true; +#endif // defined(WITH_LIGHTNING_COMPILER) } else if (isAMDILTarget(*aclutGetTargetInfo(aclBin))) { parsedOptions.oVariables->Frontend = "edg"; } aclBinaryFini(aclBin); } -#endif // !defined(WITH_LIGHTNING_COMPILER) found = true; } if (inputDevPrograms.size() == 0) {