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

Revert "Revert "Merge branch 'amd-master-next' into amd-npi-next""

Browse Source 
This reverts commit 28b17d3dbd.

Reason for revert: <INSERT REASONING HERE>

Change-Id: I92ceb171e31026ed1864704cef2fc1497b883ef9


[ROCm/hip commit: ad2d55c144]
This commit is contained in:
Vladislav Sytchenko
2020-10-05 13:20:58 -04:00
parent 28b17d3dbd
commit e4caaa2a77
111 changed files with 6800 additions and 753 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/rocclr/CMakeLists.txt
+25 -5
View File
@@ -96,6 +96,14 @@ find_package(amd_comgr REQUIRED CONFIG
message(STATUS "Code Object Manager found at ${amd_comgr_DIR}.")
find_package(LLVM REQUIRED CONFIG
PATHS
/opt/rocm/llvm
PATH_SUFFIXES
lib/cmake/llvm)
message(STATUS "llvm found at ${LLVM_DIR}.")
add_library(hip64 OBJECT
hip_context.cpp
hip_code_object.cpp
@@ -148,10 +156,9 @@ endif()
# Short-Term solution for pre-compiled headers for online compilation
# Enable pre compiled header
if(${ENABLE_HIP_PCH})
execute_process(COMMAND sh -c "${CMAKE_CURRENT_SOURCE_DIR}/../bin/hip_gen_pch.sh")
execute_process(COMMAND sh -c "${CMAKE_CURRENT_SOURCE_DIR}/../bin/hip_embed_pch.sh")
add_definitions(-DENABLE_HIP_PCH)
if(${__HIP_ENABLE_PCH})
execute_process(COMMAND sh -c "${CMAKE_CURRENT_SOURCE_DIR}/../bin/hip_embed_pch.sh ${LLVM_DIR}")
add_definitions(-D__HIP_ENABLE_PCH)
endif()
# Enable profiling API
@@ -216,7 +223,7 @@ add_library(device INTERFACE)
target_link_libraries(device INTERFACE host)
# Short-Term solution for pre-compiled headers for online compilation
if(${ENABLE_HIP_PCH})
if(${__HIP_ENABLE_PCH})
target_link_libraries(amdhip64 PRIVATE ${CMAKE_BINARY_DIR}/hip_pch.o)
endif()
@@ -227,6 +234,18 @@ endif()
# filename.
if(${BUILD_SHARED_LIBS})
target_link_libraries(amdhip64 PRIVATE amdrocclr_static Threads::Threads dl hsa-runtime64::hsa-runtime64)
add_custom_command(TARGET amdhip64 POST_BUILD COMMAND
${CMAKE_COMMAND} -E create_symlink ${PROJECT_BINARY_DIR}/lib/libamdhip64.so.${HIP_LIB_VERSION_STRING}
${PROJECT_BINARY_DIR}/lib/libhip_hcc.so.${HIP_LIB_VERSION_MAJOR})
add_custom_command(TARGET amdhip64 POST_BUILD COMMAND
${CMAKE_COMMAND} -E create_symlink ${PROJECT_BINARY_DIR}/lib/libhip_hcc.so.${HIP_LIB_VERSION_MAJOR}
${PROJECT_BINARY_DIR}/lib/libhip_hcc.so)
add_custom_command(TARGET amdhip64 POST_BUILD COMMAND
${CMAKE_COMMAND} -E copy ${PROJECT_BINARY_DIR}/.hipInfo ${PROJECT_BINARY_DIR}/lib/.hipInfo)
add_custom_command(TARGET amdhip64 POST_BUILD COMMAND ${CMAKE_COMMAND} -E copy_directory
${PROJECT_SOURCE_DIR}/include ${PROJECT_BINARY_DIR}/include)
INSTALL(PROGRAMS $<TARGET_FILE:amdhip64> DESTINATION lib COMPONENT MAIN)
else()
target_link_libraries(amdhip64 PRIVATE Threads::Threads dl hsa-runtime64::hsa-runtime64 amd_comgr)
@@ -244,6 +263,7 @@ else()
INSTALL(PROGRAMS $<TARGET_FILE:amdhip64> DESTINATION lib COMPONENT MAIN)
endif()
INSTALL(TARGETS amdhip64 host device EXPORT hip-targets DESTINATION ${LIB_INSTALL_DIR})
INSTALL(EXPORT hip-targets DESTINATION ${CONFIG_PACKAGE_INSTALL_DIR} NAMESPACE hip::)
projects/hip/rocclr/hip_code_object.cpp
-25
View File
@@ -202,19 +202,10 @@ hipError_t DynCO::populateDynGlobalVars() {
return hipErrorSharedObjectSymbolNotFound;
}
if (!dev_program->getUndefinedVarFromCodeObj(&undef_var_names)) {
DevLogPrintfError("Could not get undefined Variables for Module: 0x%x \n", module());
return hipErrorSharedObjectSymbolNotFound;
}
for (auto& elem : var_names) {
vars_.insert(std::make_pair(elem, new Var(elem, Var::DeviceVarKind::DVK_Variable, 0, 0, 0, nullptr)));
}
for (auto& elem : undef_var_names) {
vars_.insert(std::make_pair(elem, new Var(elem, Var::DeviceVarKind::DVK_Texture, 0, 0, 0, nullptr)));
}
return hipSuccess;
}
@@ -377,20 +368,4 @@ hipError_t StatCO::getStatGlobalVar(const void* hostVar, int deviceId, hipDevice
*size_ptr = dvar->size();
return hipSuccess;
}
hipError_t StatCO::getStatGlobalVarByName(std::string hostVar, int deviceId, hipModule_t hmod,
hipDeviceptr_t* dev_ptr, size_t* size_ptr) {
amd::ScopedLock lock(sclock_);
for (auto& elem : vars_) {
if ((elem.second->name() == hostVar)
&& (elem.second->module(deviceId) == hmod)) {
*dev_ptr = elem.second->device_ptr(deviceId);
*size_ptr = elem.second->device_size(deviceId);
return hipSuccess;
}
}
return hipErrorNotFound;
}
}; //namespace: hip
projects/hip/rocclr/hip_code_object.hpp
-2
View File
@@ -118,8 +118,6 @@ public:
hipError_t getStatFuncAttr(hipFuncAttributes* func_attr, const void* hostFunction, int deviceId);
hipError_t getStatGlobalVar(const void* hostVar, int deviceId, hipDeviceptr_t* dev_ptr,
size_t* size_ptr);
hipError_t getStatGlobalVarByName(std::string hostVar, int deviceId, hipModule_t hmod,
hipDeviceptr_t* dev_ptr, size_t* size_ptr);
private:
friend class ::PlatformState;
projects/hip/rocclr/hip_device.cpp
+1 -1
View File
@@ -155,7 +155,7 @@ hipError_t hipGetDeviceProperties ( hipDeviceProp_t* props, hipDevice_t device )
::strncpy(deviceProps.name, info.boardName_, 128);
deviceProps.totalGlobalMem = info.globalMemSize_;
deviceProps.sharedMemPerBlock = info.localMemSizePerCU_;
deviceProps.regsPerBlock = info.availableSGPRs_;
deviceProps.regsPerBlock = info.availableRegistersPerCU_;
deviceProps.warpSize = info.wavefrontWidth_;
deviceProps.maxThreadsPerBlock = info.maxWorkGroupSize_;
deviceProps.maxThreadsDim[0] = info.maxWorkItemSizes_[0];
projects/hip/rocclr/hip_fatbin.cpp
+5 -2
View File
@@ -12,7 +12,7 @@ FatBinaryDeviceInfo::~FatBinaryDeviceInfo() {
}
FatBinaryInfo::FatBinaryInfo(const char* fname, const void* image)
: fdesc_(-1), fsize_(0), image_(image), uri_(std::string()) {
: fdesc_(amd::Os::FDescInit()), fsize_(0), image_(image), uri_(std::string()) {
guarantee(fname || image);
if (fname != nullptr) {
@@ -41,7 +41,7 @@ FatBinaryInfo::~FatBinaryInfo() {
}
fname_ = std::string();
fdesc_ = -1;
fdesc_ = amd::Os::FDescInit();
fsize_ = 0;
image_ = nullptr;
uri_ = std::string();
@@ -64,6 +64,9 @@ hipError_t FatBinaryInfo::ExtractFatBinary(const std::vector<hip::Device*>& devi
if (!amd::Os::GetFileHandle(fname_.c_str(), &fdesc_, &fsize_)) {
return hipErrorFileNotFound;
}
if (fsize_ == 0) {
return hipErrorInvalidKernelFile;
}
// Extract the code object from file
hip_error = CodeObject::ExtractCodeObjectFromFile(fdesc_, fsize_,
projects/hip/rocclr/hip_global.cpp
+3 -1
View File
@@ -5,7 +5,9 @@
#include "hip_code_object.hpp"
#include "platform/program.hpp"
#ifdef ENABLE_HIP_PCH
#ifdef __HIP_ENABLE_PCH
extern const char __hip_pch[];
extern unsigned __hip_pch_size;
void __hipGetPCH(const char** pch, unsigned int *size) {
*pch = __hip_pch;
*size = __hip_pch_size;
projects/hip/rocclr/hip_global.hpp
-5
View File
@@ -95,11 +95,6 @@ public:
hipError_t getStatDeviceVar(DeviceVar** dvar, int deviceId);
void resize_dVar(size_t size) { dVar_.resize(size); }
//Accessor for device_ptrs.
std::string name() const { return name_; }
hipModule_t module(int deviceId) const { return nullptr; }
hipDeviceptr_t device_ptr(int deviceId) const { return dVar_[deviceId]->device_ptr(); }
size_t device_size(int deviceId) const { return dVar_[deviceId]->size(); }
FatBinaryInfo** moduleInfo() { return modules_; };
private:
projects/hip/rocclr/hip_internal.hpp
-2
View File
@@ -252,8 +252,6 @@ extern int ihipGetDevice();
extern hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags);
extern amd::Memory* getMemoryObject(const void* ptr, size_t& offset);
extern amd::Memory* getMemoryObjectWithOffset(const void* ptr, const size_t size);
extern bool CL_CALLBACK getSvarInfo(cl_program program, std::string var_name, void** var_addr,
size_t* var_size);
constexpr bool kOptionChangeable = true;
constexpr bool kNewDevProg = false;
projects/hip/rocclr/hip_memory.cpp
+37 -36
View File
@@ -124,7 +124,7 @@ hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
if (*ptr == nullptr) {
size_t free = 0, total =0;
hipMemGetInfo(&free, &total);
LogPrintfError("Allocation failed : Device memory : required :%u | free :%u | total :%u \n", sizeBytes, free, total);
LogPrintfError("Allocation failed : Device memory : required :%zu | free :%zu | total :%zu \n", sizeBytes, free, total);
return hipErrorOutOfMemory;
}
@@ -202,14 +202,14 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
}
} else {
amd::HostQueue* pQueue = &queue;
if (queueDevice != srcMemory->getContext().devices()[0]) {
if ((srcMemory->getContext().devices()[0] == dstMemory->getContext().devices()[0]) &&
(queueDevice != srcMemory->getContext().devices()[0])) {
pQueue = hip::getNullStream(srcMemory->getContext());
amd::Command* cmd = queue.getLastQueuedCommand(true);
if (cmd != nullptr) {
waitList.push_back(cmd);
}
}
command = new amd::CopyMemoryCommand(*pQueue, CL_COMMAND_COPY_BUFFER, waitList,
*srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, dOffset, sizeBytes);
}
@@ -1850,18 +1850,27 @@ hipError_t ihipMemset3D(hipPitchedPtr pitchedDevPtr,
hipExtent extent,
hipStream_t stream,
bool isAsync = false) {
if (pitchedDevPtr.pitch == extent.width) {
return ihipMemset(pitchedDevPtr.ptr, value, sizeof(int8_t), extent.width * extent.height * extent.depth, stream, isAsync);
}
// Workaround for cases when pitch > row untill fill kernel will be updated to support pitch.
// Fallback to filling one row at a time.
amd::HostQueue* queue = hip::getQueue(stream);
size_t offset = 0;
amd::Memory* memory = getMemoryObject(pitchedDevPtr.ptr, offset);
auto sizeBytes = extent.width * extent.height * extent.depth;
if (memory == nullptr) {
return hipErrorInvalidValue;
}
if (sizeBytes > memory->getSize()) {
return hipErrorInvalidValue;
}
if (pitchedDevPtr.pitch == extent.width) {
return ihipMemset(pitchedDevPtr.ptr, value, sizeof(int8_t), static_cast<size_t>(sizeBytes), stream, isAsync);
}
// Workaround for cases when pitch > row until fill kernel will be updated to support pitch.
// Fall back to filling one row at a time.
amd::HostQueue* queue = hip::getQueue(stream);
amd::Coord3D origin(offset);
amd::Coord3D region(pitchedDevPtr.xsize, pitchedDevPtr.ysize, extent.depth);
amd::BufferRect rect;
@@ -1870,34 +1879,26 @@ hipError_t ihipMemset3D(hipPitchedPtr pitchedDevPtr,
return hipErrorInvalidValue;
}
if (memory != nullptr) {
std::vector<amd::FillMemoryCommand*> commands;
std::vector<amd::FillMemoryCommand*> commands;
for (size_t slice = 0; slice < extent.depth; slice++) {
for (size_t row = 0; row < extent.height; row++) {
const size_t rowOffset = rect.offset(0, row, slice);
amd::FillMemoryCommand* command = new amd::FillMemoryCommand(*queue,
CL_COMMAND_FILL_BUFFER,
amd::Command::EventWaitList{},
*memory->asBuffer(),
&value,
sizeof(int8_t),
amd::Coord3D{rowOffset, 0, 0},
amd::Coord3D{extent.width, 1, 1});
for (size_t slice = 0; slice < extent.depth; slice++) {
for (size_t row = 0; row < extent.height; row++) {
const size_t rowOffset = rect.offset(0, row, slice);
amd::FillMemoryCommand *command = new amd::FillMemoryCommand(*queue,
CL_COMMAND_FILL_BUFFER, amd::Command::EventWaitList { },
*memory->asBuffer(), &value, sizeof(int8_t), amd::Coord3D { rowOffset,
0, 0 }, amd::Coord3D { extent.width, 1, 1 });
command->enqueue();
commands.push_back(command);
}
command->enqueue();
commands.push_back(command);
}
}
for (auto &command: commands) {
if (!isAsync) {
command->awaitCompletion();
}
command->release();
for (auto &command : commands) {
if (!isAsync) {
command->awaitCompletion();
}
} else {
return hipErrorInvalidValue;
command->release();
}
return hipSuccess;
@@ -2038,7 +2039,7 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void
memset(attributes, 0, sizeof(hipPointerAttribute_t));
if (memObj != nullptr) {
attributes->memoryType = (CL_MEM_SVM_FINE_GRAIN_BUFFER & memObj->getMemFlags())? hipMemoryTypeHost : hipMemoryTypeDevice;
attributes->memoryType = ((CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_USE_HOST_PTR) & memObj->getMemFlags())? hipMemoryTypeHost : hipMemoryTypeDevice;
if (attributes->memoryType == hipMemoryTypeHost) {
attributes->hostPointer = static_cast<char*>(memObj->getSvmPtr()) + offset;
}
projects/hip/rocclr/hip_module.cpp
+1 -1
View File
@@ -537,7 +537,7 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL
if (result != hipSuccess) {
break;
}
prevGridSize += launch.gridDim.x * launch.gridDim.y * launch.gridDim.z;
prevGridSize += globalWorkSizeX * globalWorkSizeY * globalWorkSizeZ;
}
// Sync the execution streams on all devices
projects/hip/rocclr/hip_peer.cpp
+4
View File
@@ -97,6 +97,10 @@ hipError_t hipExtGetLinkTypeAndHopCount(int device1, int device2,
uint32_t* linktype, uint32_t* hopcount) {
HIP_INIT_API(hipExtGetLinkTypeAndHopCount, device1, device2, linktype, hopcount);
if (linktype == nullptr || hopcount == nullptr ||
device1 == device2 || device1 < 0 || device2 < 0) {
HIP_RETURN(hipErrorInvalidValue);
}
// Fill out the list of LinkAttributes
std::vector<amd::Device::LinkAttrType> link_attrs;
link_attrs.push_back(std::make_pair(amd::Device::LinkAttribute::kLinkLinkType, 0));
projects/hip/rocclr/hip_platform.cpp
+13 -28
View File
@@ -80,27 +80,6 @@ extern "C" hip::FatBinaryInfo** __hipRegisterFatBinary(const void* data)
return PlatformState::instance().addFatBinary(fbwrapper->binary);
}
bool PlatformState::getShadowVarInfo(std::string var_name, hipModule_t hmod,
void** var_addr, size_t* var_size) {
amd::ScopedLock lock(lock_);
if (hipSuccess == getDynGlobalVar(var_name.c_str(), ihipGetDevice(), hmod, var_addr, var_size)) {
return true;
}
if (hipSuccess == getStatGlobalVarByName(var_name, ihipGetDevice(), hmod, var_addr, var_size)) {
return true;
}
return false;
}
bool CL_CALLBACK getSvarInfo(cl_program program, std::string var_name, void** var_addr,
size_t* var_size) {
return PlatformState::instance().getShadowVarInfo(var_name, reinterpret_cast<hipModule_t>(program),
var_addr, var_size);
}
extern "C" void __hipRegisterFunction(
hip::FatBinaryInfo** modules,
const void* hostFunction,
@@ -686,11 +665,19 @@ static inline std::uint32_t __convert_float_to_half(float a) noexcept {
return s | v;
}
extern "C" __attribute__((weak)) float __gnu_h2f_ieee(unsigned short h){
extern "C"
#if !defined(_MSC_VER)
__attribute__((weak))
#endif
float __gnu_h2f_ieee(unsigned short h){
return __convert_half_to_float((std::uint32_t) h);
}
extern "C" __attribute__((weak)) unsigned short __gnu_f2h_ieee(float f){
extern "C"
#if !defined(_MSC_VER)
__attribute__((weak))
#endif
unsigned short __gnu_f2h_ieee(float f){
return (unsigned short)__convert_float_to_half(f);
}
@@ -765,6 +752,9 @@ hipError_t PlatformState::getDynFunc(hipFunction_t* hfunc, hipModule_t hmod,
DevLogPrintfError("Cannot find the module: 0x%x", hmod);
return hipErrorNotFound;
}
if (0 == strlen(func_name)) {
return hipErrorNotFound;
}
return it->second->getDynFunc(hfunc, func_name);
}
@@ -868,11 +858,6 @@ hipError_t PlatformState::getStatGlobalVar(const void* hostVar, int deviceId, hi
return statCO_.getStatGlobalVar(hostVar, deviceId, dev_ptr, size_ptr);
}
hipError_t PlatformState::getStatGlobalVarByName(std::string hostVar, int deviceId, hipModule_t hmod,
hipDeviceptr_t* dev_ptr, size_t* size_ptr) {
return statCO_.getStatGlobalVarByName(hostVar, deviceId, hmod, dev_ptr, size_ptr);
}
void PlatformState::setupArgument(const void *arg, size_t size, size_t offset) {
auto& arguments = execStack_.top().arguments_;
projects/hip/rocclr/hip_platform.hpp
-5
View File
@@ -77,11 +77,6 @@ public:
hipError_t getStatFuncAttr(hipFuncAttributes* func_attr, const void* hostFunction, int deviceId);
hipError_t getStatGlobalVar(const void* hostVar, int deviceId, hipDeviceptr_t* dev_ptr,
size_t* size_ptr);
hipError_t getStatGlobalVarByName(std::string hostVar, int deviceId, hipModule_t hmod,
hipDeviceptr_t* dev_ptr, size_t* size_ptr);
bool getShadowVarInfo(std::string var_name, hipModule_t hmod,
void** var_addr, size_t* var_size);
//Exec Functions
void setupArgument(const void *arg, size_t size, size_t offset);