AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

SWDEV-280473 - Remove HSAIL support from the ROCm backend

Просмотр исходного кода 
In adition to removing the HSAIL logic from the ROCm backend, guard all
of the HSAIL includes in the common layer behind the WITH_COMPILER_LIB
define. This is to avoid including HSAIL headers when building with
no support for it.

In common logic replace the use of the aclType enum with the new
Program::file_type_t enum. This is essentially a local copy of the HSAIL
enum to avoid including any HSAIL headers.

Change-Id: Ica0651d1b29dfccc255cc584eb82a5cb35e1b520
Этот коммит содержится в:
Vladislav Sytchenko
2021-04-12 14:55:06 -04:00
родитель 2b133ed3de
Коммит cbeb372e46
17 изменённых файлов: 126 добавлений и 405 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
rocclr/device/device.cpp
+12
Просмотреть файл
@@ -56,9 +56,11 @@ extern void DeviceUnload();
#include "blowfish/oclcrypt.hpp"
#endif
#if defined(WITH_COMPILER_LIB)
#include "utils/bif_section_labels.hpp"
#include "utils/libUtils.h"
#include "spirv/spirvUtils.h"
#endif
#include <vector>
#include <string>
@@ -757,6 +759,7 @@ bool ClBinary::setElfTarget() {
return true;
}
#if defined(WITH_COMPILER_LIB)
std::string ClBinary::getBIFSymbol(unsigned int symbolID) const {
size_t nSymbols = 0;
// Due to PRE & POST defines in bif_section_labels.hpp conflict with
@@ -791,6 +794,7 @@ std::string ClBinary::getBIFSymbol(unsigned int symbolID) const {
}
return "";
}
#endif
void ClBinary::init(amd::option::Options* optionsObj) {
// option has higher priority than environment variable.
@@ -1089,6 +1093,7 @@ bool ClBinary::loadCompileOptions(std::string& compileOptions) const {
char* options = nullptr;
size_t sz;
compileOptions.clear();
#if defined(WITH_COMPILER_LIB)
if (elfIn_->getSymbol(amd::Elf::COMMENT, getBIFSymbol(symOpenclCompilerOptions).c_str(),
&options, &sz)) {
if (sz > 0) {
@@ -1096,6 +1101,7 @@ bool ClBinary::loadCompileOptions(std::string& compileOptions) const {
}
return true;
}
#endif
return false;
}
@@ -1103,6 +1109,7 @@ bool ClBinary::loadLinkOptions(std::string& linkOptions) const {
char* options = nullptr;
size_t sz;
linkOptions.clear();
#if defined(WITH_COMPILER_LIB)
if (elfIn_->getSymbol(amd::Elf::COMMENT, getBIFSymbol(symOpenclLinkerOptions).c_str(),
&options, &sz)) {
if (sz > 0) {
@@ -1110,17 +1117,22 @@ bool ClBinary::loadLinkOptions(std::string& linkOptions) const {
}
return true;
}
#endif
return false;
}
void ClBinary::storeCompileOptions(const std::string& compileOptions) {
#if defined(WITH_COMPILER_LIB)
elfOut()->addSymbol(amd::Elf::COMMENT, getBIFSymbol(symOpenclCompilerOptions).c_str(),
compileOptions.c_str(), compileOptions.length());
#endif
}
void ClBinary::storeLinkOptions(const std::string& linkOptions) {
#if defined(WITH_COMPILER_LIB)
elfOut()->addSymbol(amd::Elf::COMMENT, getBIFSymbol(symOpenclLinkerOptions).c_str(),
linkOptions.c_str(), linkOptions.length());
#endif
}
bool ClBinary::isSPIR() const {
rocclr/device/device.hpp
+8
Просмотреть файл
@@ -34,7 +34,9 @@
#include "devprogram.hpp"
#include "devkernel.hpp"
#include "amdocl/cl_profile_amd.h"
#if defined(WITH_COMPILER_LIB)
#include "acl.h"
#endif
#include "hwdebug.hpp"
#include "devsignal.hpp"
@@ -1030,9 +1032,11 @@ class ClBinary : public amd::HeapObject {
//! Returns TRUE if binary file was allocated
bool isBinaryAllocated() const { return (flags_ & BinaryAllocated) ? true : false; }
#if defined(WITH_COMPILER_LIB)
//! Returns BIF symbol name by symbolID,
//! returns empty string if not found or if BIF version is unsupported
std::string getBIFSymbol(unsigned int symbolID) const;
#endif
protected:
const amd::Device& dev_; //!< Device object
@@ -1445,7 +1449,9 @@ class Isa {
*/
class Device : public RuntimeObject {
protected:
#if defined(WITH_COMPILER_LIB)
typedef aclCompiler Compiler;
#endif
public:
// The structures below for MGPU launch match the device library format
@@ -1498,8 +1504,10 @@ class Device : public RuntimeObject {
);
};
#if defined(WITH_COMPILER_LIB)
virtual Compiler* compiler() const = 0;
virtual Compiler* binCompiler() const { return compiler(); }
#endif
Device();
virtual ~Device();
rocclr/device/devkernel.cpp
+6
Просмотреть файл
@@ -24,15 +24,19 @@
#include "devkernel.hpp"
#include "utils/macros.hpp"
#include "utils/options.hpp"
#if defined(WITH_COMPILER_LIB)
#include "utils/bif_section_labels.hpp"
#include "utils/libUtils.h"
#endif
#include "comgrctx.hpp"
#include <map>
#include <string>
#include <sstream>
#if defined(WITH_COMPILER_LIB)
#include "acl.h"
#endif
namespace device {
@@ -660,11 +664,13 @@ bool Kernel::createSignature(
Kernel::~Kernel() { delete signature_; }
// ================================================================================================
#if defined(WITH_COMPILER_LIB)
std::string Kernel::openclMangledName(const std::string& name) {
const oclBIFSymbolStruct* bifSym = findBIF30SymStruct(symOpenclKernel);
assert(bifSym && "symbol not found");
return std::string("&") + bifSym->str[bif::PRE] + name + bifSym->str[bif::POST];
}
#endif
// ================================================================================================
void Kernel::FindLocalWorkSize(size_t workDim, const amd::NDRange& gblWorkSize,
rocclr/device/devkernel.hpp
+4
Просмотреть файл
@@ -20,7 +20,9 @@
#pragma once
#if defined(WITH_COMPILER_LIB)
#include "aclTypes.h"
#endif
#include "platform/context.hpp"
#include "platform/object.hpp"
#include "platform/memory.hpp"
@@ -421,7 +423,9 @@ class Kernel : public amd::HeapObject {
//! Return the build log
const std::string& buildLog() const { return buildLog_; }
#if defined(WITH_COMPILER_LIB)
static std::string openclMangledName(const std::string& name);
#endif
const std::unordered_map<size_t, size_t>& patch() const { return patchReferences_; }
rocclr/device/devprogram.cpp
+51 -42
Просмотреть файл
@@ -25,8 +25,10 @@
#include "devkernel.hpp"
#include "utils/macros.hpp"
#include "utils/options.hpp"
#if defined(WITH_COMPILER_LIB)
#include "utils/bif_section_labels.hpp"
#include "utils/libUtils.h"
#endif
#include "comgrctx.hpp"
#include <algorithm>
@@ -43,8 +45,10 @@
#include <libgen.h>
#endif // defined(ATI_OS_LINUX)
#if defined(WITH_COMPILER_LIB)
#include "spirv/spirvUtils.h"
#include "acl.h"
#endif
#ifdef EARLY_INLINE
#define AMDGPU_EARLY_INLINE_ALL_OPTION " -mllvm -amdgpu-early-inline-all"
@@ -78,19 +82,23 @@ Program::Program(amd::Device& device, amd::Program& owner)
elfSectionType_(amd::Elf::LLVMIR),
compileOptions_(),
linkOptions_(),
#if defined(WITH_COMPILER_LIB)
binaryElf_(nullptr),
#endif
lastBuildOptionsArg_(),
buildStatus_(CL_BUILD_NONE),
buildError_(CL_SUCCESS),
globalVariableTotalSize_(0),
programOptions_(nullptr)
{
#if defined(WITH_COMPILER_LIB)
memset(&binOpts_, 0, sizeof(binOpts_));
binOpts_.struct_size = sizeof(binOpts_);
binOpts_.elfclass = LP64_SWITCH(ELFCLASS32, ELFCLASS64);
binOpts_.bitness = ELFDATA2LSB;
binOpts_.alloc = &::malloc;
binOpts_.dealloc = &::free;
#endif
}
// ================================================================================================
@@ -1103,7 +1111,7 @@ static void dumpCodeObject(const std::string& image) {
// ================================================================================================
bool Program::linkImplLC(amd::option::Options* options) {
#if defined(USE_COMGR_LIBRARY)
aclType continueCompileFrom = ACL_TYPE_LLVMIR_BINARY;
file_type_t continueCompileFrom = FILE_TYPE_LLVMIR_BINARY;
internal_ = (compileOptions_.find("-cl-internal-kernel") != std::string::npos) ?
true : false;
@@ -1120,11 +1128,11 @@ bool Program::linkImplLC(amd::option::Options* options) {
}
switch (continueCompileFrom) {
case ACL_TYPE_CG:
case ACL_TYPE_LLVMIR_BINARY: {
case FILE_TYPE_CG:
case FILE_TYPE_LLVMIR_BINARY: {
break;
}
case ACL_TYPE_ASM_TEXT: {
case FILE_TYPE_ASM_TEXT: {
char* section;
size_t sz;
clBinary()->elfOut()->getSection(amd::Elf::SOURCE, &section, &sz);
@@ -1139,7 +1147,7 @@ bool Program::linkImplLC(amd::option::Options* options) {
bLinkLLVMBitcode = false;
break;
}
case ACL_TYPE_ISA: {
case FILE_TYPE_ISA: {
amd::Comgr::destroy_data_set(inputs);
binary_t isaBinary = binary();
finfo_t isaFdesc = BinaryFd();
@@ -1262,7 +1270,7 @@ bool Program::linkImplLC(amd::option::Options* options) {
amd::Comgr::destroy_data_set(inputs);
if (!ret) {
if (continueCompileFrom == ACL_TYPE_ASM_TEXT) {
if (continueCompileFrom == FILE_TYPE_ASM_TEXT) {
buildLog_ += "Error: Creating the executable from ISA assembly text failed.\n";
} else {
buildLog_ += "Error: Creating the executable from LLVM IRs failed.\n";
@@ -2100,9 +2108,10 @@ bool Program::setBinary(const char* binaryIn, size_t size, const device::Program
}
// ================================================================================================
aclType Program::getCompilationStagesFromBinary(std::vector<aclType>& completeStages,
Program::file_type_t Program::getCompilationStagesFromBinary(
std::vector<Program::file_type_t>& completeStages,
bool& needOptionsCheck) {
aclType from = ACL_TYPE_DEFAULT;
Program::file_type_t from = FILE_TYPE_DEFAULT;
if (isLC()) {
#if defined(USE_COMGR_LIBRARY)
completeStages.clear();
@@ -2115,30 +2124,30 @@ aclType Program::getCompilationStagesFromBinary(std::vector<aclType>& completeSt
if (containsLlvmirText && containsOpts) {
completeStages.push_back(from);
from = ACL_TYPE_LLVMIR_BINARY;
from = FILE_TYPE_LLVMIR_BINARY;
}
if (containsShaderIsa) {
completeStages.push_back(from);
from = ACL_TYPE_ISA;
from = FILE_TYPE_ISA;
}
std::string sCurOptions = compileOptions_ + linkOptions_;
amd::option::Options curOptions;
if (!amd::option::parseAllOptions(sCurOptions, curOptions, false, isLC())) {
buildLog_ += curOptions.optionsLog();
LogError("Parsing compile options failed.");
return ACL_TYPE_DEFAULT;
return FILE_TYPE_DEFAULT;
}
switch (from) {
case ACL_TYPE_CG:
case ACL_TYPE_ISA:
case FILE_TYPE_CG:
case FILE_TYPE_ISA:
// do not check options, if LLVMIR is absent or might be absent or options are absent
if (!curOptions.oVariables->BinLLVMIR || !containsLlvmirText || !containsOpts) {
needOptionsCheck = false;
}
break;
// recompilation might be needed
case ACL_TYPE_LLVMIR_BINARY:
case ACL_TYPE_DEFAULT:
case FILE_TYPE_LLVMIR_BINARY:
case FILE_TYPE_DEFAULT:
default:
break;
}
@@ -2159,7 +2168,7 @@ aclType Program::getCompilationStagesFromBinary(std::vector<aclType>& completeSt
}
if (containsSpirv) {
completeStages.push_back(from);
from = ACL_TYPE_SPIRV_BINARY;
from = FILE_TYPE_SPIRV_BINARY;
}
bool containsSpirText = true;
errorCode = aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_SPIR, nullptr,
@@ -2169,7 +2178,7 @@ aclType Program::getCompilationStagesFromBinary(std::vector<aclType>& completeSt
}
if (containsSpirText) {
completeStages.push_back(from);
from = ACL_TYPE_SPIR_BINARY;
from = FILE_TYPE_SPIR_BINARY;
}
bool containsLlvmirText = true;
errorCode = aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_LLVMIR, nullptr,
@@ -2186,7 +2195,7 @@ aclType Program::getCompilationStagesFromBinary(std::vector<aclType>& completeSt
}
if (containsLlvmirText && containsOpts) {
completeStages.push_back(from);
from = ACL_TYPE_LLVMIR_BINARY;
from = FILE_TYPE_LLVMIR_BINARY;
}
// Checking HSAIL in .cg section
bool containsHsailText = true;
@@ -2204,11 +2213,11 @@ aclType Program::getCompilationStagesFromBinary(std::vector<aclType>& completeSt
}
if (containsBrig) {
completeStages.push_back(from);
from = ACL_TYPE_HSAIL_BINARY;
from = FILE_TYPE_HSAIL_BINARY;
}
else if (containsHsailText) {
completeStages.push_back(from);
from = ACL_TYPE_HSAIL_TEXT;
from = FILE_TYPE_HSAIL_TEXT;
}
// Checking Loader Map symbol from CG section
bool containsLoaderMap = true;
@@ -2219,7 +2228,7 @@ aclType Program::getCompilationStagesFromBinary(std::vector<aclType>& completeSt
}
if (containsLoaderMap) {
completeStages.push_back(from);
from = ACL_TYPE_CG;
from = FILE_TYPE_CG;
}
// Checking ISA in .text section
bool containsShaderIsa = true;
@@ -2230,28 +2239,28 @@ aclType Program::getCompilationStagesFromBinary(std::vector<aclType>& completeSt
}
if (containsShaderIsa) {
completeStages.push_back(from);
from = ACL_TYPE_ISA;
from = FILE_TYPE_ISA;
}
std::string sCurOptions = compileOptions_ + linkOptions_;
amd::option::Options curOptions;
if (!amd::option::parseAllOptions(sCurOptions, curOptions, false, isLC())) {
buildLog_ += curOptions.optionsLog();
LogError("Parsing compile options failed.");
return ACL_TYPE_DEFAULT;
return FILE_TYPE_DEFAULT;
}
switch (from) {
// compile from HSAIL text, no matter prev. stages and options
case ACL_TYPE_HSAIL_TEXT:
case FILE_TYPE_HSAIL_TEXT:
needOptionsCheck = false;
break;
case ACL_TYPE_HSAIL_BINARY:
case FILE_TYPE_HSAIL_BINARY:
// do not check options, if LLVMIR is absent or might be absent or options are absent
if (!curOptions.oVariables->BinLLVMIR || !containsLlvmirText || !containsOpts) {
needOptionsCheck = false;
}
break;
case ACL_TYPE_CG:
case ACL_TYPE_ISA:
case FILE_TYPE_CG:
case FILE_TYPE_ISA:
// do not check options, if LLVMIR is absent or might be absent or options are absent
if (!curOptions.oVariables->BinLLVMIR || !containsLlvmirText || !containsOpts) {
needOptionsCheck = false;
@@ -2262,8 +2271,8 @@ aclType Program::getCompilationStagesFromBinary(std::vector<aclType>& completeSt
}
break;
// recompilation might be needed
case ACL_TYPE_LLVMIR_BINARY:
case ACL_TYPE_DEFAULT:
case FILE_TYPE_LLVMIR_BINARY:
case FILE_TYPE_DEFAULT:
default:
break;
}
@@ -2273,8 +2282,8 @@ aclType Program::getCompilationStagesFromBinary(std::vector<aclType>& completeSt
}
// ================================================================================================
aclType Program::getNextCompilationStageFromBinary(amd::option::Options* options) {
aclType continueCompileFrom = ACL_TYPE_DEFAULT;
Program::file_type_t Program::getNextCompilationStageFromBinary(amd::option::Options* options) {
Program::file_type_t continueCompileFrom = FILE_TYPE_DEFAULT;
binary_t binary = this->binary();
finfo_t finfo = this->BinaryFd();
std::string uri = this->BinaryURI();
@@ -2303,7 +2312,7 @@ aclType Program::getNextCompilationStageFromBinary(amd::option::Options* options
// Calculate the next stage to compile from, based on sections in binaryElf_;
// No any validity checks here
std::vector<aclType> completeStages;
std::vector<file_type_t> completeStages;
bool needOptionsCheck = true;
continueCompileFrom = getCompilationStagesFromBinary(completeStages, needOptionsCheck);
if (!options || !needOptionsCheck) {
@@ -2312,9 +2321,9 @@ aclType Program::getNextCompilationStageFromBinary(amd::option::Options* options
bool recompile = false;
//! @todo Should we also check for ACL_TYPE_OPENCL & ACL_TYPE_LLVMIR_TEXT?
switch (continueCompileFrom) {
case ACL_TYPE_HSAIL_BINARY:
case ACL_TYPE_CG:
case ACL_TYPE_ISA: {
case FILE_TYPE_HSAIL_BINARY:
case FILE_TYPE_CG:
case FILE_TYPE_ISA: {
// Compare options loaded from binary with current ones, recompile if differ;
// If compile options are absent in binary, do not compare and recompile
if (compileOptions_.empty()) break;
@@ -2350,12 +2359,12 @@ aclType Program::getNextCompilationStageFromBinary(amd::option::Options* options
if (!amd::option::parseAllOptions(sBinOptions, binOptions, false, isLC())) {
buildLog_ += binOptions.optionsLog();
LogError("Parsing compile options from binary failed.");
return ACL_TYPE_DEFAULT;
return FILE_TYPE_DEFAULT;
}
if (!amd::option::parseAllOptions(sCurOptions, curOptions, false, isLC())) {
buildLog_ += curOptions.optionsLog();
LogError("Parsing compile options failed.");
return ACL_TYPE_DEFAULT;
return FILE_TYPE_DEFAULT;
}
if (!curOptions.equals(binOptions)) {
recompile = true;
@@ -2368,10 +2377,10 @@ aclType Program::getNextCompilationStageFromBinary(amd::option::Options* options
if (recompile) {
while (!completeStages.empty()) {
continueCompileFrom = completeStages.back();
if (continueCompileFrom == ACL_TYPE_SPIRV_BINARY ||
continueCompileFrom == ACL_TYPE_LLVMIR_BINARY ||
continueCompileFrom == ACL_TYPE_SPIR_BINARY ||
continueCompileFrom == ACL_TYPE_DEFAULT) {
if (continueCompileFrom == FILE_TYPE_SPIRV_BINARY ||
continueCompileFrom == FILE_TYPE_LLVMIR_BINARY ||
continueCompileFrom == FILE_TYPE_SPIR_BINARY ||
continueCompileFrom == FILE_TYPE_DEFAULT) {
break;
}
completeStages.pop_back();
@@ -2381,7 +2390,7 @@ aclType Program::getNextCompilationStageFromBinary(amd::option::Options* options
else {
const char* xLang = options->oVariables->XLang;
if (xLang != nullptr && strcmp(xLang, "asm") == 0) {
continueCompileFrom = ACL_TYPE_ASM_TEXT;
continueCompileFrom = FILE_TYPE_ASM_TEXT;
}
}
return continueCompileFrom;
rocclr/device/devprogram.hpp
+37 -3
Просмотреть файл
@@ -20,7 +20,9 @@
#pragma once
#if defined(WITH_COMPILER_LIB)
#include "aclTypes.h"
#endif
#include "platform/context.hpp"
#include "platform/object.hpp"
#include "platform/memory.hpp"
@@ -82,6 +84,30 @@ class Program : public amd::HeapObject {
TYPE_INTERMEDIATE // intermediate
} type_t;
//! type of the input file
typedef enum {
FILE_TYPE_DEFAULT = 0,
FILE_TYPE_OPENCL = 1,
FILE_TYPE_LLVMIR_TEXT = 2,
FILE_TYPE_LLVMIR_BINARY = 3,
FILE_TYPE_SPIR_TEXT = 4,
FILE_TYPE_SPIR_BINARY = 5,
FILE_TYPE_AMDIL_TEXT = 6,
FILE_TYPE_AMDIL_BINARY = 7,
FILE_TYPE_HSAIL_TEXT = 8,
FILE_TYPE_HSAIL_BINARY = 9,
FILE_TYPE_X86_TEXT = 10,
FILE_TYPE_X86_BINARY = 11,
FILE_TYPE_CG = 12,
FILE_TYPE_SOURCE = 13,
FILE_TYPE_ISA = 14,
FILE_TYPE_HEADER = 15,
FILE_TYPE_RSLLVMIR_BINARY = 16,
FILE_TYPE_SPIRV_BINARY = 17,
FILE_TYPE_ASM_TEXT = 18,
FILE_TYPE_LAST = 19
} file_type_t;
private:
//! The device target for this binary.
amd::SharedReference<amd::Device> device_;
@@ -109,15 +135,19 @@ class Program : public amd::HeapObject {
std::string linkOptions_; //!< link options.
//!< the option arg passed in to clCompileProgram(), clLinkProgram(),
//! or clBuildProgram(), whichever is called last
#if defined(WITH_COMPILER_LIB)
aclBinaryOptions binOpts_; //!< Binary options to create aclBinary
aclBinary* binaryElf_; //!< Binary for the new compiler library
#endif
std::string lastBuildOptionsArg_;
mutable std::string buildLog_; //!< build log.
int32_t buildStatus_; //!< build status.
int32_t buildError_; //!< build error
#if defined(WITH_COMPILER_LIB)
aclTargetInfo info_; //!< The info target for this binary.
#endif
size_t globalVariableTotalSize_;
amd::option::Options* programOptions_;
@@ -201,8 +231,10 @@ class Program : public amd::HeapObject {
size_t globalVariableTotalSize() const { return globalVariableTotalSize_; }
#if defined(WITH_COMPILER_LIB)
//! Returns the aclBinary associated with the program
aclBinary* binaryElf() const { return static_cast<aclBinary*>(binaryElf_); }
#endif
//! Returns TRUE if the program just compiled
bool isNull() const { return isNull_; }
@@ -288,8 +320,10 @@ class Program : public amd::HeapObject {
//! Release the Binary
void releaseClBinary();
#if defined(WITH_COMPILER_LIB)
//! return target info
virtual const aclTargetInfo& info() = 0;
#endif
virtual bool setKernels(
amd::option::Options* options, void* binary, size_t binSize,
@@ -314,13 +348,13 @@ class Program : public amd::HeapObject {
* also returns completeStages in a vector, which contains at least ACL_TYPE_DEFAULT,
* sets needOptionsCheck to true if options check is needed to decide whether or not to recompile
*/
aclType getCompilationStagesFromBinary(
std::vector<aclType>& completeStages,
file_type_t getCompilationStagesFromBinary(
std::vector<file_type_t>& completeStages,
bool& needOptionsCheck);
/* \brief Returns the next stage to compile from, based on sections and options in binary
*/
aclType getNextCompilationStageFromBinary(amd::option::Options* options);
file_type_t getNextCompilationStageFromBinary(amd::option::Options* options);
//! Finds the total size of all global variables in the program
bool FindGlobalVarSize(void* binary, size_t binSize);
rocclr/device/pal/paldevice.cpp
+2
Просмотреть файл
@@ -154,7 +154,9 @@ Util::GenericAllocator NullDevice::allocator_;
char* Device::platformObj_;
Pal::IPlatform* Device::platform_;
#if defined(WITH_COMPILER_LIB)
NullDevice::Compiler* NullDevice::compiler_;
#endif
AppProfile Device::appProfile_;
NullDevice::NullDevice() : amd::Device(), ipLevel_(Pal::GfxIpLevel::None), palName_(nullptr) {}
rocclr/device/pal/paldevice.hpp
+4
Просмотреть файл
@@ -53,10 +53,14 @@ namespace pal {
//! A nil device object
class NullDevice : public amd::Device {
protected:
#if defined(WITH_COMPILER_LIB)
static Compiler* compiler_;
#endif
public:
#if defined(WITH_COMPILER_LIB)
Compiler* compiler() const { return compiler_; }
#endif
public:
static bool init(void);
rocclr/device/pal/palprogram.hpp
+2
Просмотреть файл
@@ -202,7 +202,9 @@ class HSAILProgram : public device::Program {
virtual bool createBinary(amd::option::Options* options);
#if defined(WITH_COMPILER_LIB)
virtual const aclTargetInfo& info();
#endif
virtual bool setKernels(amd::option::Options* options, void* binary, size_t binSize,
amd::Os::FileDesc fdesc = amd::Os::FDescInit(), size_t foffset = 0,
rocclr/device/rocm/CMakeLists.txt
-1
Просмотреть файл
@@ -46,7 +46,6 @@ target_include_directories(oclrocm
PRIVATE
${OPENCL_DIR}
${PROJECT_SOURCE_DIR}/compiler/lib
${PROJECT_SOURCE_DIR}/compiler/lib/include
${PROJECT_SOURCE_DIR}/compiler/lib/backends/common
${PROJECT_SOURCE_DIR}/elf
${CMAKE_CURRENT_BINARY_DIR}
rocclr/device/rocm/rocdevice.cpp
-26
Просмотреть файл
@@ -94,7 +94,6 @@ extern const char* BlitSourceCode;
} // namespace device
namespace roc {
amd::Device::Compiler* NullDevice::compilerHandle_;
bool roc::Device::isHsaInitialized_ = false;
std::vector<hsa_agent_t> roc::Device::gpu_agents_;
std::vector<AgentInfo> roc::Device::cpu_agents_;
@@ -264,35 +263,10 @@ Device::~Device() {
}
bool NullDevice::initCompiler(bool isOffline) {
#if defined(WITH_COMPILER_LIB)
// Initialize the compiler handle if has already not been initialized
// This is destroyed in Device::teardown
acl_error error;
if (!compilerHandle_) {
aclCompilerOptions opts = {
sizeof(aclCompilerOptions_0_8), "libamdoclcl64.so",
NULL, NULL, NULL, NULL, NULL, NULL
};
compilerHandle_ = aclCompilerInit(&opts, &error);
if (!GPU_ENABLE_LC && error != ACL_SUCCESS) {
LogError("Error initializing the compiler handle");
return false;
}
}
#endif // defined(WITH_COMPILER_LIB)
return true;
}
bool NullDevice::destroyCompiler() {
#if defined(WITH_COMPILER_LIB)
if (compilerHandle_ != nullptr) {
acl_error error = aclCompilerFini(compilerHandle_);
if (error != ACL_SUCCESS) {
LogError("Error closing the compiler");
return false;
}
}
#endif // defined(WITH_COMPILER_LIB)
return true;
}
rocclr/device/rocm/rocdevice.hpp
-4
Просмотреть файл
@@ -120,8 +120,6 @@ class NullDevice : public amd::Device {
//! Destructor for the Null device
virtual ~NullDevice();
Compiler* compiler() const { return compilerHandle_; }
const Settings& settings() const { return static_cast<Settings&>(*settings_); }
//! Construct an HSAIL program object from the ELF assuming it is valid
@@ -234,8 +232,6 @@ class NullDevice : public amd::Device {
static bool initCompiler(bool isOffline);
//! destroy compiler instance and handle
static bool destroyCompiler();
//! Handle to the the compiler
static Compiler* compilerHandle_;
private:
static constexpr bool offlineDevice_ = true;
rocclr/device/rocm/rockernel.cpp
-113
Просмотреть файл
@@ -174,118 +174,5 @@ bool LightningKernel::init() {
}
#endif // defined(USE_COMGR_LIBRARY)
#if defined(WITH_COMPILER_LIB)
bool HSAILKernel::init() {
acl_error errorCode;
// Pull out metadata from the ELF
size_t sizeOfArgList;
aclCompiler* compileHandle = program()->rocDevice().compiler();
std::string openClKernelName("&__OpenCL_" + name() + "_kernel");
errorCode = aclQueryInfo(compileHandle, program()->binaryElf(), RT_ARGUMENT_ARRAY,
openClKernelName.c_str(), nullptr, &sizeOfArgList);
if (errorCode != ACL_SUCCESS) {
LogPrintfError("Query Info failed with error code: %d \n", errorCode);
return false;
}
std::unique_ptr<char[]> argList(new char[sizeOfArgList]);
errorCode = aclQueryInfo(compileHandle, program()->binaryElf(), RT_ARGUMENT_ARRAY,
openClKernelName.c_str(), argList.get(), &sizeOfArgList);
if (errorCode != ACL_SUCCESS) {
LogPrintfError("Query Info failed with error code: %d \n", errorCode);
return false;
}
// Set the argList
InitParameters((const aclArgData*)argList.get(), KernargSegmentByteSize());
// Set the workgroup information for the kernel
memset(&workGroupInfo_, 0, sizeof(workGroupInfo_));
workGroupInfo_.availableLDSSize_ = program()->rocDevice().info().localMemSizePerCU_;
assert(workGroupInfo_.availableLDSSize_ > 0);
workGroupInfo_.availableSGPRs_ = 104;
workGroupInfo_.availableVGPRs_ = 256;
size_t sizeOfWorkGroupSize;
errorCode = aclQueryInfo(compileHandle, program()->binaryElf(), RT_WORK_GROUP_SIZE,
openClKernelName.c_str(), nullptr, &sizeOfWorkGroupSize);
if (errorCode != ACL_SUCCESS) {
LogPrintfError("Query Info failed with error code: %d \n", errorCode);
return false;
}
errorCode = aclQueryInfo(compileHandle, program()->binaryElf(), RT_WORK_GROUP_SIZE,
openClKernelName.c_str(), workGroupInfo_.compileSize_,
&sizeOfWorkGroupSize);
if (errorCode != ACL_SUCCESS) {
LogPrintfError("Query Info failed with error code: %d \n ", errorCode);
return false;
}
uint32_t wavefront_size = 0;
hsa_status_t hsaStatus = hsa_agent_get_info(program()->rocDevice().getBackendDevice(),
HSA_AGENT_INFO_WAVEFRONT_SIZE, &wavefront_size);
if (HSA_STATUS_SUCCESS != hsaStatus) {
DevLogPrintfError("Could not get Wave Info Size: %d, failed with hsa_status: %d \n",
errorCode, hsaStatus);
return false;
}
assert(wavefront_size > 0);
// Setting it the same as used LDS.
workGroupInfo_.localMemSize_ = workgroupGroupSegmentByteSize_;
workGroupInfo_.privateMemSize_ = workitemPrivateSegmentByteSize_;
workGroupInfo_.usedLDSSize_ = workgroupGroupSegmentByteSize_;
workGroupInfo_.preferredSizeMultiple_ = wavefront_size;
// Query kernel header object to initialize the number of
// SGPR's and VGPR's used by the kernel
const void* kernelHostPtr = nullptr;
if (Device::loaderQueryHostAddress(reinterpret_cast<const void*>(kernelCodeHandle_),
&kernelHostPtr) == HSA_STATUS_SUCCESS) {
auto akc = reinterpret_cast<const amd_kernel_code_t*>(kernelHostPtr);
workGroupInfo_.usedSGPRs_ = akc->wavefront_sgpr_count;
workGroupInfo_.usedVGPRs_ = akc->workitem_vgpr_count;
} else {
workGroupInfo_.usedSGPRs_ = 0;
workGroupInfo_.usedVGPRs_ = 0;
}
workGroupInfo_.usedStackSize_ = 0;
workGroupInfo_.wavefrontPerSIMD_ = program()->rocDevice().info().maxWorkItemSizes_[0] / wavefront_size;
workGroupInfo_.wavefrontSize_ = wavefront_size;
if (workGroupInfo_.compileSize_[0] != 0) {
workGroupInfo_.size_ = workGroupInfo_.compileSize_[0] * workGroupInfo_.compileSize_[1] *
workGroupInfo_.compileSize_[2];
} else {
workGroupInfo_.size_ = program()->rocDevice().info().preferredWorkGroupSize_;
}
// Pull out printf metadata from the ELF
size_t sizeOfPrintfList;
errorCode = aclQueryInfo(compileHandle, program()->binaryElf(), RT_GPU_PRINTF_ARRAY,
openClKernelName.c_str(), nullptr, &sizeOfPrintfList);
if (errorCode != ACL_SUCCESS) {
LogPrintfError("Query Info failed with error code: %d \n", errorCode);
return false;
}
// Make sure kernel has any printf info
if (0 != sizeOfPrintfList) {
std::unique_ptr<char[]> aclPrintfList(new char[sizeOfPrintfList]);
if (!aclPrintfList) {
return false;
}
errorCode = aclQueryInfo(compileHandle, program()->binaryElf(),
RT_GPU_PRINTF_ARRAY, openClKernelName.c_str(),
aclPrintfList.get(), &sizeOfPrintfList);
if (errorCode != ACL_SUCCESS) {
return false;
}
// Set the Printf List
InitPrintf(reinterpret_cast<aclPrintfFmt*>(aclPrintfList.get()));
}
return true;
}
#endif // defined(WITH_COMPILER_LIB)
} // namespace roc
#endif // WITHOUT_HSA_BACKEND
rocclr/device/rocm/rockernel.hpp
-1
Просмотреть файл
@@ -21,7 +21,6 @@
#pragma once
#include <memory>
#include "acl.h"
#include "rocprogram.hpp"
#include "top.hpp"
#include "rocprintf.hpp"
rocclr/device/rocm/rocprogram.cpp
-210
Просмотреть файл
@@ -25,7 +25,6 @@
#include "utils/options.hpp"
#include "rockernel.hpp"
#include "utils/bif_section_labels.hpp"
#include "amd_hsa_kernel_code.h"
#include <string>
@@ -38,30 +37,6 @@
namespace roc {
#if defined(WITH_COMPILER_LIB)
static hsa_status_t GetKernelNamesCallback(hsa_executable_t exec, hsa_agent_t agent,
hsa_executable_symbol_t symbol, void* data) {
std::vector<std::string>* symNameList = reinterpret_cast<std::vector<std::string>*>(data);
hsa_symbol_kind_t sym_type;
hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &sym_type);
if (sym_type == HSA_SYMBOL_KIND_KERNEL) {
uint32_t len;
hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH, &len);
char* symName = (char*)alloca(len + 1);
hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME, symName);
symName[len] = '\0';
std::string kernelName(symName);
symNameList->push_back(kernelName);
}
return HSA_STATUS_SUCCESS;
}
#endif
static inline const char* hsa_strerror(hsa_status_t status) {
const char* str = nullptr;
if (hsa_status_string(status, &str) == HSA_STATUS_SUCCESS) {
@@ -238,199 +213,14 @@ HSAILProgram::HSAILProgram(roc::NullDevice& device, amd::Program& owner)
: roc::Program(device, owner) {}
HSAILProgram::~HSAILProgram() {
#if defined(WITH_COMPILER_LIB)
acl_error error;
// Free the elf binary
if (binaryElf_ != nullptr) {
error = aclBinaryFini(binaryElf_);
if (error != ACL_SUCCESS) {
LogWarning("Error while destroying the acl binary \n");
}
}
#endif // defined(WITH_COMPILER_LIB)
}
bool HSAILProgram::saveBinaryAndSetType(type_t type) {
#if defined(WITH_COMPILER_LIB)
void* rawBinary;
size_t size;
// Write binary to memory
if (aclWriteToMem(binaryElf_, &rawBinary, &size) != ACL_SUCCESS) {
buildLog_ += "Failed to write binary to memory \n";
return false;
}
clBinary()->saveBIFBinary((char*)rawBinary, size);
// Set the type of binary
setType(type);
// Free memory containing rawBinary
binaryElf_->binOpts.dealloc(rawBinary);
#endif // defined(WITH_COMPILER_LIB)
return true;
}
bool HSAILProgram::setKernels(amd::option::Options* options, void* binary, size_t binSize,
amd::Os::FileDesc fdesc, size_t foffset, std::string uri) {
#if defined(WITH_COMPILER_LIB)
// Stop compilation if it is an offline device - HSA runtime does not
// support ISA compiled offline
if (!device().isOnline()) {
return true;
}
size_t secSize = binSize;
void* data = binary;
// Create an executable.
hsa_status_t status = hsa_executable_create_alt(
HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, nullptr, &hsaExecutable_);
if (status != HSA_STATUS_SUCCESS) {
buildLog_ += "Error: Failed to create executable: ";
buildLog_ += hsa_strerror(status);
buildLog_ += "\n";
return false;
}
// Load the code object.
status = hsa_code_object_reader_create_from_memory(data, secSize, &hsaCodeObjectReader_);
if (status != HSA_STATUS_SUCCESS) {
buildLog_ += "Error: AMD HSA Code Object Reader create failed: ";
buildLog_ += hsa_strerror(status);
buildLog_ += "\n";
return false;
}
hsa_agent_t hsaDevice = rocDevice().getBackendDevice();
status = hsa_executable_load_agent_code_object(hsaExecutable_, hsaDevice, hsaCodeObjectReader_,
nullptr, nullptr);
if (status != HSA_STATUS_SUCCESS) {
buildLog_ += "Error: AMD HSA Code Object loading failed: ";
buildLog_ += hsa_strerror(status);
buildLog_ += "\n";
return false;
}
// Freeze the executable.
status = hsa_executable_freeze(hsaExecutable_, nullptr);
if (status != HSA_STATUS_SUCCESS) {
buildLog_ += "Error: Failed to freeze executable: ";
buildLog_ += hsa_strerror(status);
buildLog_ += "\n";
return false;
}
// Get the list of kernels
std::vector<std::string> kernelNameList;
status = hsa_executable_iterate_agent_symbols(hsaExecutable_, hsaDevice, GetKernelNamesCallback,
(void*)&kernelNameList);
if (status != HSA_STATUS_SUCCESS) {
buildLog_ += "Error: Failed to get kernel names: ";
buildLog_ += hsa_strerror(status);
buildLog_ += "\n";
return false;
}
for (auto& kernelName : kernelNameList) {
// Query symbol handle for this symbol.
hsa_executable_symbol_t kernelSymbol;
status = hsa_executable_get_symbol_by_name(hsaExecutable_, kernelName.c_str(), &hsaDevice,
&kernelSymbol);
if (status != HSA_STATUS_SUCCESS) {
buildLog_ += "Error: Failed to get executable symbol: ";
buildLog_ += hsa_strerror(status);
buildLog_ += "\n";
return false;
}
// Query code handle for this symbol.
uint64_t kernelCodeHandle;
status = hsa_executable_symbol_get_info(kernelSymbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT,
&kernelCodeHandle);
if (status != HSA_STATUS_SUCCESS) {
buildLog_ += "Error: Failed to get executable symbol info: ";
buildLog_ += hsa_strerror(status);
buildLog_ += "\n";
return false;
}
std::string openclKernelName = kernelName;
// Strip the opencl and kernel name
kernelName = kernelName.substr(strlen("&__OpenCL_"), kernelName.size());
kernelName = kernelName.substr(0, kernelName.size() - strlen("_kernel"));
aclMetadata md;
md.numHiddenKernelArgs = 0;
size_t sizeOfnumHiddenKernelArgs = sizeof(md.numHiddenKernelArgs);
acl_error errorCode = aclQueryInfo(device().compiler(), binaryElf_,
RT_NUM_KERNEL_HIDDEN_ARGS, openclKernelName.c_str(),
&md.numHiddenKernelArgs, &sizeOfnumHiddenKernelArgs);
if (errorCode != ACL_SUCCESS) {
buildLog_ +=
"Error while Finalization phase: Kernel extra arguments count querying from the ELF "
"failed\n";
return false;
}
uint32_t workgroupGroupSegmentByteSize;
status = hsa_executable_symbol_get_info(kernelSymbol,
HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE,
&workgroupGroupSegmentByteSize);
if (status != HSA_STATUS_SUCCESS) {
buildLog_ += "Error: Failed to get group segment size info: ";
buildLog_ += hsa_strerror(status);
buildLog_ += "\n";
return false;
}
uint32_t workitemPrivateSegmentByteSize;
status = hsa_executable_symbol_get_info(kernelSymbol,
HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE,
&workitemPrivateSegmentByteSize);
if (status != HSA_STATUS_SUCCESS) {
buildLog_ += "Error: Failed to get private segment size info: ";
buildLog_ += hsa_strerror(status);
buildLog_ += "\n";
return false;
}
uint32_t kernargSegmentByteSize;
status = hsa_executable_symbol_get_info(kernelSymbol,
HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE,
&kernargSegmentByteSize);
if (status != HSA_STATUS_SUCCESS) {
buildLog_ += "Error: Failed to get kernarg segment size info: ";
buildLog_ += hsa_strerror(status);
buildLog_ += "\n";
return false;
}
uint32_t kernargSegmentAlignment;
status = hsa_executable_symbol_get_info(
kernelSymbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_ALIGNMENT,
&kernargSegmentAlignment);
if (status != HSA_STATUS_SUCCESS) {
buildLog_ += "Error: Failed to get kernarg segment alignment info: ";
buildLog_ += hsa_strerror(status);
buildLog_ += "\n";
return false;
}
Kernel* aKernel = new roc::HSAILKernel(kernelName, this, kernelCodeHandle,
workgroupGroupSegmentByteSize,
workitemPrivateSegmentByteSize,
kernargSegmentByteSize, kernargSegmentAlignment);
if (!aKernel->init()) {
buildLog_ += "Error: Kernel Init Failed ";
buildLog_ += "\n";
return false;
}
aKernel->setUniformWorkGroupSize(options->oVariables->UniformWorkGroupSize);
aKernel->setInternalKernelFlag(compileOptions_.find("-cl-internal-kernel") !=
std::string::npos);
kernels()[kernelName] = aKernel;
}
#endif // defined(WITH_COMPILER_LIB)
return true;
}
rocclr/device/rocm/rocprogram.hpp
-3
Просмотреть файл
@@ -22,7 +22,6 @@
#ifndef WITHOUT_HSA_BACKEND
#include "acl.h"
#include <string>
#include <sstream>
#include <fstream>
@@ -74,8 +73,6 @@ class Program : public device::Program {
);
virtual bool createBinary(amd::option::Options* options) = 0;
virtual const aclTargetInfo& info() { return info_; }
protected:
//! Disable default copy constructor
Program(const Program&) = delete;
rocclr/device/rocm/rocsettings.hpp
-2
Просмотреть файл
@@ -22,8 +22,6 @@
#ifndef WITHOUT_HSA_BACKEND
#include "library.hpp"
/*! \addtogroup HSA OCL Stub Implementation
* @{
*/