From f9e4fcdf2cf769a29e1a90ed68cf5119c42da9a6 Mon Sep 17 00:00:00 2001 From: foreman Date: Wed, 24 Aug 2016 15:25:49 -0400 Subject: [PATCH] P4 to Git Change 1306648 by lmoriche@lmoriche_opencl_dev on 2016/08/24 15:17:25 SWDEV-94610 - [OpenCL/LC] Program Manager: Implement multi-file linking and separate compilation. Implement saving and loading programs, libraries, and executables. Affected files ... ... //depot/stg/opencl/drivers/opencl/make/clang.git/include/clang/Config/config.h#3 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/device.hpp#277 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/roccompiler.cpp#9 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rockernel.cpp#7 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocprogram.cpp#12 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocprogram.hpp#7 edit ... //depot/stg/opencl/drivers/opencl/runtime/platform/program.cpp#82 edit [ROCm/clr commit: 01ed632c6c3479d4608d700c4d6602aa37dc9642] --- projects/clr/rocclr/runtime/device/device.hpp | 2 +- .../runtime/device/rocm/roccompiler.cpp | 76 +++-- .../rocclr/runtime/device/rocm/rockernel.cpp | 4 +- .../rocclr/runtime/device/rocm/rocprogram.cpp | 277 ++++++++++++------ .../rocclr/runtime/device/rocm/rocprogram.hpp | 88 ++---- .../clr/rocclr/runtime/platform/program.cpp | 3 +- 6 files changed, 253 insertions(+), 197 deletions(-) diff --git a/projects/clr/rocclr/runtime/device/device.hpp b/projects/clr/rocclr/runtime/device/device.hpp index 32d12f8788..ee751be32d 100644 --- a/projects/clr/rocclr/runtime/device/device.hpp +++ b/projects/clr/rocclr/runtime/device/device.hpp @@ -1097,7 +1097,7 @@ public: amd::option::Options* options); //! Returns the device object, associated with this program. - const amd::Device& device() { return device_(); } + const amd::Device& device() const { return device_(); } //! Return the compiler options used to build the program. const std::string& compileOptions() const { return compileOptions_; } diff --git a/projects/clr/rocclr/runtime/device/rocm/roccompiler.cpp b/projects/clr/rocclr/runtime/device/rocm/roccompiler.cpp index 0f9042beba..7969e5bca0 100644 --- a/projects/clr/rocclr/runtime/device/rocm/roccompiler.cpp +++ b/projects/clr/rocclr/runtime/device/rocm/roccompiler.cpp @@ -41,17 +41,17 @@ HSAILProgram::compileImpl_LC(const std::string& sourceCode, amd::option::Options* options) { using namespace amd::opencl_driver; + Compiler* C = device().compiler(); std::vector inputs; - Data* src = device().compiler()->NewBufferReference(DT_CL, + Data* input = C->NewBufferReference(DT_CL, sourceCode.c_str(), sourceCode.length()); - - if (src == NULL) { + if (input == NULL) { buildLog_ += "Error while creating data from source code"; return false; } - inputs.push_back(src); + inputs.push_back(input); //Find the temp folder for the OS std::string tempFolder = amd::Os::getEnvironment("TEMP"); @@ -97,17 +97,22 @@ HSAILProgram::compileImpl_LC(const std::string& sourceCode, f.write(headers[i]->c_str(), headers[i]->length()); f.close(); - Data* inc = device().compiler()->NewFileReference(DT_CL_HEADER, - headerFileNames[i]); + Data* inc = C->NewFileReference(DT_CL_HEADER, headerFileNames[i]); + if (inc == NULL) { + buildLog_ += "Error while creating data from headers"; + return false; + } inputs.push_back(inc); } //Set the options for the compiler + std::string driverOptions(compileOptions_); + //Set the include path for the temp folder that contains the includes if(!headers.empty()) { - compileOptions_.append(" -I"); - compileOptions_.append(tempFolder); + driverOptions.append(" -I"); + driverOptions.append(tempFolder); } const char* xLang = options->oVariables->XLang; @@ -116,14 +121,14 @@ HSAILProgram::compileImpl_LC(const std::string& sourceCode, } //FIXME_Nikolay: the program manager should be setting the language - //compileOptions_.append(" -x cl"); + //driverOptions.append(" -x cl"); //FIXME_Nikolay: the program manager shouls be setting the cl-std. -Xclang //is not necessary, we add it to overridde the flag set in the comp driver. - compileOptions_.append(" -Xclang -cl-std=").append(options->oVariables->CLStd); + driverOptions.append(" -Xclang -cl-std=").append(options->oVariables->CLStd); std::ostringstream optLevel; optLevel << " -O" << options->oVariables->OptLevel; - compileOptions_.append(optLevel.str()); + driverOptions.append(optLevel.str()); //FIXME_lmoriche: has the CL option been validated? uint clVer = (options->oVariables->CLStd[2] - '0') * 100 @@ -131,56 +136,67 @@ HSAILProgram::compileImpl_LC(const std::string& sourceCode, std::pair hdr; switch(clVer) { - case 120: hdr = std::make_pair(opencl1_2_c_amdgcn, opencl1_2_c_amdgcn_size); break; - case 200: hdr = std::make_pair(opencl2_0_c_amdgcn, opencl2_0_c_amdgcn_size); break; + case 120: + hdr = std::make_pair(opencl1_2_c_amdgcn, opencl1_2_c_amdgcn_size); + break; + case 200: + hdr = std::make_pair(opencl2_0_c_amdgcn, opencl2_0_c_amdgcn_size); + break; default: buildLog_ += "Unsupported requested OpenCL C version (-cl-std).\n"; return false; } - File* pch = device().compiler()->NewTempFile(DT_CL_HEADER); + File* pch = C->NewTempFile(DT_CL_HEADER); if (pch == NULL || !pch->WriteData((const char*) hdr.first, hdr.second)) { buildLog_ += "Error while opening the opencl-c header "; return false; } - compileOptions_.append(" -Xclang -include-pch -Xclang " + pch->Name()); - compileOptions_.append(" -Xclang -fno-validate-pch"); + driverOptions.append(" -Xclang -include-pch -Xclang " + pch->Name()); + driverOptions.append(" -Xclang -fno-validate-pch"); - compileOptions_.append(hsailOptions()); + driverOptions.append(hsailOptions(options)); if (clVer >= 200) { std::stringstream opts; //Add only for CL2.0 and later opts << " -D" << "CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE=" << device().info().maxGlobalVariableSize_; - compileOptions_.append(opts.str()); + driverOptions.append(opts.str()); } - Buffer* output = device().compiler()->NewBuffer(DT_LLVM_BC); + Buffer* output = C->NewBuffer(DT_LLVM_BC); if (output == NULL) { buildLog_ += "Error while creating buffer for the LLVM bitcode"; return false; } // Tokenize the options string into a vector of strings - std::istringstream strstr(compileOptions_); + std::istringstream strstr(driverOptions); std::istream_iterator sit(strstr), end; - std::vector optionsvec(sit, end); + std::vector params(sit, end); // Compile source to IR - bool ret = device().compiler()->CompileToLLVMBitcode(inputs, output, optionsvec); - buildLog_ += device().compiler()->Output().c_str(); + bool ret = C->CompileToLLVMBitcode(inputs, output, params); + buildLog_ += C->Output(); if (!ret) { buildLog_ += "Error while compiling opencl source: Compiling CL to IR"; return false; } - // save the source code - //Create Binary - codeObjBinary_ = new CodeObjBinary(); + llvmBinary_.assign(output->Buf().data(), output->Size()); + elfSectionType_ = amd::OclElf::LLVMIR; - openCLSource_ = sourceCode; - codeObjBinary_->saveIR(std::string(output->Buf().begin(), output->Buf().end())); + if (clBinary()->saveSOURCE()) { + clBinary()->elfOut()->addSection( + amd::OclElf::SOURCE, sourceCode.data(), sourceCode.size()); + } + if (clBinary()->saveLLVMIR()) { + clBinary()->elfOut()->addSection( + amd::OclElf::LLVMIR, llvmBinary_.data(), llvmBinary_.size(), false); + // store the original compile options + clBinary()->storeCompileOptions(compileOptions_); + } return true; } #endif // defined(WITH_LIGHTNING_COMPILER) @@ -295,8 +311,8 @@ HSAILProgram::compileImpl(const std::string& sourceCode, } //Compile source to IR - this->compileOptions_.append(hsailOptions()); - + this->compileOptions_.append(hsailOptions(options)); + errorCode = g_complibApi._aclCompile(device().compiler(), binaryElf_, //"-Wf,--support_all_extensions", diff --git a/projects/clr/rocclr/runtime/device/rocm/rockernel.cpp b/projects/clr/rocclr/runtime/device/rocm/rockernel.cpp index af3c27c390..2752b46d9a 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rockernel.cpp +++ b/projects/clr/rocclr/runtime/device/rocm/rockernel.cpp @@ -814,9 +814,7 @@ bool Kernel::init_LC(){ hsa_agent_t hsaDevice = program_->hsaDevice(); // Pull out metadata from the ELF - const CodeObjBinary* codeObj = program_->codeObjBinary(); - const RuntimeMD::Program::Metadata* runtimeMD = codeObj->GetProgramMetadata(); - + const RuntimeMD::Program::Metadata* runtimeMD = program_->metadata(); if (!runtimeMD) { return false; } diff --git a/projects/clr/rocclr/runtime/device/rocm/rocprogram.cpp b/projects/clr/rocclr/runtime/device/rocm/rocprogram.cpp index 9c2c58dd7b..599a72b937 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocprogram.cpp +++ b/projects/clr/rocclr/runtime/device/rocm/rocprogram.cpp @@ -83,16 +83,16 @@ namespace roc { } HSAILProgram::~HSAILProgram() { +#if !defined(WITH_LIGHTNING_COMPILER) acl_error error; // Free the elf binary if (binaryElf_ != NULL) { -#if !defined(WITH_LIGHTNING_COMPILER) error = g_complibApi._aclBinaryFini(binaryElf_); if (error != ACL_SUCCESS) { LogWarning( "Error while destroying the acl binary \n" ); } -#endif // !defined(WITH_LIGHTNING_COMPILER) } +#endif // !defined(WITH_LIGHTNING_COMPILER) // Destroy the executable. if (hsaExecutable_.handle != 0) { hsa_executable_destroy(hsaExecutable_); @@ -108,12 +108,15 @@ namespace roc { destroyBrigModule(); destroyBrigContainer(); releaseClBinary(); - } - HSAILProgram::HSAILProgram(roc::NullDevice& device): device::Program(device), - llvmBinary_(), +#if defined(WITH_LIGHTNING_COMPILER) + delete metadata_; +#endif // defined(WITH_LIGHTNING_COMPILER) + } + + HSAILProgram::HSAILProgram(roc::NullDevice& device) + : Program(device), binaryElf_(NULL), - device_(device), brigModule_(NULL), hsaBrigContainer_(NULL) { @@ -126,12 +129,13 @@ namespace roc { binOpts_.bitness = ELFDATA2LSB; binOpts_.alloc = &::malloc; binOpts_.dealloc = &::free; + hsaProgramHandle_.handle = 0; hsaProgramCodeObject_.handle = 0; hsaExecutable_.handle = 0; #if defined(WITH_LIGHTNING_COMPILER) - codeObjBinary_ = NULL; + metadata_ = NULL; #endif // defined(WITH_LIGHTNING_COMPILER) } @@ -344,7 +348,7 @@ namespace roc { void *mem = const_cast(binary.first); acl_error errorCode; #if defined(WITH_LIGHTNING_COMPILER) - // TODO: FIXME_Wilkin + assert(!"FIXME_lmoriche: deserialize the code object, extract the metadata"); #else // !defined(WITH_LIGHTNING_COMPILER) binaryElf_ = g_complibApi._aclReadFromMem(mem, binary.second, &errorCode); if (errorCode != ACL_SUCCESS) { @@ -423,13 +427,32 @@ namespace roc { 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; + } + + *address = (char*) malloc(size); + if (!*address) { + return HSA_STATUS_ERROR_OUT_OF_RESOURCES; + } + + return HSA_STATUS_SUCCESS; + } + bool HSAILProgram::saveBinaryAndSetType(type_t type) { //Write binary to memory void *rawBinary = NULL; size_t size = 0; #if defined(WITH_LIGHTNING_COMPILER) - rawBinary = codeObjBinary_->Binary(); - size = codeObjBinary_->BinarySize(); + hsa_callback_data_t allocData = {0}; + if (hsa_code_object_serialize(hsaProgramCodeObject_, + allocFunc, allocData, + NULL, &rawBinary, &size) != HSA_STATUS_SUCCESS) { + buildLog_ += "Failed to write binary to memory \n"; + return false; + } #else // !defined(WITH_LIGHTNING_COMPILER) if (g_complibApi._aclWriteToMem(binaryElf_, &rawBinary, &size) != ACL_SUCCESS) { @@ -440,19 +463,110 @@ namespace roc { clBinary()->saveBIFBinary((char*)rawBinary, size); //Set the type of binary setType(type); -#if !defined(WITH_LIGHTNING_COMPILER) //Free memory containing rawBinary +#if !defined(WITH_LIGHTNING_COMPILER) binaryElf_->binOpts.dealloc(rawBinary); -#endif // !defined(WITH_LIGHTNING_COMPILER) +#else // defined(WITH_LIGHTNING_COMPILER) + free(rawBinary); +#endif // defined(WITH_LIGHTNING_COMPILER) return true; } +#if defined(WITH_LIGHTNING_COMPILER) + bool HSAILProgram::linkImpl_LC( + const std::vector &inputPrograms, + amd::option::Options *options, + bool createLibrary) + { + using namespace amd::opencl_driver; + Compiler* C = device().compiler(); + + std::vector inputs; + for (auto program : (const std::vector&)inputPrograms) { + if (program->llvmBinary_.empty()) { + if (program->clBinary() == NULL) { + buildLog_ += "Internal error: Input program not compiled!\n"; + return false; + } + + // 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)) { + buildLog_ += "Internal error: Setting input OCL binary failed!\n"; + return false; + } + if (!program->clBinary()->loadLlvmBinary(program->llvmBinary_, + program->elfSectionType_)) { + buildLog_ += "Internal error: Failed loading compiled binary!\n"; + return false; + } + } + + if (program->elfSectionType_ != amd::OclElf::LLVMIR) { + buildLog_ += "Error: Input binary format is not supported\n."; + return false; + } + + Data* input = C->NewBufferReference(DT_LLVM_BC, + (const char*) program->llvmBinary_.data(), + program->llvmBinary_.size()); + + if (!input) { + buildLog_ += "Internal error: Failed to open the compiled programs.\n"; + return false; + } + + inputs.push_back(input); + } + + // open the linked output + Buffer* output = C->NewBuffer(DT_LLVM_BC); + + if (!output) { + buildLog_ += "Error: Failed to open the linked program.\n"; + return false; + } + + std::vector linkOptions; + bool ret = C->LinkLLVMBitcode(inputs, output, linkOptions); + buildLog_ += C->Output(); + if (!ret) { + buildLog_ += "Error: Linking bitcode failed: linking source & IR libraries.\n"; + return false; + } + + llvmBinary_.assign(output->Buf().data(), output->Size()); + elfSectionType_ = amd::OclElf::LLVMIR; + + if (clBinary()->saveLLVMIR()) { + clBinary()->elfOut()->addSection( + amd::OclElf::LLVMIR, llvmBinary_.data(), llvmBinary_.size(), false); + // store the original link options + clBinary()->storeLinkOptions(linkOptions_); + // store the original compile options + clBinary()->storeCompileOptions(compileOptions_); + } + + // skip the rest if we are building an opencl library + if (createLibrary) { + setType(TYPE_LIBRARY); + if (!createBinary(options)) { + buildLog_ += "Internal error: creating OpenCL binary failed\n"; + return false; + } + return true; + } + + return linkImpl_LC(options); + } +#endif // defined(WITH_LIGHTNING_COMPILER) + bool HSAILProgram::linkImpl(const std::vector &inputPrograms, amd::option::Options *options, bool createLibrary) { #if defined(WITH_LIGHTNING_COMPILER) - assert(!"FIXME_Wilkin"); - return false; + return linkImpl_LC(inputPrograms, options, createLibrary); #else // !defined(WITH_LIGHTNING_COMPILER) std::vector::const_iterator it = inputPrograms.begin(); @@ -607,14 +721,14 @@ namespace roc { bool HSAILProgram::linkImpl_LC(amd::option::Options *options) { using namespace amd::opencl_driver; + Compiler* C = device().compiler(); // call LinkLLVMBitcode std::vector inputs; // open the input IR source - const std::string llvmIR = codeObjBinary_->getLlvmIR(); - Data* input = device().compiler()->NewBufferReference( - DT_LLVM_BC, llvmIR.c_str(), llvmIR.length()); + Data* input = C->NewBufferReference( + DT_LLVM_BC, llvmBinary_.data(), llvmBinary_.size()); if (!input) { buildLog_ += "Error: Failed to open the compiled program.\n"; @@ -624,13 +738,13 @@ namespace roc { inputs.push_back(input); //< must be the first input // open the bitcode libraries - Data* opencl_bc = device().compiler()->NewBufferReference(DT_LLVM_BC, + Data* opencl_bc = C->NewBufferReference(DT_LLVM_BC, (const char*) builtins_opencl_amdgcn, builtins_opencl_amdgcn_size); - Data* ocml_bc = device().compiler()->NewBufferReference(DT_LLVM_BC, + Data* ocml_bc = C->NewBufferReference(DT_LLVM_BC, (const char*) builtins_ocml_amdgcn, builtins_ocml_amdgcn_size); - Data* ockl_bc = device().compiler()->NewBufferReference(DT_LLVM_BC, + Data* ockl_bc = C->NewBufferReference(DT_LLVM_BC, (const char*) builtins_ockl_amdgcn, builtins_ockl_amdgcn_size); - Data* irif_bc = device().compiler()->NewBufferReference(DT_LLVM_BC, + Data* irif_bc = C->NewBufferReference(DT_LLVM_BC, (const char*) builtins_irif_amdgcn, builtins_irif_amdgcn_size); if (!opencl_bc || !ocml_bc || !ockl_bc || !irif_bc) { @@ -655,7 +769,7 @@ namespace roc { default: buildLog_ += "Error: Linking for this device is not supported\n"; return false; } - Data* isa_version_bc = device().compiler()->NewBufferReference(DT_LLVM_BC, + Data* isa_version_bc = C->NewBufferReference(DT_LLVM_BC, (const char*) isa_version.first, isa_version.second); if (!isa_version_bc) { @@ -684,13 +798,13 @@ namespace roc { ? std::make_pair(unsafe_math_on_amdgcn, unsafe_math_on_amdgcn_size) : std::make_pair(unsafe_math_off_amdgcn, unsafe_math_off_amdgcn_size); - Data* correctly_rounded_sqrt_bc = device().compiler()->NewBufferReference(DT_LLVM_BC, + Data* correctly_rounded_sqrt_bc = C->NewBufferReference(DT_LLVM_BC, (const char*) correctly_rounded_sqrt.first, correctly_rounded_sqrt.second); - Data* daz_opt_bc = device().compiler()->NewBufferReference(DT_LLVM_BC, + Data* daz_opt_bc = C->NewBufferReference(DT_LLVM_BC, (const char*) daz_opt.first, daz_opt.second); - Data* finite_only_bc = device().compiler()->NewBufferReference(DT_LLVM_BC, + Data* finite_only_bc = C->NewBufferReference(DT_LLVM_BC, (const char*) finite_only.first, finite_only.second); - Data* unsafe_math_bc = device().compiler()->NewBufferReference(DT_LLVM_BC, + Data* unsafe_math_bc = C->NewBufferReference(DT_LLVM_BC, (const char*) unsafe_math.first, unsafe_math.second); if (!correctly_rounded_sqrt_bc || !daz_opt_bc || !finite_only_bc || !unsafe_math_bc) { @@ -705,15 +819,15 @@ namespace roc { // open the linked output std::vector linkOptions; - Data* linked_bc = device().compiler()->NewBuffer(DT_LLVM_BC); + Data* linked_bc = C->NewBuffer(DT_LLVM_BC); if (!linked_bc) { buildLog_ += "Error: Failed to open the linked program.\n"; return false; } - bool ret = device().compiler()->LinkLLVMBitcode(inputs, linked_bc, linkOptions); - buildLog_ += device().compiler()->Output().c_str(); + bool ret = C->LinkLLVMBitcode(inputs, linked_bc, linkOptions); + buildLog_ += C->Output(); if (!ret) { buildLog_ += "Error: Linking bitcode failed: linking source & IR libraries.\n"; return false; @@ -722,13 +836,13 @@ namespace roc { inputs.clear(); inputs.push_back(linked_bc); - Buffer* out_exec = device().compiler()->NewBuffer(DT_EXECUTABLE); + Buffer* out_exec = C->NewBuffer(DT_EXECUTABLE); if (!out_exec) { buildLog_ += "Error: Failed to create the linked executable.\n"; return false; } - std::string optionsstr = options->origOptionStr + hsailOptions(); + std::string optionsstr = options->origOptionStr + hsailOptions(options); // Set the machine target optionsstr.append(" -mcpu="); @@ -743,9 +857,8 @@ namespace roc { std::istream_iterator sit(strstr), end; std::vector optionsvec(sit, end); - ret = device().compiler()->CompileAndLinkExecutable( - inputs, out_exec, optionsvec); - buildLog_ += device().compiler()->Output().c_str(); + ret = C->CompileAndLinkExecutable(inputs, out_exec, optionsvec); + buildLog_ += C->Output(); if (!ret) { buildLog_ += "Error: Creating the executable failed: Compiling LLVM IRs to exe.\n"; return false; @@ -792,12 +905,23 @@ namespace roc { return false; } - //TODO: WC - use the proper target code based on the agent - std::string target = "AMD:AMDGPU:8:0:3"; - codeObjBinary_->init( target, out_exec->Buf().data(), out_exec->Size()); - saveBinaryAndSetType(TYPE_EXECUTABLE); + // load the runtime metadata + amd::OclElf elf(ELFCLASS64, out_exec->Buf().data(), out_exec->Size(), NULL, ELF_C_READ); - buildLog_ += device().compiler()->Output(); + char* data; + size_t size; + if (!elf.getSection(amd::OclElf::RUNTIME_METADATA, &data, &size)) { + buildLog_ += "Error while access runtime metadata.\n"; + return false; + } + + metadata_ = new roc::RuntimeMD::Program::Metadata(); + if (!metadata_->ReadFrom((void *) data, size)) { + buildLog_ += "Error while parsing runtime metadata.\n"; + return false; + } + + saveBinaryAndSetType(TYPE_EXECUTABLE); // Get the list of kernels std::vector kernelNameList; @@ -912,12 +1036,12 @@ namespace roc { 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 defined(WITH_LIGHTNING_COMPILER) - if (!codeObjBinary_) -#else // !defined(WITH_LIGHTNING_COMPILER) if (!binaryElf_) -#endif // !defined(WITH_LIGHTNING_COMPILER) +#else // defined(WITH_LIGHTNING_COMPILER) + if (llvmBinary_.empty()) +#endif // defined(WITH_LIGHTNING_COMPILER) { continueCompileFrom = getNextCompilationStageFromBinary(options); } @@ -939,7 +1063,7 @@ namespace roc { return false; } #else // !defined(WITH_LIGHTNING_COMPILER) - std::string curOptions = options->origOptionStr + hsailOptions(); + std::string curOptions = options->origOptionStr + hsailOptions(options); errorCode = g_complibApi._aclCompile(device().compiler(), binaryElf_, curOptions.c_str(), continueCompileFrom, ACL_TYPE_CG, logFunction); buildLog_ += g_complibApi._aclGetCompilerLog(device().compiler()); @@ -1198,7 +1322,15 @@ namespace roc { } 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() { @@ -1218,13 +1350,13 @@ namespace roc { } } - std::string HSAILProgram::hsailOptions() { + std::string HSAILProgram::hsailOptions(amd::option::Options* options) { std::string hsailOptions; //Set options for the standard device specific options //This is just for legacy compiler code // All our devices support these options now - hsailOptions.append(" -DFP_FAST_FMAF=1"); - hsailOptions.append(" -DFP_FAST_FMA=1"); + hsailOptions.append(" -DFP_FAST_FMAF"); + hsailOptions.append(" -DFP_FAST_FMA"); if (dev().deviceInfo().gfxipVersion_ < 900) { hsailOptions.append(" -cl-denorms-are-zero"); @@ -1240,57 +1372,18 @@ namespace roc { iss.str(device().info().extensions_); while (getline(iss, token, ' ')) { if (!token.empty()) { +#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? + if (options->oVariables->CLStd[2] >= '2' + && token == "cl_khr_depth_images") continue; +#endif // defined(WITH_LIGHTHNING_COMPILER) hsailOptions.append(" -D"); hsailOptions.append(token); - hsailOptions.append("=1"); } } return hsailOptions; } - -#if defined(WITH_LIGHTNING_COMPILER) - void CodeObjBinary::init(std::string& target, void* binary, size_t binarySize) - { - target_ = target; - binary_ = binary; - binarySize_ = binarySize; - - oclElf_ = new amd::OclElf(ELFCLASS64, (char *)binary_, binarySize_, NULL, ELF_C_READ); - - // load the runtime metadata - runtimeMD_ = new roc::RuntimeMD::Program::Metadata(); - } - - void CodeObjBinary::fini() - { - if (oclElf_) { - delete oclElf_; - } - - if (runtimeMD_) { - delete runtimeMD_; - } - - target_ = ""; - binary_ = NULL; - binarySize_ = 0; - } - - const RuntimeMD::Program::Metadata* CodeObjBinary::GetProgramMetadata() const - { - char* metaData; - size_t metaSize; - if (!oclElf_->getSection(amd::OclElf::RUNTIME_METADATA, &metaData, &metaSize)) { - LogWarning( "Error while access runtime metadata section from the binary \n" ); - } - - if (!runtimeMD_->ReadFrom((void *) metaData, metaSize)) { - LogWarning( "Error while parsing runtime metadata \n" ); - } - - return runtimeMD_; - } -#endif // defined(WITH_LIGHTNING_COMPILER) #endif // WITHOUT_HSA_BACKEND } // namespace roc diff --git a/projects/clr/rocclr/runtime/device/rocm/rocprogram.hpp b/projects/clr/rocclr/runtime/device/rocm/rocprogram.hpp index 6280637a17..17520d4c3c 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocprogram.hpp +++ b/projects/clr/rocclr/runtime/device/rocm/rocprogram.hpp @@ -26,56 +26,6 @@ using namespace HSAIL_ASM; //! \namespace roc HSA Device Implementation namespace roc { -#if defined(WITH_LIGHTNING_COMPILER) - class CodeObjBinary { - public: - CodeObjBinary() - : target_(""), kernelArgAlign_(0), capFlags_(0), encryptCode_(0), - binary_(NULL), binarySize_(0), llvmIR_(""), oclElf_(NULL), runtimeMD_(NULL) {} - - void init(std::string& target, void* binary, size_t binarySize); - void fini(); - - std::string Target() const { return target_; } - uint32_t KernelArgAlign() const { return kernelArgAlign_; } - void* Binary() const { return binary_; } - size_t BinarySize() const { return binarySize_; } - - void saveIR(std::string llvmIR) { llvmIR_ = llvmIR; } - std::string getLlvmIR() const { return llvmIR_; } - - amd::OclElf* oclElf() const { return oclElf_; } - RuntimeMD::Program::Metadata* runtimeMD() const { return runtimeMD_; } - - const RuntimeMD::Program::Metadata* GetProgramMetadata() const; - - private: - enum CapFlag { - capSaveSource = 0, - capSaveLLVMIR = 1, - capSaveCG = 2, - capSaveEXE = 3, - capSaveHSAIL = 4, - capSaveISASM = 5, - capEncryted = 6 - }; - - std::string target_; // target device - uint32_t kernelArgAlign_; - uint32_t capFlags_; - uint32_t encryptCode_; - - void * binary_; //!< code object binary (ISA) - size_t binarySize_; //!< size of the code object binary - - std::string llvmIR_; //!< LLVM IR binary code - - amd::OclElf* oclElf_; //!< ELF object to access runtime metadata - - roc::RuntimeMD::Program::Metadata* runtimeMD_; //!< runtime metadata - }; -#endif // defined(WITH_LIGHTNING_COMPILER) - //! \class empty program class HSAILProgram : public device::Program { @@ -89,22 +39,20 @@ namespace roc { // Initialize Binary for GPU (used only for clCreateProgramWithBinary()). virtual bool initClBinary(char *binaryIn, size_t size); - //! Returns the aclBinary associated with the progrm + //! Returns the aclBinary associated with the program const aclBinary* binaryElf() const { return static_cast(binaryElf_); } #if defined(WITH_LIGHTNING_COMPILER) - //! Returns the code object binary associated with the progrm - const CodeObjBinary* codeObjBinary() const { //! Binary for the code object - return static_cast(codeObjBinary_); } + //! Returns the program metadata. + const RuntimeMD::Program::Metadata* metadata() const { return metadata_; } #endif // defined(WITH_LIGHTNING_COMPILER) - const std::string& HsailText() { - return hsailProgram_; - } + //! Return a typecasted GPU device + const NullDevice& dev() const + { return static_cast(device()); } - const NullDevice& dev() const { return device_; } - //! Returns the hsaBinary associated with the progrm + //! Returns the hsaBinary associated with the program hsa_agent_t hsaDevice() const { return dev().getBackendDevice(); } @@ -118,7 +66,7 @@ namespace roc { /*! \brief Compiles GPU CL program to LLVM binary (compiler frontend) * - * \return True if we successefully compiled a GPU program + * \return True if we successfully compiled a GPU program */ virtual bool compileImpl( const std::string& sourceCode, //!< the program's source code @@ -153,6 +101,11 @@ namespace roc { 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); @@ -204,15 +157,12 @@ namespace roc { //! Returns all the options to be appended while passing to the //compiler - std::string hsailOptions(); + std::string hsailOptions(amd::option::Options* options); + + // aclBinary and aclCompiler - for the compiler library + aclBinary* binaryElf_; //!< Binary for the new compiler library + aclBinaryOptions binOpts_; //!< Binary options to create aclBinary - std::string openCLSource_; //!< Original OpenCL source - std::string hsailProgram_; //!< HSAIL program after compilation. - std::string llvmBinary_; //!< LLVM IR binary code - //!< aclBinary and aclCompiler - for the compiler libray - aclBinary* binaryElf_; //!