AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

P4 to Git Change 1611775 by gandryey@gera-ocl-lc on 2018/09/27 18:02:54

Browse Source 
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


[ROCm/clr commit: 8bbcc1c38a]
This commit is contained in:
foreman
2018-09-27 18:15:24 -04:00
parent c4c871ba6c
commit dc9db4eb3e
9 changed files with 515 additions and 729 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/rocclr/runtime/device/device.hpp
+2 -1
View File
@@ -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_;
};
projects/clr/rocclr/runtime/device/devprogram.cpp
+436 -91
View File
@@ -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 <cstdio>
#include <fstream>
#include <iostream>
@@ -492,97 +494,6 @@ bool Program::linkImpl(const std::vector<device::Program*>& inputPrograms,
}
}
// ================================================================================================
bool Program::linkImplHSAIL(const std::vector<Program*>& 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<aclBinary*> 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<void*>(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<Program*>& inputPrograms,
amd::option::Options* options, bool createLibrary) {
@@ -678,6 +589,440 @@ bool Program::linkImplLC(const std::vector<Program*>& inputPrograms,
#endif // defined(WITH_LIGHTNING_COMPILER)
}
// ================================================================================================
bool Program::linkImplHSAIL(const std::vector<Program*>& 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<aclBinary*> 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<void*>(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<Data*> inputs;
std::unique_ptr<Compiler> 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, &section, &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<std::string> 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<std::string>(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<std::string> sit(strstr), end;
std::vector<std::string> 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<const char*>(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<void*>(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) {
projects/clr/rocclr/runtime/device/devprogram.hpp
+24 -15
View File
@@ -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<Program*>& 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<const std::string*>& headers,
const char** headerIncludeNames, amd::option::Options* options);
//! Compile the device program with LC path
bool compileImplLC(const std::string& sourceCode,
const std::vector<const std::string*>& headers,
const char** headerIncludeNames, amd::option::Options* options);
//! Compile the device program with HSAIL path
bool compileImplHSAIL(const std::string& sourceCode,
const std::vector<const std::string*>& headers,
const char** headerIncludeNames, amd::option::Options* options);
//! Compile the device program with HSAIL path
bool compileImplHSAIL(const std::string& sourceCode,
const std::vector<const std::string*>& headers,
const char** headerIncludeNames, amd::option::Options* options);
//! Link the device programs with HSAIL path
bool linkImplHSAIL(const std::vector<Program*>& inputPrograms,
amd::option::Options* options, bool createLibrary);
//! Link the device programs with LC path
bool linkImplLC(const std::vector<Program*>& inputPrograms,
amd::option::Options* options, bool createLibrary);
//! Link the device programs with LC path
bool linkImplLC(const std::vector<Program*>& inputPrograms,
amd::option::Options* options, bool createLibrary);
//! Link the device programs with HSAIL path
bool linkImplHSAIL(const std::vector<Program*>& 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&);
projects/clr/rocclr/runtime/device/pal/paldefs.hpp
+17 -17
View File
@@ -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},
projects/clr/rocclr/runtime/device/pal/palprogram.cpp
+24 -306
View File
@@ -244,113 +244,37 @@ inline static std::vector<std::string> 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<uint64_t>(
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<uint64_t>(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<Compiler> C(newCompilerInstance());
// call LinkLLVMBitcode
std::vector<Data*> 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<std::string> 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<std::string>(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<std::string> sit(strstr), end;
std::vector<std::string> 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;
}
projects/clr/rocclr/runtime/device/pal/palprogram.hpp
+3 -5
View File
@@ -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;
};
projects/clr/rocclr/runtime/device/pal/palsettings.cpp
+2
View File
@@ -457,6 +457,8 @@ bool Settings::create(const Pal::DeviceProperties& palProp,
}
#endif // WITH_LIGHTNING_COMPILER
gfx10Hsail_ = gfx10Plus_;
// Override current device settings
override();
projects/clr/rocclr/runtime/device/rocm/rocprogram.cpp
+4 -289
View File
@@ -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<Data*> inputs;
std::unique_ptr<Compiler> 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, &section, &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<std::string> 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<std::string>(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<std::string> sit(strstr), end;
std::vector<std::string> 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<void*>(stored_binary.first);
projects/clr/rocclr/runtime/device/rocm/rocprogram.hpp
+3 -5
View File
@@ -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)