foreman
eb852bdb53
SWDEV-144570 - Disbaling __hipHostRegister implementation to unblock pytorch tests in PAL. Affected files ... ... //depot/stg/opencl/drivers/opencl/api/hip/hip_platform.cpp#24 edit
611 líneas
19 KiB
C++
611 líneas
19 KiB
C++
/*
|
|
Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
|
|
|
|
Permission is hereby granted, free of charge, to any person obtaining a copy
|
|
of this software and associated documentation files (the "Software"), to deal
|
|
in the Software without restriction, including without limitation the rights
|
|
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
|
copies of the Software, and to permit persons to whom the Software is
|
|
furnished to do so, subject to the following conditions:
|
|
|
|
The above copyright notice and this permission notice shall be included in
|
|
all copies or substantial portions of the Software.
|
|
|
|
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
|
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
|
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
|
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
|
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
|
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
|
THE SOFTWARE.
|
|
*/
|
|
|
|
#include <hip/hip_runtime.h>
|
|
|
|
#include "hip_internal.hpp"
|
|
#include "platform/program.hpp"
|
|
#include "platform/runtime.hpp"
|
|
|
|
#include <unordered_map>
|
|
#include "elfio.hpp"
|
|
|
|
constexpr unsigned __hipFatMAGIC2 = 0x48495046; // "HIPF"
|
|
|
|
thread_local std::stack<ihipExec_t> execStack_;
|
|
PlatformState* PlatformState::platform_ = new PlatformState();
|
|
|
|
struct __CudaFatBinaryWrapper {
|
|
unsigned int magic;
|
|
unsigned int version;
|
|
void* binary;
|
|
void* dummy1;
|
|
};
|
|
|
|
#define CLANG_OFFLOAD_BUNDLER_MAGIC_STR "__CLANG_OFFLOAD_BUNDLE__"
|
|
#define HIP_AMDGCN_AMDHSA_TRIPLE "hip-amdgcn-amd-amdhsa"
|
|
#define HCC_AMDGCN_AMDHSA_TRIPLE "hcc-amdgcn-amd-amdhsa-"
|
|
|
|
struct __ClangOffloadBundleDesc {
|
|
uint64_t offset;
|
|
uint64_t size;
|
|
uint64_t tripleSize;
|
|
const char triple[1];
|
|
};
|
|
|
|
struct __ClangOffloadBundleHeader {
|
|
const char magic[sizeof(CLANG_OFFLOAD_BUNDLER_MAGIC_STR) - 1];
|
|
uint64_t numBundles;
|
|
__ClangOffloadBundleDesc desc[1];
|
|
};
|
|
|
|
hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes,
|
|
hipModule_t hmod, const char* name);
|
|
|
|
extern "C" std::vector<hipModule_t>* __hipRegisterFatBinary(const void* data)
|
|
{
|
|
HIP_INIT();
|
|
|
|
if(g_devices.empty()) {
|
|
return nullptr;
|
|
}
|
|
const __CudaFatBinaryWrapper* fbwrapper = reinterpret_cast<const __CudaFatBinaryWrapper*>(data);
|
|
if (fbwrapper->magic != __hipFatMAGIC2 || fbwrapper->version != 1) {
|
|
return nullptr;
|
|
}
|
|
std::string magic((char*)fbwrapper->binary, sizeof(CLANG_OFFLOAD_BUNDLER_MAGIC_STR) - 1);
|
|
if (magic.compare(CLANG_OFFLOAD_BUNDLER_MAGIC_STR)) {
|
|
return nullptr;
|
|
}
|
|
|
|
auto programs = new std::vector<hipModule_t>{g_devices.size()};
|
|
|
|
const auto obheader = reinterpret_cast<const __ClangOffloadBundleHeader*>(fbwrapper->binary);
|
|
const auto* desc = &obheader->desc[0];
|
|
for (uint64_t i = 0; i < obheader->numBundles; ++i,
|
|
desc = reinterpret_cast<const __ClangOffloadBundleDesc*>(
|
|
reinterpret_cast<uintptr_t>(&desc->triple[0]) + desc->tripleSize)) {
|
|
|
|
std::string triple(desc->triple, sizeof(HIP_AMDGCN_AMDHSA_TRIPLE) - 1);
|
|
if (triple.compare(HIP_AMDGCN_AMDHSA_TRIPLE))
|
|
continue;
|
|
|
|
std::string target(desc->triple + sizeof(HIP_AMDGCN_AMDHSA_TRIPLE),
|
|
desc->tripleSize - sizeof(HIP_AMDGCN_AMDHSA_TRIPLE));
|
|
|
|
const void *image = reinterpret_cast<const void*>(
|
|
reinterpret_cast<uintptr_t>(obheader) + desc->offset);
|
|
size_t size = desc->size;
|
|
|
|
for (size_t dev = 0; dev < g_devices.size(); ++dev) {
|
|
amd::Context* ctx = g_devices[dev];
|
|
|
|
if (target.compare(ctx->devices()[0]->info().name_)) {
|
|
continue;
|
|
}
|
|
|
|
amd::Program* program = new amd::Program(*ctx);
|
|
if (program == nullptr) {
|
|
return nullptr;
|
|
}
|
|
if (CL_SUCCESS == program->addDeviceProgram(*ctx->devices()[0], image, size) &&
|
|
CL_SUCCESS == program->build(ctx->devices(), nullptr, nullptr, nullptr)) {
|
|
programs->at(dev) = reinterpret_cast<hipModule_t>(as_cl(program));
|
|
}
|
|
}
|
|
}
|
|
|
|
return programs;
|
|
}
|
|
|
|
PlatformState::RegisteredVar::RegisteredVar(char* hostVar, size_t size, hipDeviceptr_t devicePtr)
|
|
: hostVar_(hostVar), size_(size), devicePtr_(devicePtr) {
|
|
amd::Memory* amd_mem_obj = nullptr;
|
|
uint32_t flags = 0;
|
|
|
|
/* Create an amd Memory object for the pointer */
|
|
amd_mem_obj
|
|
= new (*hip::getCurrentContext()) amd::Buffer(*hip::getCurrentContext(), flags, size, devicePtr_);
|
|
|
|
if (amd_mem_obj == nullptr) {
|
|
LogError("[OCL] failed to create a mem object!");
|
|
}
|
|
|
|
if (!amd_mem_obj->create(nullptr)) {
|
|
LogError("[OCL] failed to create a svm hidden buffer!");
|
|
amd_mem_obj->release();
|
|
}
|
|
|
|
/* Add the memory to the MemObjMap */
|
|
amd::MemObjMap::AddMemObj(devicePtr_, amd_mem_obj);
|
|
}
|
|
|
|
void PlatformState::registerVar(const char* hostvar,
|
|
const std::vector<RegisteredVar>& rvar) {
|
|
amd::ScopedLock lock(lock_);
|
|
vars_.insert(std::make_pair(hostvar, rvar));
|
|
}
|
|
|
|
void PlatformState::registerFunction(const void* hostFunction,
|
|
const std::vector<hipFunction_t>& funcs) {
|
|
amd::ScopedLock lock(lock_);
|
|
functions_.insert(std::make_pair(hostFunction, funcs));
|
|
}
|
|
|
|
hipFunction_t PlatformState::getFunc(const void* hostFunction, int deviceId) {
|
|
amd::ScopedLock lock(lock_);
|
|
const auto it = functions_.find(hostFunction);
|
|
if (it != functions_.cend()) {
|
|
return it->second[deviceId];
|
|
} else {
|
|
return nullptr;
|
|
}
|
|
}
|
|
|
|
bool PlatformState::getGlobalVar(const void* hostVar, int deviceId,
|
|
hipDeviceptr_t* dev_ptr, size_t* size_ptr) {
|
|
amd::ScopedLock lock(lock_);
|
|
const auto it = vars_.find(hostVar);
|
|
if (it != vars_.cend()) {
|
|
*size_ptr = it->second[deviceId].getvarsize();
|
|
*dev_ptr = it->second[deviceId].getdeviceptr();
|
|
return true;
|
|
} else {
|
|
return false;
|
|
}
|
|
}
|
|
|
|
void PlatformState::setupArgument(const void *arg, size_t size, size_t offset) {
|
|
auto& arguments = execStack_.top().arguments_;
|
|
|
|
if (arguments.size() < offset + size) {
|
|
arguments.resize(offset + size);
|
|
}
|
|
|
|
::memcpy(&arguments[offset], arg, size);
|
|
}
|
|
|
|
void PlatformState::configureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem,
|
|
hipStream_t stream) {
|
|
execStack_.push(ihipExec_t{gridDim, blockDim, sharedMem, stream});
|
|
}
|
|
|
|
void PlatformState::popExec(ihipExec_t& exec) {
|
|
exec = std::move(execStack_.top());
|
|
execStack_.pop();
|
|
}
|
|
|
|
extern "C" void __hipRegisterFunction(
|
|
std::vector<hipModule_t>* modules,
|
|
const void* hostFunction,
|
|
char* deviceFunction,
|
|
const char* deviceName,
|
|
unsigned int threadLimit,
|
|
uint3* tid,
|
|
uint3* bid,
|
|
dim3* blockDim,
|
|
dim3* gridDim,
|
|
int* wSize)
|
|
{
|
|
HIP_INIT();
|
|
|
|
std::vector<hipFunction_t> functions{g_devices.size()};
|
|
|
|
for (size_t deviceId=0; deviceId < g_devices.size(); ++deviceId) {
|
|
hipFunction_t function = nullptr;
|
|
if (hipSuccess == hipModuleGetFunction(&function, modules->at(deviceId), deviceName) &&
|
|
function != nullptr) {
|
|
functions[deviceId] = function;
|
|
}
|
|
else {
|
|
// tprintf(DB_FB, "__hipRegisterFunction cannot find kernel %s for"
|
|
// " device %d\n", deviceName, deviceId);
|
|
}
|
|
}
|
|
|
|
PlatformState::instance().registerFunction(hostFunction, functions);
|
|
}
|
|
|
|
// Registers a device-side global variable.
|
|
// For each global variable in device code, there is a corresponding shadow
|
|
// global variable in host code. The shadow host variable is used to keep
|
|
// track of the value of the device side global variable between kernel
|
|
// executions.
|
|
extern "C" void __hipRegisterVar(
|
|
std::vector<hipModule_t>* modules, // The device modules containing code object
|
|
char* var, // The shadow variable in host code
|
|
char* hostVar, // Variable name in host code
|
|
char* deviceVar, // Variable name in device code
|
|
int ext, // Whether this variable is external
|
|
int size, // Size of the variable
|
|
int constant, // Whether this variable is constant
|
|
int global) // Unknown, always 0
|
|
{
|
|
#if 0
|
|
HIP_INIT();
|
|
|
|
size_t sym_size = 0;
|
|
std::vector<PlatformState::RegisteredVar> global_vars{g_devices.size()};
|
|
|
|
for (size_t deviceId=0; deviceId < g_devices.size(); ++deviceId) {
|
|
hipDeviceptr_t device_ptr = nullptr;
|
|
if((hipSuccess == hipModuleGetGlobal(&device_ptr, &sym_size, modules->at(deviceId),
|
|
hostVar)) && (device_ptr != nullptr)) {
|
|
|
|
if (static_cast<size_t>(size) != sym_size) {
|
|
LogError("[OCL] Size Mismatch with the HSA Symbol retrieved \n");
|
|
}
|
|
|
|
global_vars[deviceId] = PlatformState::RegisteredVar(hostVar, sym_size, device_ptr);
|
|
|
|
} else {
|
|
LogError("[OCL] __hipRegisterVar cannot find kernel for device \n");
|
|
}
|
|
}
|
|
|
|
PlatformState::instance().registerVar(hostVar, global_vars);
|
|
#endif
|
|
}
|
|
|
|
extern "C" void __hipUnregisterFatBinary(std::vector<hipModule_t>* modules)
|
|
{
|
|
HIP_INIT();
|
|
|
|
std::for_each(modules->begin(), modules->end(), [](hipModule_t module){
|
|
if (module != nullptr) {
|
|
as_amd(reinterpret_cast<cl_program>(module))->release();
|
|
}
|
|
});
|
|
delete modules;
|
|
}
|
|
|
|
extern "C" hipError_t hipConfigureCall(
|
|
dim3 gridDim,
|
|
dim3 blockDim,
|
|
size_t sharedMem,
|
|
hipStream_t stream)
|
|
{
|
|
HIP_INIT_API(gridDim, blockDim, sharedMem, stream);
|
|
|
|
PlatformState::instance().configureCall(gridDim, blockDim, sharedMem, stream);
|
|
|
|
HIP_RETURN(hipSuccess);
|
|
}
|
|
|
|
extern "C" hipError_t hipSetupArgument(
|
|
const void *arg,
|
|
size_t size,
|
|
size_t offset)
|
|
{
|
|
HIP_INIT_API(arg, size, offset);
|
|
|
|
PlatformState::instance().setupArgument(arg, size, offset);
|
|
|
|
HIP_RETURN(hipSuccess);
|
|
}
|
|
|
|
extern "C" hipError_t hipLaunchByPtr(const void *hostFunction)
|
|
{
|
|
HIP_INIT_API(hostFunction);
|
|
|
|
int deviceId = ihipGetDevice();
|
|
hipFunction_t func = PlatformState::instance().getFunc(hostFunction, deviceId);
|
|
if (func == nullptr) {
|
|
HIP_RETURN(hipErrorUnknown);
|
|
}
|
|
|
|
ihipExec_t exec;
|
|
PlatformState::instance().popExec(exec);
|
|
|
|
size_t size = exec.arguments_.size();
|
|
void *extra[] = {
|
|
HIP_LAUNCH_PARAM_BUFFER_POINTER, &exec.arguments_[0],
|
|
HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
|
|
HIP_LAUNCH_PARAM_END
|
|
};
|
|
|
|
HIP_RETURN(hipModuleLaunchKernel(func,
|
|
exec.gridDim_.x, exec.gridDim_.y, exec.gridDim_.z,
|
|
exec.blockDim_.x, exec.blockDim_.y, exec.blockDim_.z,
|
|
exec.sharedMem_, exec.hStream_, nullptr, extra));
|
|
}
|
|
|
|
hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName) {
|
|
size_t size = 0;
|
|
if(!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), devPtr, &size)) {
|
|
HIP_RETURN(hipErrorUnknown);
|
|
}
|
|
HIP_RETURN(hipSuccess);
|
|
}
|
|
|
|
hipError_t hipGetSymbolSize(size_t* sizePtr, const void* symbolName) {
|
|
hipDeviceptr_t devPtr = nullptr;
|
|
if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), &devPtr, sizePtr)) {
|
|
HIP_RETURN(hipErrorUnknown);
|
|
}
|
|
HIP_RETURN(hipSuccess);
|
|
}
|
|
|
|
#if defined(ATI_OS_LINUX)
|
|
|
|
namespace hip_impl {
|
|
|
|
struct dl_phdr_info {
|
|
ELFIO::Elf64_Addr dlpi_addr;
|
|
const char *dlpi_name;
|
|
const ELFIO::Elf64_Phdr *dlpi_phdr;
|
|
ELFIO::Elf64_Half dlpi_phnum;
|
|
};
|
|
|
|
extern "C" int dl_iterate_phdr(
|
|
int (*callback) (struct dl_phdr_info *info, size_t size, void *data), void *data
|
|
);
|
|
|
|
struct Symbol {
|
|
std::string name;
|
|
ELFIO::Elf64_Addr value = 0;
|
|
ELFIO::Elf_Xword size = 0;
|
|
ELFIO::Elf_Half sect_idx = 0;
|
|
uint8_t bind = 0;
|
|
uint8_t type = 0;
|
|
uint8_t other = 0;
|
|
};
|
|
|
|
inline Symbol read_symbol(const ELFIO::symbol_section_accessor& section, unsigned int idx) {
|
|
assert(idx < section.get_symbols_num());
|
|
|
|
Symbol r;
|
|
section.get_symbol(idx, r.name, r.value, r.size, r.bind, r.type, r.sect_idx, r.other);
|
|
|
|
return r;
|
|
}
|
|
|
|
template <typename P>
|
|
inline ELFIO::section* find_section_if(ELFIO::elfio& reader, P p) {
|
|
const auto it = find_if(reader.sections.begin(), reader.sections.end(), std::move(p));
|
|
|
|
return it != reader.sections.end() ? *it : nullptr;
|
|
}
|
|
|
|
std::vector<std::pair<uintptr_t, std::string>> function_names_for(const ELFIO::elfio& reader,
|
|
ELFIO::section* symtab) {
|
|
std::vector<std::pair<uintptr_t, std::string>> r;
|
|
ELFIO::symbol_section_accessor symbols{reader, symtab};
|
|
|
|
for (auto i = 0u; i != symbols.get_symbols_num(); ++i) {
|
|
auto tmp = read_symbol(symbols, i);
|
|
|
|
if (tmp.type == STT_FUNC && tmp.sect_idx != SHN_UNDEF && !tmp.name.empty()) {
|
|
r.emplace_back(tmp.value, tmp.name);
|
|
}
|
|
}
|
|
|
|
return r;
|
|
}
|
|
|
|
const std::vector<std::pair<uintptr_t, std::string>>& function_names_for_process() {
|
|
static constexpr const char self[] = "/proc/self/exe";
|
|
|
|
static std::vector<std::pair<uintptr_t, std::string>> r;
|
|
static std::once_flag f;
|
|
|
|
std::call_once(f, []() {
|
|
ELFIO::elfio reader;
|
|
|
|
if (reader.load(self)) {
|
|
const auto it = find_section_if(
|
|
reader, [](const ELFIO::section* x) { return x->get_type() == SHT_SYMTAB; });
|
|
|
|
if (it) r = function_names_for(reader, it);
|
|
}
|
|
});
|
|
|
|
return r;
|
|
}
|
|
|
|
|
|
const std::unordered_map<uintptr_t, std::string>& function_names()
|
|
{
|
|
static std::unordered_map<uintptr_t, std::string> r{
|
|
function_names_for_process().cbegin(),
|
|
function_names_for_process().cend()};
|
|
static std::once_flag f;
|
|
|
|
std::call_once(f, []() {
|
|
dl_iterate_phdr([](dl_phdr_info* info, size_t, void*) {
|
|
ELFIO::elfio reader;
|
|
|
|
if (reader.load(info->dlpi_name)) {
|
|
const auto it = find_section_if(
|
|
reader, [](const ELFIO::section* x) { return x->get_type() == SHT_SYMTAB; });
|
|
|
|
if (it) {
|
|
auto n = function_names_for(reader, it);
|
|
|
|
for (auto&& f : n) f.first += info->dlpi_addr;
|
|
|
|
r.insert(make_move_iterator(n.begin()), make_move_iterator(n.end()));
|
|
}
|
|
}
|
|
return 0;
|
|
},
|
|
nullptr);
|
|
});
|
|
|
|
return r;
|
|
}
|
|
|
|
std::vector<char> bundles_for_process() {
|
|
static constexpr const char self[] = "/proc/self/exe";
|
|
static constexpr const char kernel_section[] = ".kernel";
|
|
std::vector<char> r;
|
|
|
|
ELFIO::elfio reader;
|
|
|
|
if (reader.load(self)) {
|
|
auto it = find_section_if(
|
|
reader, [](const ELFIO::section* x) { return x->get_name() == kernel_section; });
|
|
|
|
if (it) r.insert(r.end(), it->get_data(), it->get_data() + it->get_size());
|
|
}
|
|
|
|
return r;
|
|
}
|
|
|
|
const std::vector<hipModule_t>& modules() {
|
|
static std::vector<hipModule_t> r;
|
|
static std::once_flag f;
|
|
|
|
std::call_once(f, []() {
|
|
static std::vector<std::vector<char>> bundles{bundles_for_process()};
|
|
|
|
dl_iterate_phdr(
|
|
[](dl_phdr_info* info, std::size_t, void*) {
|
|
ELFIO::elfio tmp;
|
|
if (tmp.load(info->dlpi_name)) {
|
|
const auto it = find_section_if(
|
|
tmp, [](const ELFIO::section* x) { return x->get_name() == ".kernel"; });
|
|
|
|
if (it) bundles.emplace_back(it->get_data(), it->get_data() + it->get_size());
|
|
}
|
|
return 0;
|
|
},
|
|
nullptr);
|
|
|
|
for (auto&& bundle : bundles) {
|
|
if (bundle.empty()) {
|
|
continue;
|
|
}
|
|
std::string magic(&bundle[0], sizeof(CLANG_OFFLOAD_BUNDLER_MAGIC_STR) - 1);
|
|
if (magic.compare(CLANG_OFFLOAD_BUNDLER_MAGIC_STR))
|
|
continue;
|
|
|
|
const auto obheader = reinterpret_cast<const __ClangOffloadBundleHeader*>(&bundle[0]);
|
|
const auto* desc = &obheader->desc[0];
|
|
for (uint64_t i = 0; i < obheader->numBundles; ++i,
|
|
desc = reinterpret_cast<const __ClangOffloadBundleDesc*>(
|
|
reinterpret_cast<uintptr_t>(&desc->triple[0]) + desc->tripleSize)) {
|
|
|
|
std::string triple(desc->triple, sizeof(HCC_AMDGCN_AMDHSA_TRIPLE) - 1);
|
|
if (triple.compare(HCC_AMDGCN_AMDHSA_TRIPLE))
|
|
continue;
|
|
|
|
std::string target(desc->triple + sizeof(HCC_AMDGCN_AMDHSA_TRIPLE),
|
|
desc->tripleSize - sizeof(HCC_AMDGCN_AMDHSA_TRIPLE));
|
|
|
|
if (!target.compare(hip::getCurrentContext()->devices()[0]->info().name_)) {
|
|
hipModule_t module;
|
|
if (hipSuccess == hipModuleLoadData(&module, reinterpret_cast<const void*>(
|
|
reinterpret_cast<uintptr_t>(obheader) + desc->offset)))
|
|
r.push_back(module);
|
|
break;
|
|
}
|
|
}
|
|
}
|
|
});
|
|
|
|
return r;
|
|
}
|
|
|
|
const std::unordered_map<uintptr_t, hipFunction_t>& functions()
|
|
{
|
|
static std::unordered_map<uintptr_t, hipFunction_t> r;
|
|
static std::once_flag f;
|
|
|
|
std::call_once(f, []() {
|
|
for (auto&& function : function_names()) {
|
|
for (auto&& module : modules()) {
|
|
hipFunction_t f;
|
|
if (hipSuccess == hipModuleGetFunction(&f, module, function.second.c_str()))
|
|
r[function.first] = f;
|
|
}
|
|
}
|
|
});
|
|
|
|
return r;
|
|
}
|
|
|
|
|
|
void hipLaunchKernelGGLImpl(
|
|
uintptr_t function_address,
|
|
const dim3& numBlocks,
|
|
const dim3& dimBlocks,
|
|
uint32_t sharedMemBytes,
|
|
hipStream_t stream,
|
|
void** kernarg)
|
|
{
|
|
HIP_INIT();
|
|
|
|
const auto it = functions().find(function_address);
|
|
if (it == functions().cend())
|
|
assert(0);
|
|
|
|
hipModuleLaunchKernel(it->second,
|
|
numBlocks.x, numBlocks.y, numBlocks.z,
|
|
dimBlocks.x, dimBlocks.y, dimBlocks.z,
|
|
sharedMemBytes, stream, nullptr, kernarg);
|
|
}
|
|
|
|
}
|
|
|
|
// conversion routines between float and half precision
|
|
static inline std::uint32_t f32_as_u32(float f) { union { float f; std::uint32_t u; } v; v.f = f; return v.u; }
|
|
static inline float u32_as_f32(std::uint32_t u) { union { float f; std::uint32_t u; } v; v.u = u; return v.f; }
|
|
static inline int clamp_int(int i, int l, int h) { return std::min(std::max(i, l), h); }
|
|
|
|
// half float, the f16 is in the low 16 bits of the input argument
|
|
static inline float __convert_half_to_float(std::uint32_t a) noexcept {
|
|
std::uint32_t u = ((a << 13) + 0x70000000U) & 0x8fffe000U;
|
|
std::uint32_t v = f32_as_u32(u32_as_f32(u) * 0x1.0p+112f) + 0x38000000U;
|
|
u = (a & 0x7fff) != 0 ? v : u;
|
|
return u32_as_f32(u) * 0x1.0p-112f;
|
|
}
|
|
|
|
// float half with nearest even rounding
|
|
// The lower 16 bits of the result is the bit pattern for the f16
|
|
static inline std::uint32_t __convert_float_to_half(float a) noexcept {
|
|
std::uint32_t u = f32_as_u32(a);
|
|
int e = static_cast<int>((u >> 23) & 0xff) - 127 + 15;
|
|
std::uint32_t m = ((u >> 11) & 0xffe) | ((u & 0xfff) != 0);
|
|
std::uint32_t i = 0x7c00 | (m != 0 ? 0x0200 : 0);
|
|
std::uint32_t n = ((std::uint32_t)e << 12) | m;
|
|
std::uint32_t s = (u >> 16) & 0x8000;
|
|
int b = clamp_int(1-e, 0, 13);
|
|
std::uint32_t d = (0x1000 | m) >> b;
|
|
d |= (d << b) != (0x1000 | m);
|
|
std::uint32_t v = e < 1 ? d : n;
|
|
v = (v >> 2) + (((v & 0x7) == 3) | ((v & 0x7) > 5));
|
|
v = e > 30 ? 0x7c00 : v;
|
|
v = e == 143 ? i : v;
|
|
return s | v;
|
|
}
|
|
|
|
extern "C" float __gnu_h2f_ieee(unsigned short h){
|
|
return __convert_half_to_float((std::uint32_t) h);
|
|
}
|
|
|
|
extern "C" unsigned short __gnu_f2h_ieee(float f){
|
|
return (unsigned short)__convert_float_to_half(f);
|
|
}
|
|
|
|
#endif // defined(ATI_OS_LINUX)
|