From 8bbcc1c38a3b99bec5299d361e5f23130ccb8c57 Mon Sep 17 00:00:00 2001 From: foreman Date: Thu, 27 Sep 2018 18:15:24 -0400 Subject: [PATCH] P4 to Git Change 1611775 by gandryey@gera-ocl-lc on 2018/09/27 18:02:54 SWDEV-79445 - OCL generic changes and code clean-up Program compilation clean-up. Step#6: - Move the second linkImpl() method to the abstraciton layer - Create the new setKernel virtual interface for the backend specific setup Affected files ... ... //depot/stg/opencl/drivers/opencl/runtime/device/device.hpp#320 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/devprogram.cpp#8 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/devprogram.hpp#6 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/paldefs.hpp#41 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palprogram.cpp#79 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palprogram.hpp#35 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palsettings.cpp#57 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocprogram.cpp#91 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocprogram.hpp#40 edit --- rocclr/runtime/device/device.hpp | 3 +- rocclr/runtime/device/devprogram.cpp | 527 ++++++++++++++++++---- rocclr/runtime/device/devprogram.hpp | 39 +- rocclr/runtime/device/pal/paldefs.hpp | 34 +- rocclr/runtime/device/pal/palprogram.cpp | 330 +------------- rocclr/runtime/device/pal/palprogram.hpp | 8 +- rocclr/runtime/device/pal/palsettings.cpp | 2 + rocclr/runtime/device/rocm/rocprogram.cpp | 293 +----------- rocclr/runtime/device/rocm/rocprogram.hpp | 8 +- 9 files changed, 515 insertions(+), 729 deletions(-) diff --git a/rocclr/runtime/device/device.hpp b/rocclr/runtime/device/device.hpp index 36dfe95d8e..b2b82adb50 100644 --- a/rocclr/runtime/device/device.hpp +++ b/rocclr/runtime/device/device.hpp @@ -518,7 +518,8 @@ class Settings : public amd::HeapObject { uint reportFMAF_ : 1; //!< Report FP_FAST_FMAF define in CL program uint reportFMA_ : 1; //!< Report FP_FAST_FMA define in CL program uint singleFpDenorm_ : 1; //!< Support Single FP Denorm - uint reserved_ : 22; + uint gfx10Hsail_ : 1 ; //!< GFX10 HSAIL path + uint reserved_ : 21; }; uint value_; }; diff --git a/rocclr/runtime/device/devprogram.cpp b/rocclr/runtime/device/devprogram.cpp index bf861ae81a..1d8dca35fb 100644 --- a/rocclr/runtime/device/devprogram.cpp +++ b/rocclr/runtime/device/devprogram.cpp @@ -13,9 +13,11 @@ #if defined(WITH_LIGHTNING_COMPILER) #include "driver/AmdCompiler.h" +#include "libraries.amdgcn.inc" #include "opencl1.2-c.amdgcn.inc" #include "opencl2.0-c.amdgcn.inc" #endif // !defined(WITH_LIGHTNING_COMPILER) + #include #include #include @@ -492,97 +494,6 @@ bool Program::linkImpl(const std::vector& inputPrograms, } } -// ================================================================================================ -bool Program::linkImplHSAIL(const std::vector& inputPrograms, - amd::option::Options* options, bool createLibrary) { -#if defined(WITH_COMPILER_LIB) || !defined(WITH_LIGHTNING_COMPILER) - acl_error errorCode; - - // For each program we need to extract the LLVMIR and create - // aclBinary for each - std::vector binaries_to_link; - - for (auto program : inputPrograms) { - // 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) { - bool spirv = false; - size_t boolSize = sizeof(bool); - errorCode = aclQueryInfo( - device().compiler(), binaryElf_, RT_CONTAINS_SPIRV, nullptr, &spirv, &boolSize); - if (errorCode != ACL_SUCCESS) { - spirv = false; - } - if (spirv) { - errorCode = aclCompile( - device().compiler(), binaryElf_, options->origOptionStr.c_str(), - ACL_TYPE_SPIRV_BINARY, ACL_TYPE_LLVMIR_BINARY, nullptr); - buildLog_ += aclGetCompilerLog(device().compiler()); - if (errorCode != ACL_SUCCESS) { - buildLog_ += "Error while linking: Could not load SPIR-V"; - return false; - } - } - else { - 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); - } - - errorCode = aclLink(device().compiler(), binaries_to_link[0], binaries_to_link.size() - 1, - binaries_to_link.size() > 1 ? &binaries_to_link[1] : nullptr, - ACL_TYPE_LLVMIR_BINARY, "-create-library", nullptr); - if (errorCode != ACL_SUCCESS) { - buildLog_ += aclGetCompilerLog(device().compiler()); - buildLog_ += "Error while linking : aclLink failed"; - 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); - buildLog_ += aclGetCompilerLog(device().compiler()); - return true; - } - // Now call linkImpl with the new options - return linkImpl(options); -#else - return false; -#endif // defined(WITH_COMPILER_LIB) || !defined(WITH_LIGHTNING_COMPILER) -} - // ================================================================================================ bool Program::linkImplLC(const std::vector& inputPrograms, amd::option::Options* options, bool createLibrary) { @@ -678,6 +589,440 @@ bool Program::linkImplLC(const std::vector& inputPrograms, #endif // defined(WITH_LIGHTNING_COMPILER) } +// ================================================================================================ +bool Program::linkImplHSAIL(const std::vector& inputPrograms, + amd::option::Options* options, bool createLibrary) { +#if defined(WITH_COMPILER_LIB) || !defined(WITH_LIGHTNING_COMPILER) + acl_error errorCode; + + // For each program we need to extract the LLVMIR and create + // aclBinary for each + std::vector binaries_to_link; + + for (auto program : inputPrograms) { + // 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) { + bool spirv = false; + size_t boolSize = sizeof(bool); + errorCode = aclQueryInfo( + device().compiler(), binaryElf_, RT_CONTAINS_SPIRV, nullptr, &spirv, &boolSize); + if (errorCode != ACL_SUCCESS) { + spirv = false; + } + if (spirv) { + errorCode = aclCompile( + device().compiler(), binaryElf_, options->origOptionStr.c_str(), + ACL_TYPE_SPIRV_BINARY, ACL_TYPE_LLVMIR_BINARY, nullptr); + buildLog_ += aclGetCompilerLog(device().compiler()); + if (errorCode != ACL_SUCCESS) { + buildLog_ += "Error while linking: Could not load SPIR-V"; + return false; + } + } + else { + 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); + } + + errorCode = aclLink(device().compiler(), binaries_to_link[0], binaries_to_link.size() - 1, + binaries_to_link.size() > 1 ? &binaries_to_link[1] : nullptr, + ACL_TYPE_LLVMIR_BINARY, "-create-library", nullptr); + if (errorCode != ACL_SUCCESS) { + buildLog_ += aclGetCompilerLog(device().compiler()); + buildLog_ += "Error while linking : aclLink failed"; + 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); + buildLog_ += aclGetCompilerLog(device().compiler()); + return true; + } + + // Now call linkImpl with the new options + return linkImpl(options); +#else + return false; +#endif // defined(WITH_COMPILER_LIB) || !defined(WITH_LIGHTNING_COMPILER) +} + +// ================================================================================================ +bool Program::linkImpl(amd::option::Options* options) { + if (isLC()) { + return linkImplLC(options); + } + else { + return linkImplHSAIL(options); + } +} + +// ================================================================================================ +bool Program::linkImplLC(amd::option::Options* options) { +#if defined(WITH_LIGHTNING_COMPILER) + using namespace amd::opencl_driver; + internal_ = (compileOptions_.find("-cl-internal-kernel") != std::string::npos) ? + true : false; + std::vector inputs; + std::unique_ptr C(newCompilerInstance()); + bool bLinkLLVMBitcode = true; + aclType continueCompileFrom = llvmBinary_.empty() ? + getNextCompilationStageFromBinary(options) : ACL_TYPE_LLVMIR_BINARY; + + switch (continueCompileFrom) { + case ACL_TYPE_CG: + case ACL_TYPE_LLVMIR_BINARY: { + break; + } + case ACL_TYPE_ASM_TEXT: { + char* section; + size_t sz; + clBinary()->elfOut()->getSection(amd::OclElf::SOURCE, §ion, &sz); + Data* input = C->NewBufferReference(DT_ASSEMBLY, section, sz); + if (!input) { + buildLog_ += "Error: Failed to open the assembler text.\n"; + return false; + } + inputs.push_back(input); + bLinkLLVMBitcode = false; + 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; + } + default: + buildLog_ += "Error while Codegen phase: the binary is incomplete \n"; + return false; + } + + // call LinkLLVMBitcode + if (bLinkLLVMBitcode) { + // open the input IR source + Data* input = C->NewBufferReference(DT_LLVM_BC, llvmBinary_.data(), llvmBinary_.size()); + + if (!input) { + buildLog_ += "Error: Failed to open the compiled program.\n"; + return false; + } + + inputs.push_back(input); // must be the first input + // open the bitcode libraries + Data* opencl_bc = + C->NewBufferReference(DT_LLVM_BC, (const char*)opencl_amdgcn, opencl_amdgcn_size); + Data* ocml_bc = C->NewBufferReference(DT_LLVM_BC, (const char*)ocml_amdgcn, ocml_amdgcn_size); + Data* ockl_bc = C->NewBufferReference(DT_LLVM_BC, (const char*)ockl_amdgcn, ockl_amdgcn_size); + + if (!opencl_bc || !ocml_bc || !ockl_bc) { + buildLog_ += "Error: Failed to open the bitcode library.\n"; + return false; + } + + inputs.push_back(opencl_bc); // depends on oclm & ockl + inputs.push_back(ockl_bc); + inputs.push_back(ocml_bc); + + // open the control functions + auto isa_version = get_oclc_isa_version(device().info().gfxipVersion_); + if (!isa_version.first) { + buildLog_ += "Error: Linking for this device is not supported\n"; + return false; + } + + Data* isa_version_bc = + C->NewBufferReference(DT_LLVM_BC, (const char*)isa_version.first, isa_version.second); + + if (!isa_version_bc) { + buildLog_ += "Error: Failed to open the control functions.\n"; + return false; + } + + inputs.push_back(isa_version_bc); + + auto correctly_rounded_sqrt = + get_oclc_correctly_rounded_sqrt(options->oVariables->FP32RoundDivideSqrt); + Data* correctly_rounded_sqrt_bc = C->NewBufferReference(DT_LLVM_BC, correctly_rounded_sqrt.first, + correctly_rounded_sqrt.second); + + auto daz_opt = get_oclc_daz_opt(options->oVariables->DenormsAreZero || + AMD_GPU_FORCE_SINGLE_FP_DENORM == 0 || + (device().info().gfxipVersion_ < 900 && AMD_GPU_FORCE_SINGLE_FP_DENORM < 0)); + Data* daz_opt_bc = C->NewBufferReference(DT_LLVM_BC, daz_opt.first, daz_opt.second); + + auto finite_only = get_oclc_finite_only(options->oVariables->FiniteMathOnly || + options->oVariables->FastRelaxedMath); + Data* finite_only_bc = C->NewBufferReference(DT_LLVM_BC, finite_only.first, finite_only.second); + + auto unsafe_math = get_oclc_unsafe_math(options->oVariables->UnsafeMathOpt || + options->oVariables->FastRelaxedMath); + Data* unsafe_math_bc = C->NewBufferReference(DT_LLVM_BC, unsafe_math.first, unsafe_math.second); + + if (!correctly_rounded_sqrt_bc || !daz_opt_bc || !finite_only_bc || !unsafe_math_bc) { + buildLog_ += "Error: Failed to open the control functions.\n"; + return false; + } + + inputs.push_back(correctly_rounded_sqrt_bc); + inputs.push_back(daz_opt_bc); + inputs.push_back(finite_only_bc); + inputs.push_back(unsafe_math_bc); + + // open the linked output + std::vector linkOptions; + Buffer* linked_bc = C->NewBuffer(DT_LLVM_BC); + + if (!linked_bc) { + buildLog_ += "Error: Failed to open the linked program.\n"; + return false; + } + + // NOTE: The linkOptions parameter is also used to identy cached code object. This parameter + // should not contain any dyanamically generated filename. + bool ret = device().cacheCompilation()->linkLLVMBitcode( + C.get(), inputs, linked_bc, linkOptions, buildLog_); + buildLog_ += C->Output(); + if (!ret) { + buildLog_ += "Error: Linking bitcode failed: linking source & IR libraries.\n"; + return false; + } + + if (options->isDumpFlagSet(amd::option::DUMP_BC_LINKED)) { + std::ofstream f(options->getDumpFileName("_linked.bc").c_str(), + std::ios::binary | std::ios::trunc); + if (f.is_open()) { + f.write(linked_bc->Buf().data(), linked_bc->Size()); + f.close(); + } + else { + buildLog_ += "Warning: opening the file to dump the linked IR failed.\n"; + } + } + + inputs.clear(); + inputs.push_back(linked_bc); + } + + Buffer* out_exec = C->NewBuffer(DT_EXECUTABLE); + if (!out_exec) { + buildLog_ += "Error: Failed to create the linked executable.\n"; + return false; + } + + std::string codegenOptions(options->llvmOptions); + + // Set the machine target + codegenOptions.append(" -mcpu="); + codegenOptions.append(machineTarget_); + + // Set xnack option if needed + if (xnackEnabled_) { + codegenOptions.append(" -mxnack"); + } + + // Set the -O# + std::ostringstream optLevel; + optLevel << "-O" << options->oVariables->OptLevel; + codegenOptions.append(" ").append(optLevel.str()); + + // Pass clang options + std::ostringstream ostrstr; + std::copy(options->clangOptions.begin(), options->clangOptions.end(), + std::ostream_iterator(ostrstr, " ")); + codegenOptions.append(" ").append(ostrstr.str()); + + // Set whole program mode + codegenOptions.append(" -mllvm -amdgpu-internalize-symbols -mllvm -amdgpu-early-inline-all"); + + // Tokenize the options string into a vector of strings + std::istringstream strstr(codegenOptions); + std::istream_iterator sit(strstr), end; + std::vector params(sit, end); + + // NOTE: The params is also used to identy cached code object. This parameter + // should not contain any dyanamically generated filename. + bool ret = device().cacheCompilation()->compileAndLinkExecutable(C.get(), inputs, out_exec, params, + buildLog_); + buildLog_ += C->Output(); + if (!ret) { + if (continueCompileFrom == ACL_TYPE_ASM_TEXT) { + buildLog_ += "Error: Creating the executable from ISA assembly text failed.\n"; + } + else { + buildLog_ += "Error: Creating the executable from LLVM IRs failed.\n"; + } + return false; + } + + if (options->isDumpFlagSet(amd::option::DUMP_O)) { + std::ofstream f(options->getDumpFileName(".so").c_str(), std::ios::binary | std::ios::trunc); + if (f.is_open()) { + f.write(out_exec->Buf().data(), out_exec->Size()); + f.close(); + } + else { + buildLog_ += "Warning: opening the file to dump the code object failed.\n"; + } + } + + if (options->isDumpFlagSet(amd::option::DUMP_ISA)) { + std::string name = options->getDumpFileName(".s"); + File* dump = C->NewFile(DT_INTERNAL, name); + if (!C->DumpExecutableAsText(out_exec, dump)) { + buildLog_ += "Warning: failed to dump code object.\n"; + } + } + + // Call the device layer to setup all available kernels on the actual device + if (!setKernels(options, out_exec->Buf().data(), out_exec->Size())) { + return false; + } + + // Save the binary and type + clBinary()->saveBIFBinary(reinterpret_cast(out_exec->Buf().data()), out_exec->Size()); + setType(TYPE_EXECUTABLE); + + return true; +#else + return false; +#endif // defined(WITH_LIGHTNING_COMPILER) +} + +// ================================================================================================ +bool Program::linkImplHSAIL(amd::option::Options* options) { +#if defined(WITH_COMPILER_LIB) || !defined(WITH_LIGHTNING_COMPILER) + acl_error errorCode; + bool finalize = true; + internal_ = (compileOptions_.find("-cl-internal-kernel") != std::string::npos) ? true : false; + // If !binaryElf_ then program must have been created using clCreateProgramWithBinary + aclType continueCompileFrom = (!binaryElf_) ? + getNextCompilationStageFromBinary(options) : ACL_TYPE_LLVMIR_BINARY; + + switch (continueCompileFrom) { + case ACL_TYPE_SPIRV_BINARY: + case ACL_TYPE_SPIR_BINARY: + // 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 + ProcessOptions(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; + } + + if (finalize) { + std::string fin_options(options->origOptionStr + ProcessOptions(options)); + // Append an option so that we can selectively enable a SCOption on CZ + // whenever IOMMUv2 is enabled. + if (device().isFineGrainedSystem(true)) { + fin_options.append(" -sc-xnack-iommu"); + } + if (device().settings().gfx10Hsail_) { + if (GPU_FORCE_WAVE_SIZE_32) { + fin_options.append(" -force-wave-size-32"); + } + if (xnackEnabled_) { + fin_options.append(" -xnack"); + } + } + + 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 binSize; + void* binary = const_cast(aclExtractSection( + device().compiler(), binaryElf_, &binSize, aclTEXT, &errorCode)); + if (errorCode != ACL_SUCCESS) { + buildLog_ += "Error: cannot extract ISA from compiled binary.\n"; + return false; + } + + // Call the device layer to setup all available kernels on the actual device + if (!setKernels(options, binary, binSize)) { + return false; + } + + // Save the binary in the interface class + saveBinaryAndSetType(TYPE_EXECUTABLE); + buildLog_ += aclGetCompilerLog(device().compiler()); + + return true; +#else + return false; +#endif // defined(WITH_COMPILER_LIB) || !defined(WITH_LIGHTNING_COMPILER) +} + // ================================================================================================ bool Program::initClBinary() { if (clBinary_ == nullptr) { diff --git a/rocclr/runtime/device/devprogram.hpp b/rocclr/runtime/device/devprogram.hpp index 2a9f1f1d67..179b56f720 100644 --- a/rocclr/runtime/device/devprogram.hpp +++ b/rocclr/runtime/device/devprogram.hpp @@ -214,7 +214,7 @@ class Program : public amd::HeapObject { ); //! Link the device program. - virtual bool linkImpl(amd::option::Options* options) = 0; + virtual bool linkImpl(amd::option::Options* options); //! Link the device programs. virtual bool linkImpl(const std::vector& inputPrograms, amd::option::Options* options, @@ -238,6 +238,9 @@ class Program : public amd::HeapObject { virtual bool isElf(const char* bin) const = 0; + virtual bool setKernels( + amd::option::Options* options, void* binary, size_t binSize) { return true; } + //! Returns all the options to be appended while passing to the compiler library std::string ProcessOptions(amd::option::Options* options); @@ -269,23 +272,29 @@ class Program : public amd::HeapObject { bool FindGlobalVarSize(void* binary, size_t binSize); private: - //! Compile the device program with LC path - bool compileImplLC(const std::string& sourceCode, - const std::vector& headers, - const char** headerIncludeNames, amd::option::Options* options); + //! Compile the device program with LC path + bool compileImplLC(const std::string& sourceCode, + const std::vector& headers, + const char** headerIncludeNames, amd::option::Options* options); - //! Compile the device program with HSAIL path - bool compileImplHSAIL(const std::string& sourceCode, - const std::vector& headers, - const char** headerIncludeNames, amd::option::Options* options); + //! Compile the device program with HSAIL path + bool compileImplHSAIL(const std::string& sourceCode, + const std::vector& headers, + const char** headerIncludeNames, amd::option::Options* options); - //! Link the device programs with HSAIL path - bool linkImplHSAIL(const std::vector& inputPrograms, - amd::option::Options* options, bool createLibrary); + //! Link the device programs with LC path + bool linkImplLC(const std::vector& inputPrograms, + amd::option::Options* options, bool createLibrary); - //! Link the device programs with LC path - bool linkImplLC(const std::vector& inputPrograms, - amd::option::Options* options, bool createLibrary); + //! Link the device programs with HSAIL path + bool linkImplHSAIL(const std::vector& inputPrograms, + amd::option::Options* options, bool createLibrary); + + //! Link the device program with LC path + bool linkImplLC(amd::option::Options* options); + + //! Link the device program with HSAIL path + bool linkImplHSAIL(amd::option::Options* options); //! Disable default copy constructor Program(const Program&); diff --git a/rocclr/runtime/device/pal/paldefs.hpp b/rocclr/runtime/device/pal/paldefs.hpp index 5cfecc2721..5ed6752712 100644 --- a/rocclr/runtime/device/pal/paldefs.hpp +++ b/rocclr/runtime/device/pal/paldefs.hpp @@ -128,26 +128,26 @@ struct AMDDeviceInfo { }; static const AMDDeviceInfo DeviceInfo[] = { - /* Unknown */ {"", "unknown", "", 4, 16, 1, 256, 64 * Ki, 32, 0, false}, - /* Tahiti */ {"", "tahiti", "", 4, 16, 1, 256, 64 * Ki, 32, 600, false}, - /* Pitcairn */ {"", "pitcairn", "", 4, 16, 1, 256, 64 * Ki, 32, 600, false}, - /* Capeverde */ {"", "bonaire", "", 4, 16, 1, 256, 64 * Ki, 32, 700, false}, - /* Oland */ {"", "oland", "", 4, 16, 1, 256, 64 * Ki, 32, 600, false}, - /* Hainan */ {"", "hainan", "", 4, 16, 1, 256, 64 * Ki, 32, 600, false}, + /* Unknown */ {"", "unknown", "", 4, 16, 1, 256, 64 * Ki, 32, 0, 0, false}, + /* Tahiti */ {"", "tahiti", "", 4, 16, 1, 256, 64 * Ki, 32, 600, 600, false}, + /* Pitcairn */ {"", "pitcairn", "", 4, 16, 1, 256, 64 * Ki, 32, 600, 600, false}, + /* Capeverde */ {"", "bonaire", "", 4, 16, 1, 256, 64 * Ki, 32, 700, 700, false}, + /* Oland */ {"", "oland", "", 4, 16, 1, 256, 64 * Ki, 32, 600, 600, false}, + /* Hainan */ {"", "hainan", "", 4, 16, 1, 256, 64 * Ki, 32, 600, 600, false}, - /* Bonaire */ {"Bonaire", "bonaire", "", 4, 16, 1, 256, 64 * Ki, 32, 700, false}, - /* Hawaii */ {"Hawaii", "hawaii", "", 4, 16, 1, 256, 64 * Ki, 32, 701, false}, - /* Hawaii */ {"", "grenada", "", 4, 16, 1, 256, 64 * Ki, 32, 701, false}, - /* Hawaii */ {"", "maui", "", 4, 16, 1, 256, 64 * Ki, 32, 701, false}, + /* Bonaire */ {"Bonaire", "bonaire", "", 4, 16, 1, 256, 64 * Ki, 32, 700, 700, false}, + /* Hawaii */ {"Hawaii", "hawaii", "", 4, 16, 1, 256, 64 * Ki, 32, 701, 701, false}, + /* Hawaii */ {"", "grenada", "", 4, 16, 1, 256, 64 * Ki, 32, 701, 701, false}, + /* Hawaii */ {"", "maui", "", 4, 16, 1, 256, 64 * Ki, 32, 701, 701, false}, - /* Kalindi */ {"Kalindi", "kalindi", "", 4, 16, 1, 256, 64 * Ki, 32, 702, false}, - /* Godavari */ {"Mullins", "mullins", "", 4, 16, 1, 256, 64 * Ki, 32, 702, false}, - /* Spectre */ {"Spectre", "spectre", "", 4, 16, 1, 256, 64 * Ki, 32, 701, false}, - /* Spooky */ {"Spooky", "spooky", "", 4, 16, 1, 256, 64 * Ki, 32, 701, false}, + /* Kalindi */ {"Kalindi", "kalindi", "", 4, 16, 1, 256, 64 * Ki, 32, 702, 702, false}, + /* Godavari */ {"Mullins", "mullins", "", 4, 16, 1, 256, 64 * Ki, 32, 702, 702, false}, + /* Spectre */ {"Spectre", "spectre", "", 4, 16, 1, 256, 64 * Ki, 32, 701, 701, false}, + /* Spooky */ {"Spooky", "spooky", "", 4, 16, 1, 256, 64 * Ki, 32, 701, 701, false}, - /* Carrizo */ {"Carrizo", "carrizo", "", 4, 16, 1, 256, 64 * Ki, 32, 801, false}, - /* Bristol */ {"Bristol Ridge", "carrizo", "", 4, 16, 1, 256, 64 * Ki, 32, 801, false}, - /* Stoney */ {"Stoney", "stoney", "", 4, 16, 1, 256, 64 * Ki, 32, 810, false}, + /* Carrizo */ {"Carrizo", "carrizo", "", 4, 16, 1, 256, 64 * Ki, 32, 801, 801,false}, + /* Bristol */ {"Bristol Ridge", "carrizo", "", 4, 16, 1, 256, 64 * Ki, 32, 801, 801, false}, + /* Stoney */ {"Stoney", "stoney", "", 4, 16, 1, 256, 64 * Ki, 32, 810, 810, false}, /* Iceland */ {"Iceland", "iceland", "gfx802", 4, 16, 1, 256, 64 * Ki, 32, 802, 800, false}, /* Tonga */ {"Tonga", "tonga", "gfx802", 4, 16, 1, 256, 64 * Ki, 32, 802, 800, false}, diff --git a/rocclr/runtime/device/pal/palprogram.cpp b/rocclr/runtime/device/pal/palprogram.cpp index cc976a521e..024766be4c 100644 --- a/rocclr/runtime/device/pal/palprogram.cpp +++ b/rocclr/runtime/device/pal/palprogram.cpp @@ -244,113 +244,37 @@ inline static std::vector splitSpaceSeparatedString(char* str) { return vec; } -bool HSAILProgram::linkImpl(amd::option::Options* options) { +bool HSAILProgram::setKernels(amd::option::Options* options, void* binary, size_t binSize) { #if defined(WITH_LIGHTNING_COMPILER) assert(!"Should not reach here"); return false; #else // !defined(WITH_LIGHTNING_COMPILER) - acl_error errorCode; - aclType continueCompileFrom = ACL_TYPE_LLVMIR_BINARY; - bool finalize = true; - bool hsaLoad = true; - internal_ = (compileOptions_.find("-cl-internal-kernel") != std::string::npos) ? true : false; - - - // If !binaryElf_ then program must have been created using clCreateProgramWithBinary - if (!binaryElf_) { - continueCompileFrom = getNextCompilationStageFromBinary(options); - } - switch (continueCompileFrom) { - case ACL_TYPE_SPIRV_BINARY: - case ACL_TYPE_SPIR_BINARY: - // 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 + ProcessOptions(options); - errorCode = aclCompile(dev().compiler(), binaryElf_, curOptions.c_str(), continueCompileFrom, - ACL_TYPE_CG, nullptr); - buildLog_ += aclGetCompilerLog(dev().compiler()); - if (errorCode != ACL_SUCCESS) { - buildLog_ += "Error: BRIG code generation failed.\n"; - return false; - } - break; - } - case ACL_TYPE_CG: - break; - case ACL_TYPE_ISA: - finalize = false; - break; - default: - buildLog_ += - "Error: The binary is incorrect or incomplete. Finalization to ISA couldn't be " - "performed.\n"; - return false; - } - if (finalize) { - std::string fin_options(options->origOptionStr + ProcessOptions(options)); - // Append an option so that we can selectively enable a SCOption on CZ - // whenever IOMMUv2 is enabled. - if (dev().settings().svmFineGrainSystem_) { - fin_options.append(" -sc-xnack-iommu"); - } - if (dev().settings().gfx10Plus_) { - if (GPU_FORCE_WAVE_SIZE_32) { - fin_options.append(" -force-wave-size-32"); - } - if (dev().hwInfo()->xnackEnabled_) { - fin_options.append(" -xnack"); - } - } - - errorCode = aclCompile(dev().compiler(), binaryElf_, fin_options.c_str(), ACL_TYPE_CG, - ACL_TYPE_ISA, nullptr); - buildLog_ += aclGetCompilerLog(dev().compiler()); - if (errorCode != ACL_SUCCESS) { - buildLog_ += "Error: BRIG finalization to ISA failed.\n"; - return false; - } - } // ACL_TYPE_CG stage is not performed for offline compilation hsa_agent_t agent; agent.handle = 1; - if (hsaLoad) { - executable_ = loader_->CreateExecutable(HSA_PROFILE_FULL, NULL); - if (executable_ == nullptr) { - buildLog_ += "Error: Executable for AMD HSA Code Object isn't created.\n"; - return false; - } - size_t size = 0; - hsa_code_object_t code_object; - code_object.handle = reinterpret_cast( - aclExtractSection(dev().compiler(), binaryElf_, &size, aclTEXT, &errorCode)); - if (errorCode != ACL_SUCCESS) { - buildLog_ += "Error: Extracting AMD HSA Code Object from binary failed.\n"; - return false; - } - hsa_status_t status = executable_->LoadCodeObject(agent, code_object, nullptr); - if (status != HSA_STATUS_SUCCESS) { - buildLog_ += "Error: AMD HSA Code Object loading failed.\n"; - return false; - } - status = executable_->Freeze(nullptr); - if (status != HSA_STATUS_SUCCESS) { - buildLog_ += "Error: AMD HSA Code Object freeze failed.\n"; - return false; - } + executable_ = loader_->CreateExecutable(HSA_PROFILE_FULL, nullptr); + if (executable_ == nullptr) { + buildLog_ += "Error: Executable for AMD HSA Code Object isn't created.\n"; + return false; } + size_t size = binSize; + hsa_code_object_t code_object; + code_object.handle = reinterpret_cast(binary); + + hsa_status_t status = executable_->LoadCodeObject(agent, code_object, nullptr); + if (status != HSA_STATUS_SUCCESS) { + buildLog_ += "Error: AMD HSA Code Object loading failed.\n"; + return false; + } + status = executable_->Freeze(nullptr); + if (status != HSA_STATUS_SUCCESS) { + buildLog_ += "Error: AMD HSA Code Object freeze failed.\n"; + return false; + } + size_t kernelNamesSize = 0; - errorCode = aclQueryInfo(dev().compiler(), binaryElf_, RT_KERNEL_NAMES, nullptr, nullptr, - &kernelNamesSize); + acl_error errorCode = aclQueryInfo(dev().compiler(), binaryElf_, RT_KERNEL_NAMES, + nullptr, nullptr, &kernelNamesSize); if (errorCode != ACL_SUCCESS) { buildLog_ += "Error: Querying of kernel names size from the binary failed.\n"; return false; @@ -400,10 +324,6 @@ bool HSAILProgram::linkImpl(amd::option::Options* options) { } DestroySegmentCpuAccess(); - - // Save the binary in the interface class - saveBinaryAndSetType(TYPE_EXECUTABLE); - buildLog_ += aclGetCompilerLog(dev().compiler()); return true; #endif // !defined(WITH_LIGHTNING_COMPILER) } @@ -736,205 +656,7 @@ bool LightningProgram::createBinary(amd::option::Options* options) { return true; } -bool LightningProgram::linkImpl(amd::option::Options* options) { - using namespace amd::opencl_driver; - internal_ = (compileOptions_.find("-cl-internal-kernel") != std::string::npos) ? true : false; - - aclType continueCompileFrom = - llvmBinary_.empty() ? getNextCompilationStageFromBinary(options) : ACL_TYPE_LLVMIR_BINARY; - - if (continueCompileFrom == ACL_TYPE_ISA) { - binary_t isa = binary(); - if ((isa.first != NULL) && (isa.second > 0)) { - return setKernels(options, (void*)isa.first, isa.second); - } else { - buildLog_ += "Error: code object is empty \n"; - return false; - } - return true; - } - if (continueCompileFrom != ACL_TYPE_LLVMIR_BINARY) { - buildLog_ += "Error while Codegen phase: the binary is incomplete \n"; - return false; - } - - std::unique_ptr C(newCompilerInstance()); - // call LinkLLVMBitcode - std::vector inputs; - - // open the input IR source - Data* input = C->NewBufferReference(DT_LLVM_BC, llvmBinary_.data(), llvmBinary_.size()); - - if (!input) { - buildLog_ += "Error: Failed to open the compiled program.\n"; - return false; - } - - inputs.push_back(input); //< must be the first input - - // open the bitcode libraries - Data* opencl_bc = - C->NewBufferReference(DT_LLVM_BC, (const char*)opencl_amdgcn, opencl_amdgcn_size); - Data* ocml_bc = C->NewBufferReference(DT_LLVM_BC, (const char*)ocml_amdgcn, ocml_amdgcn_size); - Data* ockl_bc = C->NewBufferReference(DT_LLVM_BC, (const char*)ockl_amdgcn, ockl_amdgcn_size); - - if (!opencl_bc || !ocml_bc || !ockl_bc) { - buildLog_ += "Error: Failed to open the bitcode library.\n"; - return false; - } - - inputs.push_back(opencl_bc); // depends on oclm & ockl - inputs.push_back(ockl_bc); - inputs.push_back(ocml_bc); - - // open the control functions - auto isa_version = get_oclc_isa_version(dev().hwInfo()->gfxipVersionLC_); - if (!isa_version.first) { - buildLog_ += "Error: Linking for this device is not supported\n"; - return false; - } - - Data* isa_version_bc = - C->NewBufferReference(DT_LLVM_BC, (const char*)isa_version.first, isa_version.second); - - if (!isa_version_bc) { - buildLog_ += "Error: Failed to open the control functions.\n"; - return false; - } - - inputs.push_back(isa_version_bc); - - auto correctly_rounded_sqrt = - get_oclc_correctly_rounded_sqrt(options->oVariables->FP32RoundDivideSqrt); - Data* correctly_rounded_sqrt_bc = C->NewBufferReference(DT_LLVM_BC, correctly_rounded_sqrt.first, - correctly_rounded_sqrt.second); - - auto daz_opt = get_oclc_daz_opt(options->oVariables->DenormsAreZero || - AMD_GPU_FORCE_SINGLE_FP_DENORM == 0 || - (dev().hwInfo()->gfxipVersionLC_ < 900 && - AMD_GPU_FORCE_SINGLE_FP_DENORM < 0)); - Data* daz_opt_bc = C->NewBufferReference(DT_LLVM_BC, daz_opt.first, daz_opt.second); - - auto finite_only = get_oclc_finite_only(options->oVariables->FiniteMathOnly || - options->oVariables->FastRelaxedMath); - Data* finite_only_bc = C->NewBufferReference(DT_LLVM_BC, finite_only.first, finite_only.second); - - auto unsafe_math = get_oclc_unsafe_math(options->oVariables->UnsafeMathOpt || - options->oVariables->FastRelaxedMath); - Data* unsafe_math_bc = C->NewBufferReference(DT_LLVM_BC, unsafe_math.first, unsafe_math.second); - - if (!correctly_rounded_sqrt_bc || !daz_opt_bc || !finite_only_bc || !unsafe_math_bc) { - buildLog_ += "Error: Failed to open the control functions.\n"; - return false; - } - - inputs.push_back(correctly_rounded_sqrt_bc); - inputs.push_back(daz_opt_bc); - inputs.push_back(finite_only_bc); - inputs.push_back(unsafe_math_bc); - - // open the linked output - std::vector linkOptions; - amd::opencl_driver::Buffer* linked_bc = C->NewBuffer(DT_LLVM_BC); - - if (!linked_bc) { - buildLog_ += "Error: Failed to open the linked program.\n"; - return false; - } - - // NOTE: The linkOptions parameter is also used to identy cached code object. This parameter - // should not contain any dyanamically generated filename. - bool ret = - dev().cacheCompilation()->linkLLVMBitcode(C.get(), inputs, linked_bc, linkOptions, buildLog_); - buildLog_ += C->Output(); - if (!ret) { - buildLog_ += "Error: Linking bitcode failed: linking source & IR libraries.\n"; - return false; - } - - if (options->isDumpFlagSet(amd::option::DUMP_BC_LINKED)) { - std::ofstream f(options->getDumpFileName("_linked.bc").c_str(), - std::ios::binary | std::ios::trunc); - if (f.is_open()) { - f.write(linked_bc->Buf().data(), linked_bc->Size()); - f.close(); - } else { - buildLog_ += "Warning: opening the file to dump the linked IR failed.\n"; - } - } - - inputs.clear(); - inputs.push_back(linked_bc); - - amd::opencl_driver::Buffer* out_exec = C->NewBuffer(DT_EXECUTABLE); - if (!out_exec) { - buildLog_ += "Error: Failed to create the linked executable.\n"; - return false; - } - - std::string codegenOptions(options->llvmOptions); - - // Set the machine target - std::ostringstream mCPU; - mCPU << " -mcpu=gfx" << dev().hwInfo()->gfxipVersionLC_; - codegenOptions.append(mCPU.str()); - - // Set xnack option if needed - if (dev().hwInfo()->xnackEnabled_) { - codegenOptions.append(" -mxnack"); - } - - // Set the -O# - std::ostringstream optLevel; - optLevel << "-O" << options->oVariables->OptLevel; - codegenOptions.append(" ").append(optLevel.str()); - - // Pass clang options - std::ostringstream ostrstr; - std::copy(options->clangOptions.begin(), options->clangOptions.end(), - std::ostream_iterator(ostrstr, " ")); - codegenOptions.append(" ").append(ostrstr.str()); - - // Set whole program mode - codegenOptions.append(" -mllvm -amdgpu-internalize-symbols -mllvm -amdgpu-early-inline-all"); - - // Tokenize the options string into a vector of strings - std::istringstream strstr(codegenOptions); - std::istream_iterator sit(strstr), end; - std::vector params(sit, end); - - // NOTE: The params is also used to identy cached code object. This parameter - // should not contain any dyanamically generated filename. - ret = dev().cacheCompilation()->compileAndLinkExecutable(C.get(), inputs, out_exec, params, - buildLog_); - buildLog_ += C->Output(); - if (!ret) { - buildLog_ += "Error: Creating the executable failed: Compiling LLVM IRs to exeutable\n"; - return false; - } - - if (options->isDumpFlagSet(amd::option::DUMP_O)) { - std::ofstream f(options->getDumpFileName(".so").c_str(), std::ios::binary | std::ios::trunc); - if (f.is_open()) { - f.write(out_exec->Buf().data(), out_exec->Size()); - f.close(); - } else { - buildLog_ += "Warning: opening the file to dump the code object failed.\n"; - } - } - - if (options->isDumpFlagSet(amd::option::DUMP_ISA)) { - std::string name = options->getDumpFileName(".s"); - File* dump = C->NewFile(DT_INTERNAL, name); - if (!C->DumpExecutableAsText(out_exec, dump)) { - buildLog_ += "Warning: failed to dump code object.\n"; - } - } - - return setKernels(options, out_exec->Buf().data(), out_exec->Size()); -} - -bool LightningProgram::setKernels(amd::option::Options* options, void* binary, size_t size) { +bool LightningProgram::setKernels(amd::option::Options* options, void* binary, size_t binSize) { hsa_agent_t agent; agent.handle = 1; @@ -960,7 +682,7 @@ bool LightningProgram::setKernels(amd::option::Options* options, void* binary, s } // Find the size of global variables from the binary - if (!FindGlobalVarSize(binary, size)) { + if (!FindGlobalVarSize(binary, binSize)) { return false; } @@ -1010,10 +732,6 @@ bool LightningProgram::setKernels(amd::option::Options* options, void* binary, s DestroySegmentCpuAccess(); - // Save the binary and type - clBinary()->saveBIFBinary((char*)binary, size); - setType(TYPE_EXECUTABLE); - return true; } diff --git a/rocclr/runtime/device/pal/palprogram.hpp b/rocclr/runtime/device/pal/palprogram.hpp index 78a0140358..c1daa131cb 100644 --- a/rocclr/runtime/device/pal/palprogram.hpp +++ b/rocclr/runtime/device/pal/palprogram.hpp @@ -183,8 +183,6 @@ class HSAILProgram : public device::Program { bool saveBinaryAndSetType(type_t type); - virtual bool linkImpl(amd::option::Options* options); - virtual bool createBinary(amd::option::Options* options); virtual const aclTargetInfo& info(const char* str = ""); @@ -193,6 +191,8 @@ class HSAILProgram : public device::Program { return amd::isElfMagic(bin); } + virtual bool setKernels(amd::option::Options* options, void* binary, size_t binSize) override; + //! Destroys CPU allocations in the code segment void DestroySegmentCpuAccess() const { if (codeSegment_ != nullptr) { codeSegment_->DestroyCpuAccess(); } } @@ -242,9 +242,7 @@ class LightningProgram : public HSAILProgram { virtual ~LightningProgram() {} protected: - virtual bool linkImpl(amd::option::Options* options) override; - - bool setKernels(amd::option::Options* options, void* binary, size_t size); + virtual bool setKernels(amd::option::Options* options, void* binary, size_t binSize) override; virtual bool createBinary(amd::option::Options* options) override; }; diff --git a/rocclr/runtime/device/pal/palsettings.cpp b/rocclr/runtime/device/pal/palsettings.cpp index 5f035c77a6..7493e91618 100644 --- a/rocclr/runtime/device/pal/palsettings.cpp +++ b/rocclr/runtime/device/pal/palsettings.cpp @@ -457,6 +457,8 @@ bool Settings::create(const Pal::DeviceProperties& palProp, } #endif // WITH_LIGHTNING_COMPILER + gfx10Hsail_ = gfx10Plus_; + // Override current device settings override(); diff --git a/rocclr/runtime/device/rocm/rocprogram.cpp b/rocclr/runtime/device/rocm/rocprogram.cpp index 088c3e83ae..0abf24d977 100644 --- a/rocclr/runtime/device/rocm/rocprogram.cpp +++ b/rocclr/runtime/device/rocm/rocprogram.cpp @@ -195,76 +195,15 @@ bool HSAILProgram::saveBinaryAndSetType(type_t type) { return true; } -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 + ProcessOptions(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; - } +bool HSAILProgram::setKernels(amd::option::Options* options, void* binary, size_t binSize) { // 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; - } + size_t secSize = binSize; + void* data = binary; // Create an executable. hsa_status_t status = hsa_executable_create_alt( @@ -349,7 +288,7 @@ bool HSAILProgram::linkImpl(amd::option::Options* options) { md.numHiddenKernelArgs = 0; size_t sizeOfnumHiddenKernelArgs = sizeof(md.numHiddenKernelArgs); - errorCode = aclQueryInfo(device().compiler(), binaryElf_, + acl_error errorCode = aclQueryInfo(device().compiler(), binaryElf_, RT_NUM_KERNEL_HIDDEN_ARGS, openclKernelName.c_str(), &md.numHiddenKernelArgs, &sizeOfnumHiddenKernelArgs); if (errorCode != ACL_SUCCESS) { @@ -415,8 +354,6 @@ bool HSAILProgram::linkImpl(amd::option::Options* options) { std::string::npos); kernels()[kernelName] = aKernel; } - saveBinaryAndSetType(TYPE_EXECUTABLE); - buildLog_ += aclGetCompilerLog(device().compiler()); return true; } #endif // defined(WITH_COMPILER_LIB) @@ -457,234 +394,12 @@ bool LightningProgram::saveBinaryAndSetType(type_t type, void* rawBinary, size_t return true; } -bool LightningProgram::linkImpl(amd::option::Options* options) { - acl_error errorCode; - aclType continueCompileFrom = ACL_TYPE_LLVMIR_BINARY; - using namespace amd::opencl_driver; - std::vector inputs; - std::unique_ptr C(newCompilerInstance()); - bool bLinkLLVMBitcode = true; - if (llvmBinary_.empty()) { - continueCompileFrom = getNextCompilationStageFromBinary(options); - } - switch (continueCompileFrom) { - case ACL_TYPE_CG: - case ACL_TYPE_LLVMIR_BINARY: { - break; - } - case ACL_TYPE_ASM_TEXT: { - char* section; - size_t sz; - clBinary()->elfOut()->getSection(amd::OclElf::SOURCE, §ion, &sz); - Data* input = C->NewBufferReference(DT_ASSEMBLY, section, sz); - if (!input) { - buildLog_ += "Error: Failed to open the assembler text.\n"; - return false; - } - inputs.push_back(input); - bLinkLLVMBitcode = false; - 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; - } - default: - buildLog_ += "Error while Codegen phase: the binary is incomplete \n"; - return false; - } - - // call LinkLLVMBitcode - if (bLinkLLVMBitcode) { - // open the input IR source - Data* input = C->NewBufferReference(DT_LLVM_BC, llvmBinary_.data(), llvmBinary_.size()); - - if (!input) { - buildLog_ += "Error: Failed to open the compiled program.\n"; - return false; - } - - inputs.push_back(input); //< must be the first input - - // open the bitcode libraries - Data* opencl_bc = - C->NewBufferReference(DT_LLVM_BC, (const char*)opencl_amdgcn, opencl_amdgcn_size); - Data* ocml_bc = C->NewBufferReference(DT_LLVM_BC, (const char*)ocml_amdgcn, ocml_amdgcn_size); - Data* ockl_bc = C->NewBufferReference(DT_LLVM_BC, (const char*)ockl_amdgcn, ockl_amdgcn_size); - - if (!opencl_bc || !ocml_bc || !ockl_bc) { - buildLog_ += "Error: Failed to open the bitcode library.\n"; - return false; - } - - inputs.push_back(opencl_bc); // depends on oclm & ockl - inputs.push_back(ockl_bc); - inputs.push_back(ocml_bc); - - // open the control functions - auto isa_version = get_oclc_isa_version(dev().deviceInfo().gfxipVersion_); - if (!isa_version.first) { - buildLog_ += "Error: Linking for this device is not supported\n"; - return false; - } - - Data* isa_version_bc = - C->NewBufferReference(DT_LLVM_BC, (const char*)isa_version.first, isa_version.second); - - if (!isa_version_bc) { - buildLog_ += "Error: Failed to open the control functions.\n"; - return false; - } - - inputs.push_back(isa_version_bc); - - auto correctly_rounded_sqrt = - get_oclc_correctly_rounded_sqrt(options->oVariables->FP32RoundDivideSqrt); - Data* correctly_rounded_sqrt_bc = C->NewBufferReference(DT_LLVM_BC, correctly_rounded_sqrt.first, - correctly_rounded_sqrt.second); - - auto daz_opt = get_oclc_daz_opt(options->oVariables->DenormsAreZero || - AMD_GPU_FORCE_SINGLE_FP_DENORM == 0 || - (dev().deviceInfo().gfxipVersion_ < 900 && - AMD_GPU_FORCE_SINGLE_FP_DENORM < 0)); - Data* daz_opt_bc = C->NewBufferReference(DT_LLVM_BC, daz_opt.first, daz_opt.second); - - auto finite_only = get_oclc_finite_only(options->oVariables->FiniteMathOnly || - options->oVariables->FastRelaxedMath); - Data* finite_only_bc = C->NewBufferReference(DT_LLVM_BC, finite_only.first, finite_only.second); - - auto unsafe_math = get_oclc_unsafe_math(options->oVariables->UnsafeMathOpt || - options->oVariables->FastRelaxedMath); - Data* unsafe_math_bc = C->NewBufferReference(DT_LLVM_BC, unsafe_math.first, unsafe_math.second); - - if (!correctly_rounded_sqrt_bc || !daz_opt_bc || !finite_only_bc || !unsafe_math_bc) { - buildLog_ += "Error: Failed to open the control functions.\n"; - return false; - } - - inputs.push_back(correctly_rounded_sqrt_bc); - inputs.push_back(daz_opt_bc); - inputs.push_back(finite_only_bc); - inputs.push_back(unsafe_math_bc); - - // open the linked output - std::vector linkOptions; - Buffer* linked_bc = C->NewBuffer(DT_LLVM_BC); - - if (!linked_bc) { - buildLog_ += "Error: Failed to open the linked program.\n"; - return false; - } - - // NOTE: The linkOptions parameter is also used to identy cached code object. This parameter - // should not contain any dyanamically generated filename. - bool ret = - dev().cacheCompilation()->linkLLVMBitcode(C.get(), inputs, linked_bc, linkOptions, buildLog_); - buildLog_ += C->Output(); - if (!ret) { - buildLog_ += "Error: Linking bitcode failed: linking source & IR libraries.\n"; - return false; - } - - if (options->isDumpFlagSet(amd::option::DUMP_BC_LINKED)) { - std::ofstream f(options->getDumpFileName("_linked.bc").c_str(), std::ios::trunc); - if (f.is_open()) { - f.write(linked_bc->Buf().data(), linked_bc->Size()); - } - else { - buildLog_ += "Warning: opening the file to dump the linked IR failed.\n"; - } - } - - inputs.clear(); - inputs.push_back(linked_bc); - } - - Buffer* out_exec = C->NewBuffer(DT_EXECUTABLE); - if (!out_exec) { - buildLog_ += "Error: Failed to create the linked executable.\n"; - return false; - } - - std::string codegenOptions(options->llvmOptions); - - // Set the machine target - codegenOptions.append(" -mcpu="); - codegenOptions.append(dev().deviceInfo().machineTarget_); - - // Set xnack option if needed - if (dev().deviceInfo().xnackEnabled_) { - codegenOptions.append(" -mxnack"); - } - - // Set the -O# - std::ostringstream optLevel; - optLevel << "-O" << options->oVariables->OptLevel; - codegenOptions.append(" ").append(optLevel.str()); - - // Pass clang options - std::ostringstream ostrstr; - std::copy(options->clangOptions.begin(), options->clangOptions.end(), - std::ostream_iterator(ostrstr, " ")); - codegenOptions.append(" ").append(ostrstr.str()); - - // Set whole program mode - codegenOptions.append(" -mllvm -amdgpu-internalize-symbols -mllvm -amdgpu-early-inline-all"); - - // Tokenize the options string into a vector of strings - std::istringstream strstr(codegenOptions); - std::istream_iterator sit(strstr), end; - std::vector params(sit, end); - - // NOTE: The params is also used to identy cached code object. This parameter - // should not contain any dyanamically generated filename. - bool ret = dev().cacheCompilation()->compileAndLinkExecutable(C.get(), inputs, out_exec, params, - buildLog_); - buildLog_ += C->Output(); - if (!ret) { - if (continueCompileFrom == ACL_TYPE_ASM_TEXT) { - buildLog_ += "Error: Creating the executable from ISA assembly text failed.\n"; - } else { - buildLog_ += "Error: Creating the executable from LLVM IRs failed.\n"; - } - return false; - } - - if (options->isDumpFlagSet(amd::option::DUMP_O)) { - std::ofstream f(options->getDumpFileName(".so").c_str(), std::ios::trunc); - if (f.is_open()) { - f.write(out_exec->Buf().data(), out_exec->Size()); - } else { - buildLog_ += "Warning: opening the file to dump the code object failed.\n"; - } - } - - if (options->isDumpFlagSet(amd::option::DUMP_ISA)) { - std::string name = options->getDumpFileName(".s"); - File* dump = C->NewFile(DT_INTERNAL, name); - if (!C->DumpExecutableAsText(out_exec, dump)) { - buildLog_ += "Warning: failed to dump code object.\n"; - } - } - - return setKernels(options, out_exec->Buf().data(), out_exec->Size()); -} - bool LightningProgram::setKernels(amd::option::Options* options, void* binary, size_t binSize) { // Find the size of global variables from the binary if (!FindGlobalVarSize(binary, binSize)) { return false; } - saveBinaryAndSetType(TYPE_EXECUTABLE, binary, binSize); - //Load the stored copy of the ELF binary. binary_t stored_binary = this->binary(); binary = const_cast(stored_binary.first); diff --git a/rocclr/runtime/device/rocm/rocprogram.hpp b/rocclr/runtime/device/rocm/rocprogram.hpp index 2348566e64..89f2f88d2b 100644 --- a/rocclr/runtime/device/rocm/rocprogram.hpp +++ b/rocclr/runtime/device/rocm/rocprogram.hpp @@ -90,10 +90,10 @@ class HSAILProgram : public roc::Program { virtual ~HSAILProgram(); protected: - virtual bool linkImpl(amd::option::Options* options) final; - virtual bool createBinary(amd::option::Options* options) final; + virtual bool setKernels(amd::option::Options* options, void* binary, size_t binSize) override; + private: std::string codegenOptions(amd::option::Options* options); @@ -108,8 +108,6 @@ public: virtual ~LightningProgram() {} protected: - virtual bool linkImpl(amd::option::Options* options) final; - virtual bool createBinary(amd::option::Options* options) final; bool saveBinaryAndSetType(type_t type) { return true; } @@ -117,7 +115,7 @@ protected: private: bool saveBinaryAndSetType(type_t type, void* rawBinary, size_t size); - bool setKernels(amd::option::Options* options, void* binary, size_t binSize); + virtual bool setKernels(amd::option::Options* options, void* binary, size_t binSize) override; }; #endif // defined(WITH_LIGHTNING_COMPILER)