From 73f931bdbd067b32ee75709d1532de143d93987c Mon Sep 17 00:00:00 2001 From: Laurent Morichetti Date: Thu, 23 May 2019 18:03:32 -0700 Subject: [PATCH 01/19] Add support for code object v3 Use the code object manager library to parse the code object metadata. Both code object v2 and v3 formats are now supported for HCC generated binaries. --- CMakeLists.txt | 15 +++ src/hip_module.cpp | 82 +++++++++++---- src/program_state.inl | 236 ++++++++++++++++++++++++++++-------------- 3 files changed, 237 insertions(+), 96 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index ba855ad86b..8701fe5635 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -271,6 +271,21 @@ if(HIP_PLATFORM STREQUAL "hcc") target_link_libraries(hiprtc PUBLIC stdc++fs) endif() + + if(HIP_PLATFORM STREQUAL "hcc") + find_package(amd_comgr REQUIRED CONFIG + PATHS + /opt/rocm/ + PATH_SUFFIXES + cmake/amd_comgr + lib/cmake/amd_comgr + ) + MESSAGE(STATUS "Code Object Manager found at ${amd_comgr_DIR}.") + endif() + + target_link_libraries(hip_hcc PRIVATE amd_comgr) + target_link_libraries(hip_hcc_static PRIVATE amd_comgr) + string(REPLACE " " ";" HCC_CXX_FLAGS_LIST ${HCC_CXX_FLAGS}) foreach(TARGET hip_hcc hip_hcc_static) target_include_directories(${TARGET} SYSTEM INTERFACE $/include>;${HSA_PATH}/include) diff --git a/src/hip_module.cpp b/src/hip_module.cpp index e8a8801e98..ed8db42c86 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -55,6 +55,18 @@ THE SOFTWARE. using namespace ELFIO; using namespace std; +struct amd_kernel_code_v3_t { + uint32_t group_segment_fixed_size; + uint32_t private_segment_fixed_size; + uint8_t reserved0[8]; + int64_t kernel_code_entry_byte_offset; + uint8_t reserved1[24]; + uint32_t compute_pgm_rsrc1; + uint32_t compute_pgm_rsrc2; + uint16_t kernel_code_properties; + uint8_t reserved2[6]; +}; + // calculate MD5 checksum inline std::string checksum(size_t size, const char *source) { // FNV-1a hashing, 64-bit version @@ -206,10 +218,20 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, aql.grid_size_x = globalWorkSizeX; aql.grid_size_y = globalWorkSizeY; aql.grid_size_z = globalWorkSizeZ; - aql.group_segment_size = - f->_header->workgroup_group_segment_byte_size + sharedMemBytes; - aql.private_segment_size = - f->_header->workitem_private_segment_byte_size; + bool is_code_object_v3 = f->_name.find(".kd") != std::string::npos; + if (is_code_object_v3) { + const auto* header = + reinterpret_cast(f->_header); + aql.group_segment_size = + header->group_segment_fixed_size + sharedMemBytes; + aql.private_segment_size = + header->private_segment_fixed_size; + } else { + aql.group_segment_size = + f->_header->workgroup_group_segment_byte_size + sharedMemBytes; + aql.private_segment_size = + f->_header->workitem_private_segment_byte_size; + } aql.kernel_object = f->_object; aql.setup = 3 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS; aql.header = @@ -462,6 +484,12 @@ hipError_t ihipModuleGetFunction(hipFunction_t* func, hipModule_t hmod, const ch auto kernel = find_kernel_by_name(hmod->executable, name, agent); + if (kernel.handle == 0u) { + std::string name_str(name); + name_str.append(".kd"); + kernel = find_kernel_by_name(hmod->executable, name_str.c_str(), agent); + } + if (kernel.handle == 0u) return hipErrorNotFound; // TODO: refactor the whole ihipThisThat, which is a mess and yields the @@ -486,7 +514,11 @@ hipError_t hipModuleGetFunctionEx(hipFunction_t* hfunc, hipModule_t hmod, } namespace { -hipFuncAttributes make_function_attributes(const amd_kernel_code_t& header) { +const amd_kernel_code_v3_t *header_v3(const ihipModuleSymbol_t& kd) { + return reinterpret_cast(kd._header); +} + +hipFuncAttributes make_function_attributes(const ihipModuleSymbol_t& kd) { hipFuncAttributes r{}; hipDeviceProp_t prop{}; @@ -495,16 +527,31 @@ hipFuncAttributes make_function_attributes(const amd_kernel_code_t& header) { // available per CU, therefore we hardcode it to 64 KiRegisters. prop.regsPerBlock = prop.regsPerBlock ? prop.regsPerBlock : 64 * 1024; - r.localSizeBytes = header.workitem_private_segment_byte_size; - r.sharedSizeBytes = header.workgroup_group_segment_byte_size; + bool is_code_object_v3 = kd._name.find(".kd") != std::string::npos; + if (is_code_object_v3) { + r.localSizeBytes = header_v3(kd)->private_segment_fixed_size; + r.sharedSizeBytes = header_v3(kd)->group_segment_fixed_size; + } else { + r.localSizeBytes = kd._header->workitem_private_segment_byte_size; + r.sharedSizeBytes = kd._header->workgroup_group_segment_byte_size; + } r.maxDynamicSharedSizeBytes = prop.sharedMemPerBlock - r.sharedSizeBytes; - r.numRegs = header.workitem_vgpr_count; + if (is_code_object_v3) { + r.numRegs = ((header_v3(kd)->compute_pgm_rsrc1 & 0x3F) + 1) << 2; + } else { + r.numRegs = kd._header->workitem_vgpr_count; + } r.maxThreadsPerBlock = r.numRegs ? std::min(prop.maxThreadsPerBlock, prop.regsPerBlock / r.numRegs) : prop.maxThreadsPerBlock; - r.binaryVersion = - header.amd_machine_version_major * 10 + - header.amd_machine_version_minor; + if (is_code_object_v3) { + r.binaryVersion = 0; // FIXME: should it be the ISA version or code + // object format version? + } else { + r.binaryVersion = + kd._header->amd_machine_version_major * 10 + + kd._header->amd_machine_version_minor; + } r.ptxVersion = prop.major * 10 + prop.minor; // HIP currently presents itself as PTX 3.0. return r; @@ -520,11 +567,10 @@ hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func) auto agent = this_agent(); auto kd = get_program_state().kernel_descriptor(reinterpret_cast(func), agent); - const auto header = kd->_header; - if (!header) throw runtime_error{"Ill-formed Kernel_descriptor."}; + if (!kd->_header) throw runtime_error{"Ill-formed Kernel_descriptor."}; - *attr = make_function_attributes(*header); + *attr = make_function_attributes(*kd); return hipSuccess; } @@ -555,11 +601,9 @@ hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) { (*module)->executable = get_program_state().load_executable( content.data(), content.size(), (*module)->executable, this_agent()); - istringstream elf{content}; - ELFIO::elfio reader; - if (reader.load(elf)) { - program_state_impl::read_kernarg_metadata(reader, (*module)->kernargs); - } + + std::vector blob(content.cbegin(), content.cend()); + program_state_impl::read_kernarg_metadata(blob, (*module)->kernargs); // compute the hash of the code object (*module)->hash = checksum(content.length(), content.data()); diff --git a/src/program_state.inl b/src/program_state.inl index f1397b3fe9..639eac9228 100644 --- a/src/program_state.inl +++ b/src/program_state.inl @@ -17,6 +17,7 @@ #include #include #include +#include #include @@ -540,9 +541,13 @@ public: std::call_once(functions[agent].first, [this](hsa_agent_t aa) { for (auto&& function : get_function_names()) { - const auto it = get_kernels(aa).find(function.second); + auto it = get_kernels(aa).find(function.second); - if (it == get_kernels(aa).cend()) continue; + if (it == get_kernels(aa).cend()) { + it = get_kernels(aa).find(function.second + ".kd"); + if (it == get_kernels(aa).cend()) + continue; + } for (auto&& kernel_symbol : it->second) { functions[aa].second.emplace( @@ -556,92 +561,172 @@ public: } static - std::size_t parse_args( - const std::string& metadata, - std::size_t f, - std::size_t l, + std::string metadata_to_string(const amd_comgr_metadata_node_t& md) { + std::string str; + size_t size; + + if (amd_comgr_get_metadata_string(md, &size, NULL) + == AMD_COMGR_STATUS_SUCCESS) { + str.resize(size - 1); + amd_comgr_get_metadata_string(md, &size, &str[0]); + } + return str; + } + + static + void parse_args( + const amd_comgr_metadata_node_t& args_md, + bool is_code_object_v3, std::vector>& size_align) { - if (f == l) return f; - if (!size_align.empty()) return l; + size_t arg_count = 0; + if (amd_comgr_get_metadata_list_size(args_md, &arg_count) + != AMD_COMGR_STATUS_SUCCESS) + return; - do { - static constexpr size_t size_sz{5}; - f = metadata.find("Size:", f) + size_sz; + for (size_t i = 0; i < arg_count; ++i) { + amd_comgr_metadata_node_t arg_md; - if (l <= f) return f; + if (amd_comgr_index_list_metadata(args_md, i, &arg_md) + != AMD_COMGR_STATUS_SUCCESS) + return; - auto size = std::strtoul(&metadata[f], nullptr, 10); + amd_comgr_metadata_node_t arg_size_md; + if (amd_comgr_metadata_lookup(arg_md, + is_code_object_v3 ? ".size" : "Size", + &arg_size_md) + != AMD_COMGR_STATUS_SUCCESS) + return; - static constexpr size_t align_sz{6}; - f = metadata.find("Align:", f) + align_sz; + size_t arg_size = std::stoul(metadata_to_string(arg_size_md)); - char* l{}; - auto align = std::strtoul(&metadata[f], &l, 10); + if (amd_comgr_destroy_metadata(arg_size_md) + != AMD_COMGR_STATUS_SUCCESS) + return; - f += (l - &metadata[f]) + 1; + size_t arg_align; - size_align.emplace_back(size, align); - } while (true); + if (is_code_object_v3) { + amd_comgr_metadata_node_t arg_offset_md; + if (amd_comgr_metadata_lookup(arg_md, ".offset", &arg_offset_md) + != AMD_COMGR_STATUS_SUCCESS) + return; + + size_t arg_offset + = std::stoul(metadata_to_string(arg_offset_md)); + + if (amd_comgr_destroy_metadata(arg_offset_md) + != AMD_COMGR_STATUS_SUCCESS) + return; + + arg_align = 1; + while (arg_offset && (arg_offset & 1) == 0) { + arg_offset >>= 1; + arg_align <<= 1; + } + } else { + amd_comgr_metadata_node_t arg_align_md; + if (amd_comgr_metadata_lookup(arg_md, "Align", &arg_align_md) + != AMD_COMGR_STATUS_SUCCESS) + return; + + arg_align = std::stoul(metadata_to_string(arg_align_md)); + + if (amd_comgr_destroy_metadata(arg_align_md) + != AMD_COMGR_STATUS_SUCCESS) + return; + } + + size_align.emplace_back(arg_size, arg_align); + + if (amd_comgr_destroy_metadata(arg_md) + != AMD_COMGR_STATUS_SUCCESS) + return; + } } static void read_kernarg_metadata( - ELFIO::elfio& reader, + const std::vector& blob, std::unordered_map< std::string, std::vector>>& kernargs) { - // TODO: this is inefficient. - auto it = find_section_if(reader, [](const ELFIO::section* x) { - return x->get_type() == SHT_NOTE; - }); + amd_comgr_data_t dataIn; + amd_comgr_status_t status; - if (!it) return; + if (amd_comgr_create_data(AMD_COMGR_DATA_KIND_RELOCATABLE, &dataIn) + != AMD_COMGR_STATUS_SUCCESS) + return; - const ELFIO::note_section_accessor acc{reader, it}; - for (decltype(acc.get_notes_num()) i = 0; i != acc.get_notes_num(); ++i) { - ELFIO::Elf_Word type{}; - std::string name{}; - void* desc{}; - ELFIO::Elf_Word desc_size{}; + if (amd_comgr_set_data(dataIn, blob.size(), blob.data()) + != AMD_COMGR_STATUS_SUCCESS) + return; - acc.get_note(i, type, name, desc, desc_size); + amd_comgr_metadata_node_t metadata; + if (amd_comgr_get_data_metadata(dataIn, &metadata) + != AMD_COMGR_STATUS_SUCCESS) + return; - if (name != "AMD") continue; // TODO: switch to using NT_AMD_AMDGPU_HSA_METADATA. - - std::string tmp{ - static_cast(desc), static_cast(desc) + desc_size}; - - auto dx = tmp.find("Kernels:"); - - if (dx == std::string::npos) continue; - - static constexpr decltype(tmp.size()) kernels_sz{8}; - dx += kernels_sz; - - do { - dx = tmp.find("Name:", dx); - - if (dx == std::string::npos) break; - - static constexpr decltype(tmp.size()) name_sz{5}; - dx = tmp.find_first_not_of(" '", dx + name_sz); - - auto fn = tmp.substr(dx, tmp.find_first_of("'\n", dx) - dx); - dx += fn.size(); - - auto dx1 = tmp.find("CodeProps", dx); - dx = tmp.find("Args:", dx); - - if (dx1 < dx) { - dx = dx1; - continue; - } - if (dx == std::string::npos) break; - - static constexpr decltype(tmp.size()) args_sz{5}; - dx = parse_args(tmp, dx + args_sz, dx1, kernargs[fn]); - } while (true); + bool is_code_object_v3 = false; + amd_comgr_metadata_node_t kernels_md; + if (amd_comgr_metadata_lookup(metadata, "Kernels", &kernels_md) + != AMD_COMGR_STATUS_SUCCESS) { + if (amd_comgr_metadata_lookup(metadata, + "amdhsa.kernels", + &kernels_md) + != AMD_COMGR_STATUS_SUCCESS) + return; + is_code_object_v3 = true; } + + size_t kernel_count = 0; + if (amd_comgr_get_metadata_list_size(kernels_md, &kernel_count) + != AMD_COMGR_STATUS_SUCCESS) + return; + + for (size_t i = 0; i < kernel_count; i++) { + amd_comgr_metadata_node_t kernel_md; + + if (amd_comgr_index_list_metadata(kernels_md, i, &kernel_md) + != AMD_COMGR_STATUS_SUCCESS) + continue; + + amd_comgr_metadata_node_t name_md; + if (amd_comgr_metadata_lookup(kernel_md, + is_code_object_v3 ? ".name" : "Name", + &name_md) + != AMD_COMGR_STATUS_SUCCESS) + continue; + + std::string kernel_name_str = metadata_to_string(name_md); + + if (amd_comgr_destroy_metadata(name_md) + != AMD_COMGR_STATUS_SUCCESS) + continue; + + if (is_code_object_v3) + kernel_name_str.append(".kd"); + + + amd_comgr_metadata_node_t args_md; + if (amd_comgr_metadata_lookup(kernel_md, + is_code_object_v3 ? ".args" : "Args", + &args_md) + != AMD_COMGR_STATUS_SUCCESS) + continue; + + parse_args(args_md, is_code_object_v3, kernargs[kernel_name_str]); + + if (amd_comgr_destroy_metadata(args_md) != AMD_COMGR_STATUS_SUCCESS + || amd_comgr_destroy_metadata(kernel_md) + != AMD_COMGR_STATUS_SUCCESS) + continue; + } + + if (amd_comgr_destroy_metadata(kernels_md) != AMD_COMGR_STATUS_SUCCESS + || amd_comgr_destroy_metadata(metadata) != AMD_COMGR_STATUS_SUCCESS) + return; + + amd_comgr_release_data(dataIn); } const std::unordered_mapsecond); if (it1 == get_kernargs().end()) { - hip_throw(std::runtime_error{ - "Missing metadata for __global__ function: " + it->second}); + it1 = get_kernargs().find(it->second + ".kd"); + if (it1 == get_kernargs().end()) { + hip_throw(std::runtime_error{ + "Missing metadata for __global__ function: " + it->second}); + } } return it1->second; From 4239cfcf02d31c88a0973d32cdef75e296359ef8 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Sat, 4 May 2019 12:57:36 -0400 Subject: [PATCH 02/19] moving agent_globals_impl into hip_module --- include/hip/hcc_detail/hip_runtime_api.h | 147 ++++---------------- src/hip_module.cpp | 165 +++++++++++++++++++++-- 2 files changed, 177 insertions(+), 135 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index d870963101..17a42e2a16 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -2590,6 +2590,7 @@ hipError_t hipModuleGetFunction(hipFunction_t* function, hipModule_t module, con hipError_t hipFuncGetAttributes(struct hipFuncAttributes* attr, const void* func); +#if 0 struct Agent_global { Agent_global() : name(nullptr), address(nullptr), byte_cnt(0) {} @@ -2629,6 +2630,7 @@ struct Agent_global { hipDeviceptr_t address; uint32_t byte_cnt; }; +#endif #if !__HIP_VDI__ #if defined(__cplusplus) @@ -2636,134 +2638,35 @@ struct Agent_global { #endif namespace hip_impl { -hsa_executable_t executable_for(hipModule_t); -const char* hash_for(hipModule_t); + class agent_globals_impl; + class agent_globals { + public: + agent_globals(); + ~agent_globals(); + agent_globals(const agent_globals&) = delete; -template -std::pair read_global_description( - ForwardIterator f, ForwardIterator l, const char* name) { - const auto it = std::find_if(f, l, [=](const Agent_global& x) { - return strcmp(x.name, name) == 0; - }); + hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes, + hipModule_t hmod, const char* name); + hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes, + const char* name); + private: + agent_globals_impl* impl; + }; - return it == l ? - std::make_pair(nullptr, 0u) : std::make_pair(it->address, it->byte_cnt); -} - -std::vector read_agent_globals(hsa_agent_t agent, - hsa_executable_t executable); -hsa_agent_t this_agent(); - - -class agent_globals_impl { -private: - std::pair< - std::mutex, - std::unordered_map< - std::string, std::vector>> globals_from_module; - - std::unordered_map< - hsa_agent_t, - std::pair< - std::once_flag, - std::vector>> globals_from_process; - -public: - - hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes, - hipModule_t hmod, const char* name) { - // the key of the map would the hash of code object associated with the - // hipModule_t instance - std::string key(hash_for(hmod)); - - if (globals_from_module.second.count(key) == 0) { - std::lock_guard lck{globals_from_module.first}; - - if (globals_from_module.second.count(key) == 0) { - globals_from_module.second.emplace( - key, read_agent_globals(this_agent(), executable_for(hmod))); - } - } - - const auto it0 = globals_from_module.second.find(key); - if (it0 == globals_from_module.second.cend()) { - hip_throw( - std::runtime_error{"agent_globals data structure corrupted."}); - } - - std::tie(*dptr, *bytes) = read_global_description(it0->second.cbegin(), - it0->second.cend(), name); - - return *dptr ? hipSuccess : hipErrorNotFound; + inline + __attribute__((visibility("hidden"))) + agent_globals& get_agent_globals() { + static agent_globals ag; + return ag; } + extern "C" + inline + __attribute__((visibility("hidden"))) hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes, - const char* name) { - - auto agent = this_agent(); - - std::call_once(globals_from_process[agent].first, [this](hsa_agent_t aa) { - std::vector tmp0; - for (auto&& executable : hip_impl::get_program_state().executables(aa)) { - auto tmp1 = read_agent_globals(aa, executable); - tmp0.insert(tmp0.end(), make_move_iterator(tmp1.begin()), - make_move_iterator(tmp1.end())); - } - globals_from_process[aa].second = move(move(tmp0)); - }, agent); - - const auto it = globals_from_process.find(agent); - - if (it == globals_from_process.cend()) return hipErrorNotInitialized; - - std::tie(*dptr, *bytes) = read_global_description(it->second.second.cbegin(), - it->second.second.cend(), name); - - return *dptr ? hipSuccess : hipErrorNotFound; + const char* name) { + return get_agent_globals().read_agent_global_from_process(dptr, bytes, name); } - -}; - -class agent_globals { -public: - agent_globals() : impl(new agent_globals_impl()) { - if (!impl) - hip_throw( - std::runtime_error{"Error when constructing agent global data structures."}); - } - ~agent_globals() { delete impl; } - - hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes, - hipModule_t hmod, const char* name) { - return impl->read_agent_global_from_module(dptr, bytes, hmod, name); - } - - hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes, - const char* name) { - return impl->read_agent_global_from_process(dptr, bytes, name); - } - -private: - agent_globals_impl* impl; -}; - -inline -__attribute__((visibility("hidden"))) -agent_globals& get_agent_globals() { - static agent_globals ag; - return ag; -} - - -extern "C" -inline -__attribute__((visibility("hidden"))) -hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes, - const char* name) { - return get_agent_globals().read_agent_global_from_process(dptr, bytes, name); -} - - } // Namespace hip_impl. #if defined(__cplusplus) diff --git a/src/hip_module.cpp b/src/hip_module.cpp index e8a8801e98..2b41988ac5 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -289,16 +289,6 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, 0)); } -hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, - hipModule_t hmod, const char* name) { - HIP_INIT_API(hipModuleGetGlobal, dptr, bytes, hmod, name); - if (!dptr || !bytes || !hmod) return hipErrorInvalidValue; - - if (!name) return hipErrorNotInitialized; - - return hip_impl::get_agent_globals().read_agent_global_from_module(dptr, bytes, hmod, name); -} - namespace hip_impl { hsa_executable_t executable_for(hipModule_t hmod) { return hmod->executable; @@ -323,10 +313,159 @@ namespace hip_impl { return currentDevice->_hsaAgent; } + + struct Agent_global { + Agent_global() : name(nullptr), address(nullptr), byte_cnt(0) {} + Agent_global(const char* name, hipDeviceptr_t address, uint32_t byte_cnt) + : name(nullptr), address(address), byte_cnt(byte_cnt) { + if (name) + this->name = strdup(name); + } + + Agent_global& operator=(Agent_global&& t) { + if (this == &t) return *this; + + if (name) free(name); + name = t.name; + address = t.address; + byte_cnt = t.byte_cnt; + + t.name = nullptr; + t.address = nullptr; + t.byte_cnt = 0; + + return *this; + } + + Agent_global(Agent_global&& t) + : name(nullptr), address(nullptr), byte_cnt(0) { + *this = std::move(t); + } + + // not needed, delete them to prevent bugs + Agent_global(const Agent_global&) = delete; + Agent_global& operator=(Agent_global& t) = delete; + + ~Agent_global() { if (name) free(name); } + + char* name; + hipDeviceptr_t address; + uint32_t byte_cnt; + }; + + template + std::pair read_global_description( + ForwardIterator f, ForwardIterator l, const char* name) { + const auto it = std::find_if(f, l, [=](const Agent_global& x) { + return strcmp(x.name, name) == 0; + }); + + return it == l ? + std::make_pair(nullptr, 0u) : std::make_pair(it->address, it->byte_cnt); + } + + std::vector read_agent_globals(hsa_agent_t agent, + hsa_executable_t executable); + class agent_globals_impl { + private: + std::pair< + std::mutex, + std::unordered_map< + std::string, std::vector>> globals_from_module; + + std::unordered_map< + hsa_agent_t, + std::pair< + std::once_flag, + std::vector>> globals_from_process; + + public: + + hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes, + hipModule_t hmod, const char* name) { + // the key of the map would the hash of code object associated with the + // hipModule_t instance + std::string key(hash_for(hmod)); + + if (globals_from_module.second.count(key) == 0) { + std::lock_guard lck{globals_from_module.first}; + + if (globals_from_module.second.count(key) == 0) { + globals_from_module.second.emplace( + key, read_agent_globals(this_agent(), executable_for(hmod))); + } + } + + const auto it0 = globals_from_module.second.find(key); + if (it0 == globals_from_module.second.cend()) { + hip_throw( + std::runtime_error{"agent_globals data structure corrupted."}); + } + + std::tie(*dptr, *bytes) = read_global_description(it0->second.cbegin(), + it0->second.cend(), name); + + return *dptr ? hipSuccess : hipErrorNotFound; + } + + hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes, + const char* name) { + + auto agent = this_agent(); + + std::call_once(globals_from_process[agent].first, [this](hsa_agent_t aa) { + std::vector tmp0; + for (auto&& executable : hip_impl::get_program_state().executables(aa)) { + auto tmp1 = read_agent_globals(aa, executable); + tmp0.insert(tmp0.end(), make_move_iterator(tmp1.begin()), + make_move_iterator(tmp1.end())); + } + globals_from_process[aa].second = move(move(tmp0)); + }, agent); + + const auto it = globals_from_process.find(agent); + + if (it == globals_from_process.cend()) return hipErrorNotInitialized; + + std::tie(*dptr, *bytes) = read_global_description(it->second.second.cbegin(), + it->second.second.cend(), name); + + return *dptr ? hipSuccess : hipErrorNotFound; + } + + }; + + agent_globals::agent_globals() : impl(new agent_globals_impl()) { + if (!impl) + hip_throw( + std::runtime_error{"Error when constructing agent global data structures."}); + } + agent_globals::~agent_globals() { delete impl; } + + hipError_t agent_globals::read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes, + hipModule_t hmod, const char* name) { + return impl->read_agent_global_from_module(dptr, bytes, hmod, name); + } + + hipError_t agent_globals::read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes, + const char* name) { + return impl->read_agent_global_from_process(dptr, bytes, name); + } + } // Namespace hip_impl. +hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, + hipModule_t hmod, const char* name) { + HIP_INIT_API(hipModuleGetGlobal, dptr, bytes, hmod, name); + if (!dptr || !bytes || !hmod) return hipErrorInvalidValue; + + if (!name) return hipErrorNotInitialized; + + return hip_impl::get_agent_globals().read_agent_global_from_module(dptr, bytes, hmod, name); +} + namespace { -inline void track(const Agent_global& x, hsa_agent_t agent) { +inline void track(const hip_impl::Agent_global& x, hsa_agent_t agent) { tprintf(DB_MEM, " add variable '%s' with ptr=%p size=%u to tracker\n", x.name, x.address, x.byte_cnt); @@ -347,7 +486,7 @@ inline void track(const Agent_global& x, hsa_agent_t agent) { } -template > +template > inline hsa_status_t copy_agent_global_variables(hsa_executable_t, hsa_agent_t agent, hsa_executable_symbol_t x, void* out) { using namespace hip_impl; @@ -358,7 +497,7 @@ inline hsa_status_t copy_agent_global_variables(hsa_executable_t, hsa_agent_t ag hsa_executable_symbol_get_info(x, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &t); if (t == HSA_SYMBOL_KIND_VARIABLE) { - Agent_global tmp(name(x).c_str(), address(x), size(x)); + hip_impl::Agent_global tmp(name(x).c_str(), address(x), size(x)); static_cast(out)->push_back(std::move(tmp)); track(static_cast(out)->back(),agent); From 80fec2b477c52c1fd28f1b9f6e1b56a3d81d738e Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Sat, 4 May 2019 14:35:27 -0400 Subject: [PATCH 03/19] remove executables() from program_state --- include/hip/hcc_detail/program_state.hpp | 10 ++-------- src/hip_module.cpp | 2 +- src/program_state.cpp | 4 ---- 3 files changed, 3 insertions(+), 13 deletions(-) diff --git a/include/hip/hcc_detail/program_state.hpp b/include/hip/hcc_detail/program_state.hpp index 19db8f9e0d..5741c22b4b 100644 --- a/include/hip/hcc_detail/program_state.hpp +++ b/include/hip/hcc_detail/program_state.hpp @@ -31,9 +31,6 @@ THE SOFTWARE. #include #include #include -#include -#include -#include struct ihipModuleSymbol_t; using hipFunction_t = ihipModuleSymbol_t*; @@ -72,6 +69,7 @@ class program_state { public: program_state(); ~program_state(); + program_state(const program_state&) = delete; hipFunction_t kernel_descriptor(std::uintptr_t, hsa_agent_t); @@ -83,12 +81,8 @@ public: void* global_addr_by_name(const char* name); - // to fix later - const std::vector& executables(hsa_agent_t agent); - - program_state(const program_state&) = delete; - private: + friend class agent_globals_impl; program_state_impl* impl; }; diff --git a/src/hip_module.cpp b/src/hip_module.cpp index 2b41988ac5..4ff6a987fd 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -415,7 +415,7 @@ namespace hip_impl { std::call_once(globals_from_process[agent].first, [this](hsa_agent_t aa) { std::vector tmp0; - for (auto&& executable : hip_impl::get_program_state().executables(aa)) { + for (auto&& executable : hip_impl::get_program_state().impl->get_executables(aa)) { auto tmp1 = read_agent_globals(aa, executable); tmp0.insert(tmp0.end(), make_move_iterator(tmp1.begin()), make_move_iterator(tmp1.end())); diff --git a/src/program_state.cpp b/src/program_state.cpp index 6783f85b9b..e226e1af66 100644 --- a/src/program_state.cpp +++ b/src/program_state.cpp @@ -45,10 +45,6 @@ hsa_executable_t program_state::load_executable(const char* data, return impl->load_executable(data, data_size, executable, agent); } -const std::vector& program_state::executables(hsa_agent_t agent) { - return impl->get_executables(agent); -} - hipFunction_t program_state::kernel_descriptor(std::uintptr_t function_address, hsa_agent_t agent) { auto& kd = impl->kernel_descriptor(function_address, agent); From 00824be34c0021f662f9feeae172475275b6cba2 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Sat, 4 May 2019 15:08:28 -0400 Subject: [PATCH 04/19] move executable_cache into program_state.cpp --- include/hip/hcc_detail/program_state.hpp | 27 -------- src/hip_hcc.cpp | 9 --- src/program_state.cpp | 87 +++++++++++++----------- src/program_state.inl | 24 +++++++ 4 files changed, 71 insertions(+), 76 deletions(-) diff --git a/include/hip/hcc_detail/program_state.hpp b/include/hip/hcc_detail/program_state.hpp index 5741c22b4b..b3656e1e48 100644 --- a/include/hip/hcc_detail/program_state.hpp +++ b/include/hip/hcc_detail/program_state.hpp @@ -30,39 +30,12 @@ THE SOFTWARE. #include #include #include -#include struct ihipModuleSymbol_t; using hipFunction_t = ihipModuleSymbol_t*; -namespace std { -template<> -struct hash { - size_t operator()(hsa_agent_t x) const { - return hash{}(x.handle); - } -}; - -template<> -struct hash { - size_t operator()(hsa_isa_t x) const { - return hash{}(x.handle); - } -}; -} // namespace std - -inline constexpr bool operator==(hsa_agent_t x, hsa_agent_t y) { - return x.handle == y.handle; -} -inline constexpr bool operator==(hsa_isa_t x, hsa_isa_t y) { - return x.handle == y.handle; -} - namespace hip_impl { -[[noreturn]] -void hip_throw(const std::exception&); - class kernargs_size_align; class program_state_impl; class program_state { diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index 648b53dab7..583017a30a 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -2492,13 +2492,4 @@ namespace hip_impl { #endif } - std::mutex executables_cache_mutex; - - std::vector& executables_cache( - std::string elf, hsa_isa_t isa, hsa_agent_t agent) { - static std::unordered_map>>> cache; - return cache[elf][isa][agent]; - } } // Namespace hip_impl. diff --git a/src/program_state.cpp b/src/program_state.cpp index e226e1af66..906584bc38 100644 --- a/src/program_state.cpp +++ b/src/program_state.cpp @@ -1,4 +1,6 @@ #include "../include/hip/hcc_detail/program_state.hpp" +// contains implementation of program_state_impl +#include "program_state.inl" #include @@ -7,53 +9,58 @@ #include #include -// contains implementation of program_state_impl -#include "program_state.inl" - namespace hip_impl { -std::size_t kernargs_size_align::kernargs_size_align::size(std::size_t n) const{ - return (*reinterpret_cast>*>(handle))[n].first; -} + std::size_t kernargs_size_align::kernargs_size_align::size(std::size_t n) const{ + return (*reinterpret_cast>*>(handle))[n].first; + } -std::size_t kernargs_size_align::alignment(std::size_t n) const{ - return (*reinterpret_cast>*>(handle))[n].second; -} + std::size_t kernargs_size_align::alignment(std::size_t n) const{ + return (*reinterpret_cast>*>(handle))[n].second; + } -program_state::program_state() : - impl(new program_state_impl) { - if (!impl) hip_throw(std::runtime_error { - "Unknown error when constructing program state."}); -} + program_state::program_state() : impl(new program_state_impl) { + if (!impl) hip_throw(std::runtime_error { + "Unknown error when constructing program state."}); + } -program_state::~program_state() { - delete(impl); -} + program_state::~program_state() { + delete(impl); + } -void* program_state::global_addr_by_name(const char* name) { - const auto it = impl->get_globals().find(name); - if (it == impl->get_globals().end()) - return nullptr; - else - return it->second; -} + void* program_state::global_addr_by_name(const char* name) { + const auto it = impl->get_globals().find(name); + if (it == impl->get_globals().end()) + return nullptr; + else + return it->second; + } -hsa_executable_t program_state::load_executable(const char* data, - const size_t data_size, - hsa_executable_t executable, - hsa_agent_t agent) { - return impl->load_executable(data, data_size, executable, agent); -} + hsa_executable_t program_state::load_executable(const char* data, + const size_t data_size, + hsa_executable_t executable, + hsa_agent_t agent) { + return impl->load_executable(data, data_size, executable, agent); + } -hipFunction_t program_state::kernel_descriptor(std::uintptr_t function_address, - hsa_agent_t agent) { - auto& kd = impl->kernel_descriptor(function_address, agent); - return kd; -} + hipFunction_t program_state::kernel_descriptor(std::uintptr_t function_address, + hsa_agent_t agent) { + auto& kd = impl->kernel_descriptor(function_address, agent); + return kd; + } -kernargs_size_align program_state::get_kernargs_size_align(std::uintptr_t kernel) { - kernargs_size_align t; - t.handle = reinterpret_cast(&impl->kernargs_size_align(kernel)); - return t; -} + kernargs_size_align program_state::get_kernargs_size_align(std::uintptr_t kernel) { + kernargs_size_align t; + t.handle = reinterpret_cast(&impl->kernargs_size_align(kernel)); + return t; + } + + std::mutex executables_cache_mutex; + std::vector& executables_cache( + std::string elf, hsa_isa_t isa, hsa_agent_t agent) { + static std::unordered_map>>> cache; + return cache[elf][isa][agent]; + } }; diff --git a/src/program_state.inl b/src/program_state.inl index f1397b3fe9..92b2c74597 100644 --- a/src/program_state.inl +++ b/src/program_state.inl @@ -26,12 +26,36 @@ #include #include #include +#include #include #include #include #include #include +namespace std { +template<> +struct hash { + size_t operator()(hsa_agent_t x) const { + return hash{}(x.handle); + } +}; + +template<> +struct hash { + size_t operator()(hsa_isa_t x) const { + return hash{}(x.handle); + } +}; +} // namespace std + +inline constexpr bool operator==(hsa_agent_t x, hsa_agent_t y) { + return x.handle == y.handle; +} +inline constexpr bool operator==(hsa_isa_t x, hsa_isa_t y) { + return x.handle == y.handle; +} + namespace hip_impl { [[noreturn]] From fc08f2973597e0904b8b2d33d3021a359ce24da5 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Sat, 4 May 2019 17:51:13 -0400 Subject: [PATCH 05/19] replace std::vector for kernarg --- .../hip/hcc_detail/functional_grid_launch.hpp | 12 ++++---- include/hip/hcc_detail/program_state.hpp | 14 ++++++++++ src/program_state.cpp | 28 +++++++++++++++++++ src/program_state.inl | 5 ++++ 4 files changed, 53 insertions(+), 6 deletions(-) diff --git a/include/hip/hcc_detail/functional_grid_launch.hpp b/include/hip/hcc_detail/functional_grid_launch.hpp index 8082188489..3e464a3326 100644 --- a/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/include/hip/hcc_detail/functional_grid_launch.hpp @@ -57,10 +57,10 @@ template < std::size_t n, typename... Ts, typename std::enable_if::type* = nullptr> -inline std::vector make_kernarg( +inline hip_impl::kernarg make_kernarg( const std::tuple&, const kernargs_size_align&, - std::vector kernarg) { + hip_impl::kernarg kernarg) { return kernarg; } @@ -68,10 +68,10 @@ template < std::size_t n, typename... Ts, typename std::enable_if::type* = nullptr> -inline std::vector make_kernarg( +inline hip_impl::kernarg make_kernarg( const std::tuple& formals, const kernargs_size_align& size_align, - std::vector kernarg) { + hip_impl::kernarg kernarg) { using T = typename std::tuple_element>::type; static_assert( @@ -96,7 +96,7 @@ inline std::vector make_kernarg( } template -inline std::vector make_kernarg( +inline hip_impl::kernarg make_kernarg( void (*kernel)(Formals...), std::tuple actuals) { static_assert(sizeof...(Formals) == sizeof...(Actuals), "The count of formal arguments must match the count of actuals."); @@ -104,7 +104,7 @@ inline std::vector make_kernarg( if (sizeof...(Formals) == 0) return {}; std::tuple to_formals{std::move(actuals)}; - std::vector kernarg; + hip_impl::kernarg kernarg; kernarg.reserve(sizeof(to_formals)); auto& ps = hip_impl::get_program_state(); diff --git a/include/hip/hcc_detail/program_state.hpp b/include/hip/hcc_detail/program_state.hpp index b3656e1e48..c64b64fde8 100644 --- a/include/hip/hcc_detail/program_state.hpp +++ b/include/hip/hcc_detail/program_state.hpp @@ -36,6 +36,20 @@ using hipFunction_t = ihipModuleSymbol_t*; namespace hip_impl { +struct kernarg_impl; +class kernarg { +public: + kernarg(); + kernarg(kernarg&&); + ~kernarg(); + std::uint8_t* data(); + std::size_t size(); + void reserve(std::size_t); + void resize(std::size_t); +private: + kernarg_impl* impl; +}; + class kernargs_size_align; class program_state_impl; class program_state { diff --git a/src/program_state.cpp b/src/program_state.cpp index 906584bc38..dbd7d3ebc4 100644 --- a/src/program_state.cpp +++ b/src/program_state.cpp @@ -10,6 +10,34 @@ #include namespace hip_impl { + + kernarg::kernarg() : impl(new kernarg_impl) { + } + + kernarg::kernarg(kernarg&& k) : impl(k.impl) { + k.impl = nullptr; + } + + kernarg::~kernarg() { + if (impl) + delete(impl); + } + + std::uint8_t* kernarg::data() { + return impl->v.data(); + } + + std::size_t kernarg::size() { + return impl->v.size(); + } + + void kernarg::reserve(std::size_t c) { + impl->v.reserve(c); + } + + void kernarg::resize(std::size_t c) { + impl->v.resize(c); + } std::size_t kernargs_size_align::kernargs_size_align::size(std::size_t n) const{ return (*reinterpret_cast>*>(handle))[n].first; diff --git a/src/program_state.inl b/src/program_state.inl index 92b2c74597..9fc590ef75 100644 --- a/src/program_state.inl +++ b/src/program_state.inl @@ -743,4 +743,9 @@ public: } }; // class program_state_impl +struct kernarg_impl { + std::vector v; +}; + + }; From cec09269247dee5811103330e2b1959837d32bb2 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Mon, 6 May 2019 12:03:09 -0400 Subject: [PATCH 06/19] fix breakage due to compiling in C++17 --- include/hip/hcc_detail/helpers.hpp | 11 ++++++++++- 1 file changed, 10 insertions(+), 1 deletion(-) diff --git a/include/hip/hcc_detail/helpers.hpp b/include/hip/hcc_detail/helpers.hpp index 1916945c1a..9366885f8b 100644 --- a/include/hip/hcc_detail/helpers.hpp +++ b/include/hip/hcc_detail/helpers.hpp @@ -84,12 +84,21 @@ struct is_callable_impl()(std::de // Not callable. template struct is_callable_impl : std::false_type {}; -#else +#elif (__cplusplus < 201703L) template struct is_callable_impl : std::false_type {}; template struct is_callable_impl > > : std::true_type {}; +#else + +// C++17 + +template +struct is_callable_impl : std::false_type {}; + +template +struct is_callable_impl > > : std::true_type {}; #endif template struct is_callable : is_callable_impl {}; From 54f94ed02f6fa331d91d351dc04c9632b35497a9 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Fri, 10 May 2019 17:23:23 -0400 Subject: [PATCH 07/19] remove code_object_bundle.hpp, clean up the old Agent_global --- .../hip/hcc_detail/functional_grid_launch.hpp | 1 - include/hip/hcc_detail/hip_runtime_api.h | 42 ------------------- 2 files changed, 43 deletions(-) diff --git a/include/hip/hcc_detail/functional_grid_launch.hpp b/include/hip/hcc_detail/functional_grid_launch.hpp index 3e464a3326..b09d0d2c6c 100644 --- a/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/include/hip/hcc_detail/functional_grid_launch.hpp @@ -22,7 +22,6 @@ THE SOFTWARE. #pragma once -#include "code_object_bundle.hpp" #include "concepts.hpp" #include "helpers.hpp" #include "program_state.hpp" diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 17a42e2a16..5dd7227f61 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -2590,48 +2590,6 @@ hipError_t hipModuleGetFunction(hipFunction_t* function, hipModule_t module, con hipError_t hipFuncGetAttributes(struct hipFuncAttributes* attr, const void* func); -#if 0 -struct Agent_global { - - Agent_global() : name(nullptr), address(nullptr), byte_cnt(0) {} - Agent_global(const char* name, hipDeviceptr_t address, uint32_t byte_cnt) - : name(nullptr), address(address), byte_cnt(byte_cnt) { - if (name) - this->name = strdup(name); - } - - Agent_global& operator=(Agent_global&& t) { - if (this == &t) return *this; - - if (name) free(name); - name = t.name; - address = t.address; - byte_cnt = t.byte_cnt; - - t.name = nullptr; - t.address = nullptr; - t.byte_cnt = 0; - - return *this; - } - - Agent_global(Agent_global&& t) - : name(nullptr), address(nullptr), byte_cnt(0) { - *this = std::move(t); - } - - // not needed, delete them to prevent bugs - Agent_global(const Agent_global&) = delete; - Agent_global& operator=(Agent_global& t) = delete; - - ~Agent_global() { if (name) free(name); } - - char* name; - hipDeviceptr_t address; - uint32_t byte_cnt; -}; -#endif - #if !__HIP_VDI__ #if defined(__cplusplus) } // extern "C" From cc0f3445bb1a0652f951df3dc21f8df3f269b2ef Mon Sep 17 00:00:00 2001 From: Zuhaib Khan Date: Tue, 28 May 2019 16:57:51 -0400 Subject: [PATCH 08/19] Structured hipFloatComplex as typedef of float2, and hipDoubleComplex as typedef of double2. --- include/hip/hcc_detail/hip_complex.h | 220 +++++++++++---------------- 1 file changed, 92 insertions(+), 128 deletions(-) diff --git a/include/hip/hcc_detail/hip_complex.h b/include/hip/hcc_detail/hip_complex.h index 128e2d670b..75930c469e 100644 --- a/include/hip/hcc_detail/hip_complex.h +++ b/include/hip/hcc_detail/hip_complex.h @@ -120,51 +120,102 @@ THE SOFTWARE. ret.y = lhs.y * rhs; \ return ret; \ } -#define MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ComplexT, T) \ - explicit __device__ __host__ ComplexT(T val) : x(val), y(val) {} \ - __device__ __host__ ComplexT(T val1, T val2) : x(val1), y(val2) {} #endif -struct hipFloatComplex { -#ifdef __cplusplus - public: - typedef float value_type; - __device__ __host__ hipFloatComplex() : x(0.0f), y(0.0f) {} - explicit __device__ __host__ hipFloatComplex(float x) : x(x), y(0.0f) {} - __device__ __host__ hipFloatComplex(float x, float y) : x(x), y(y) {} - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, signed short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, signed int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, double) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, signed long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, signed long long) -#endif - float x, y; -} __attribute__((aligned(8))); +typedef float2 hipFloatComplex; + +__device__ __host__ static inline float hipCrealf(hipFloatComplex z) { return z.x; } + +__device__ __host__ static inline float hipCimagf(hipFloatComplex z) { return z.y; } + +__device__ __host__ static inline hipFloatComplex make_hipFloatComplex(float a, float b) { + hipFloatComplex z; + z.x = a; + z.y = b; + return z; +} + +__device__ __host__ static inline hipFloatComplex hipConjf(hipFloatComplex z) { + hipFloatComplex ret; + ret.x = z.x; + ret.y = -z.y; + return ret; +} + +__device__ __host__ static inline float hipCsqabsf(hipFloatComplex z) { + return z.x * z.x + z.y * z.y; +} + +__device__ __host__ static inline hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q) { + return make_hipFloatComplex(p.x + q.x, p.y + q.y); +} + +__device__ __host__ static inline hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q) { + return make_hipFloatComplex(p.x - q.x, p.y - q.y); +} + +__device__ __host__ static inline hipFloatComplex hipCmulf(hipFloatComplex p, hipFloatComplex q) { + return make_hipFloatComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y); +} + +__device__ __host__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q) { + float sqabs = hipCsqabsf(q); + hipFloatComplex ret; + ret.x = (p.x * q.x + p.y * q.y) / sqabs; + ret.y = (p.y * q.x - p.x * q.y) / sqabs; + return ret; +} + +__device__ __host__ static inline float hipCabsf(hipFloatComplex z) { return sqrtf(hipCsqabsf(z)); } + + +typedef double2 hipDoubleComplex; + +__device__ __host__ static inline double hipCreal(hipDoubleComplex z) { return z.x; } + +__device__ __host__ static inline double hipCimag(hipDoubleComplex z) { return z.y; } + +__device__ __host__ static inline hipDoubleComplex make_hipDoubleComplex(double a, double b) { + hipDoubleComplex z; + z.x = a; + z.y = b; + return z; +} + +__device__ __host__ static inline hipDoubleComplex hipConj(hipDoubleComplex z) { + hipDoubleComplex ret; + ret.x = z.x; + ret.y = z.y; + return ret; +} + +__device__ __host__ static inline double hipCsqabs(hipDoubleComplex z) { + return z.x * z.x + z.y * z.y; +} + +__device__ __host__ static inline hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q) { + return make_hipDoubleComplex(p.x + q.x, p.y + q.y); +} + +__device__ __host__ static inline hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q) { + return make_hipDoubleComplex(p.x - q.x, p.y - q.y); +} + +__device__ __host__ static inline hipDoubleComplex hipCmul(hipDoubleComplex p, hipDoubleComplex q) { + return make_hipDoubleComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y); +} + +__device__ __host__ static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q) { + double sqabs = hipCsqabs(q); + hipDoubleComplex ret; + ret.x = (p.x * q.x + p.y * q.y) / sqabs; + ret.y = (p.y * q.x - p.x * q.y) / sqabs; + return ret; +} + +__device__ __host__ static inline double hipCabs(hipDoubleComplex z) { return sqrtf(hipCsqabs(z)); } -struct hipDoubleComplex { -#ifdef __cplusplus - public: - typedef double value_type; - __device__ __host__ hipDoubleComplex() : x(0.0f), y(0.0f) {} - explicit __device__ __host__ hipDoubleComplex(double x) : x(x), y(0.0f) {} - __device__ __host__ hipDoubleComplex(double x, double y) : x(x), y(y) {} - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, signed short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, signed int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, float) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, signed long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, signed long long) -#endif - double x, y; -} __attribute__((aligned(16))); #if __cplusplus @@ -214,93 +265,6 @@ COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned long long) #endif -__device__ __host__ static inline float hipCrealf(hipFloatComplex z) { return z.x; } - -__device__ __host__ static inline float hipCimagf(hipFloatComplex z) { return z.y; } - -__device__ __host__ static inline hipFloatComplex make_hipFloatComplex(float a, float b) { - hipFloatComplex z; - z.x = a; - z.y = b; - return z; -} - -__device__ __host__ static inline hipFloatComplex hipConjf(hipFloatComplex z) { - hipFloatComplex ret; - ret.x = z.x; - ret.y = -z.y; - return ret; -} - -__device__ __host__ static inline float hipCsqabsf(hipFloatComplex z) { - return z.x * z.x + z.y * z.y; -} - -__device__ __host__ static inline hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q) { - return make_hipFloatComplex(p.x + q.x, p.y + q.y); -} - -__device__ __host__ static inline hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q) { - return make_hipFloatComplex(p.x - q.x, p.y - q.y); -} - -__device__ __host__ static inline hipFloatComplex hipCmulf(hipFloatComplex p, hipFloatComplex q) { - return make_hipFloatComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y); -} - -__device__ __host__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q) { - float sqabs = hipCsqabsf(q); - hipFloatComplex ret; - ret.x = (p.x * q.x + p.y * q.y) / sqabs; - ret.y = (p.y * q.x - p.x * q.y) / sqabs; - return ret; -} - -__device__ __host__ static inline float hipCabsf(hipFloatComplex z) { return sqrtf(hipCsqabsf(z)); } - -__device__ __host__ static inline double hipCreal(hipDoubleComplex z) { return z.x; } - -__device__ __host__ static inline double hipCimag(hipDoubleComplex z) { return z.y; } - -__device__ __host__ static inline hipDoubleComplex make_hipDoubleComplex(double a, double b) { - hipDoubleComplex z; - z.x = a; - z.y = b; - return z; -} - -__device__ __host__ static inline hipDoubleComplex hipConj(hipDoubleComplex z) { - hipDoubleComplex ret; - ret.x = z.x; - ret.y = z.y; - return ret; -} - -__device__ __host__ static inline double hipCsqabs(hipDoubleComplex z) { - return z.x * z.x + z.y * z.y; -} - -__device__ __host__ static inline hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q) { - return make_hipDoubleComplex(p.x + q.x, p.y + q.y); -} - -__device__ __host__ static inline hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q) { - return make_hipDoubleComplex(p.x - q.x, p.y - q.y); -} - -__device__ __host__ static inline hipDoubleComplex hipCmul(hipDoubleComplex p, hipDoubleComplex q) { - return make_hipDoubleComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y); -} - -__device__ __host__ static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q) { - double sqabs = hipCsqabs(q); - hipDoubleComplex ret; - ret.x = (p.x * q.x + p.y * q.y) / sqabs; - ret.y = (p.y * q.x - p.x * q.y) / sqabs; - return ret; -} - -__device__ __host__ static inline double hipCabs(hipDoubleComplex z) { return sqrtf(hipCsqabs(z)); } typedef hipFloatComplex hipComplex; From f34654f8aa4926a140cd263700aa06822afd3d86 Mon Sep 17 00:00:00 2001 From: Aryan Salmanpour Date: Thu, 30 May 2019 18:04:05 -0400 Subject: [PATCH 09/19] Header change for new hip API hipExtLaunchMultiKernelMultiDevice --- include/hip/hcc_detail/hip_runtime_api.h | 21 +++++++++++++++++++++ 1 file changed, 21 insertions(+) diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index ba8d2a21c5..1b332cdb85 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -2906,6 +2906,21 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor( hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( int* numBlocks, const void* f, int blockSize, size_t dynamicSMemSize, unsigned int flags); +/** + * @brief Launches kernels on multiple devices and guarantees all specified kernels are dispatched + * on respective streams before enqueuing any other work on the specified streams from any other threads + * + * + * @param [in] hipLaunchParams List of launch parameters, one per device. + * @param [in] numDevices Size of the launchParamsList array. + * @param [in] flags Flags to control launch behavior. + * + * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue + */ +hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList, + int numDevices, unsigned int flags); + + // doxygen end Version Management /** @@ -3262,6 +3277,12 @@ inline hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchP return hipLaunchCooperativeKernelMultiDevice(launchParamsList, numDevices, flags); } +template +inline hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList, + unsigned int numDevices, unsigned int flags = 0) { + return hipExtLaunchMultiKernelMultiDevice(launchParamsList, numDevices, flags); +} + /* * @brief Unbinds the textuer bound to @p tex From bd4d1838f21c6065b2fced47ee7314b425cb9cdb Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 31 May 2019 16:39:33 +0530 Subject: [PATCH 10/19] [dtests] Temporarily disable hiprtc tests Change-Id: I87c0c01837e7b59b11d99fb94d679a765f914da5 --- tests/src/hiprtc/hiprtcGetLoweredName.cpp | 2 +- tests/src/hiprtc/hiprtcGetTypeName.cpp | 2 +- tests/src/hiprtc/saxpy.cpp | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/tests/src/hiprtc/hiprtcGetLoweredName.cpp b/tests/src/hiprtc/hiprtcGetLoweredName.cpp index e3fa057a81..d9e6d71f93 100644 --- a/tests/src/hiprtc/hiprtcGetLoweredName.cpp +++ b/tests/src/hiprtc/hiprtcGetLoweredName.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. */ /* HIT_START * BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc - * TEST: %t + * TEST: %t EXCLUDE_HIP_PLATFORM hcc * HIT_END */ #include diff --git a/tests/src/hiprtc/hiprtcGetTypeName.cpp b/tests/src/hiprtc/hiprtcGetTypeName.cpp index b0348408f3..7f49ea984b 100644 --- a/tests/src/hiprtc/hiprtcGetTypeName.cpp +++ b/tests/src/hiprtc/hiprtcGetTypeName.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. */ /* HIT_START * BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc - * TEST: %t + * TEST: %t EXCLUDE_HIP_PLATFORM hcc * HIT_END */ diff --git a/tests/src/hiprtc/saxpy.cpp b/tests/src/hiprtc/saxpy.cpp index 5f9dc7a125..437420266d 100644 --- a/tests/src/hiprtc/saxpy.cpp +++ b/tests/src/hiprtc/saxpy.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. */ /* HIT_START * BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc - * TEST: %t + * TEST: %t EXCLUDE_HIP_PLATFORM hcc * HIT_END */ From 80bcf0785a7a6bf9a66ac5dfef20d53f5966aad2 Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Fri, 31 May 2019 12:07:58 -0400 Subject: [PATCH 11/19] Add device_builtin_texture_type attribute to texture type for hip-clang This is required to support texture type for hip-clang. --- include/hip/hcc_detail/hip_texture_types.h | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/include/hip/hcc_detail/hip_texture_types.h b/include/hip/hcc_detail/hip_texture_types.h index 0a68b507e8..b229f4e696 100644 --- a/include/hip/hcc_detail/hip_texture_types.h +++ b/include/hip/hcc_detail/hip_texture_types.h @@ -45,10 +45,15 @@ THE SOFTWARE. * * * * *******************************************************************************/ +#if __HIP__ +#define __HIP_TEXTURE_ATTRIB __attribute__((device_builtin_texture_type)) +#else +#define __HIP_TEXTURE_ATTRIB +#endif template -struct texture : public textureReference { +struct __HIP_TEXTURE_ATTRIB texture : public textureReference { texture(int norm = 0, enum hipTextureFilterMode fMode = hipFilterModePoint, enum hipTextureAddressMode aMode = hipAddressModeClamp) { normalized = norm; From a489f583bb699b066a7c58116c11169138a56389 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 31 May 2019 22:30:24 +0530 Subject: [PATCH 12/19] Fix wrong grid dim shown in trace --- src/hip_module.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/hip_module.cpp b/src/hip_module.cpp index e8a8801e98..0b98af5a82 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -189,7 +189,7 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, lp.dynamic_group_mem_bytes = sharedMemBytes; // TODO - this should be part of preLaunchKernel. hStream = ihipPreLaunchKernel( - hStream, dim3(globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ), + hStream, dim3(globalWorkSizeX/localWorkSizeX, globalWorkSizeY/localWorkSizeY, globalWorkSizeZ/localWorkSizeZ), dim3(localWorkSizeX, localWorkSizeY, localWorkSizeZ), &lp, f->_name.c_str()); From 71f6bf4e678b5124df762e89197e921a793fc6c3 Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Fri, 31 May 2019 23:58:59 -0400 Subject: [PATCH 13/19] Fix default HIP_VDI_HOME There is soft link /opt/rocm/bin/.hipVersion, therefore when hipcc is executed as /opt/rocm/bin/hipcc, it will set HIP_VDI_HOME to /opt/rocm, which is incorrect. Check ../lib/bitcode instead to identify HIP_VDI_HOME. --- bin/hipcc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/bin/hipcc b/bin/hipcc index a438f0afe5..58c8fe45c9 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -108,7 +108,7 @@ $HIP_RUNTIME= $hipConfig{'HIP_RUNTIME'}; # If using VDI runtime, need to find HIP_VDI_HOME if ($HIP_RUNTIME eq "VDI" and !defined $HIP_VDI_HOME) { my $hipcc_dir = dirname($0); - if (-e "$hipcc_dir/.hipVersion") { + if (-e "$hipcc_dir/../lib/bitcode") { $HIP_VDI_HOME = abs_path($hipcc_dir . "/.."); } else { $HIP_VDI_HOME = "/opt/rocm/hip"; From 9c03a5f948dacfa8127069c371eb135324b78b11 Mon Sep 17 00:00:00 2001 From: cdevadas Date: Thu, 2 May 2019 16:10:07 +0530 Subject: [PATCH 14/19] Runtime changes to append implicit kernel arguments. Appended 48 empty bytes to the kernarg area at runtime. The implicit arguments are enabled primarily for the hostcall services and it is completely abstracted from the user code. Enabled it for both hip-clang and hip-hcc. --- src/hip_module.cpp | 34 ++++++++++++++++++---------------- 1 file changed, 18 insertions(+), 16 deletions(-) diff --git a/src/hip_module.cpp b/src/hip_module.cpp index e8a8801e98..91544b82a8 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -55,6 +55,10 @@ THE SOFTWARE. using namespace ELFIO; using namespace std; +// For HIP implicit kernargs. +static const size_t HIP_IMPLICIT_KERNARG_SIZE = 48; +static const size_t HIP_IMPLICIT_KERNARG_ALIGNMENT = 8; + // calculate MD5 checksum inline std::string checksum(size_t size, const char *source) { // FNV-1a hashing, 64-bit version @@ -146,33 +150,28 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, ihipDevice_t* currentDevice = ihipGetDevice(deviceId); hsa_agent_t gpuAgent = (hsa_agent_t)currentDevice->_hsaAgent; - void* config[5] = {0}; - size_t kernArgSize; - - std::vector tmp{}; + std::vector kernargs{}; if (kernelParams) { if (extra) return hipErrorInvalidValue; for (auto&& x : f->_kernarg_layout) { const auto p{static_cast(*kernelParams)}; - tmp.insert( - tmp.cend(), + kernargs.insert( + kernargs.cend(), round_up_to_next_multiple_nonnegative( - tmp.size(), x.second) - tmp.size(), + kernargs.size(), x.second) - kernargs.size(), '\0'); - tmp.insert(tmp.cend(), p, p + x.first); + kernargs.insert(kernargs.cend(), p, p + x.first); ++kernelParams; } - config[1] = static_cast(tmp.data()); - - kernArgSize = tmp.size(); } else if (extra) { - memcpy(config, extra, sizeof(size_t) * 5); - if (config[0] == HIP_LAUNCH_PARAM_BUFFER_POINTER && - config[2] == HIP_LAUNCH_PARAM_BUFFER_SIZE && config[4] == HIP_LAUNCH_PARAM_END) { - kernArgSize = *(size_t*)(config[3]); + if (extra[0] == HIP_LAUNCH_PARAM_BUFFER_POINTER && + extra[2] == HIP_LAUNCH_PARAM_BUFFER_SIZE && extra[4] == HIP_LAUNCH_PARAM_END) { + auto args = (char*)extra[1]; + size_t argSize = *(size_t*)(extra[3]); + kernargs.insert(kernargs.end(), args, args+argSize); } else { return hipErrorNotInitialized; } @@ -181,6 +180,9 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, return hipErrorInvalidValue; } + // Insert 48-bytes at the end for implicit kernel arguments and fill with value zero. + size_t padSize = (~kernargs.size() + 1) & (HIP_IMPLICIT_KERNARG_ALIGNMENT - 1); + kernargs.insert(kernargs.end(), padSize + HIP_IMPLICIT_KERNARG_SIZE, 0); /* Kernel argument preparation. @@ -230,7 +232,7 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, hc::completion_future cf; - lp.av->dispatch_hsa_kernel(&aql, config[1] /* kernarg*/, kernArgSize, + lp.av->dispatch_hsa_kernel(&aql, kernargs.data(), kernargs.size(), (startEvent || stopEvent) ? &cf : nullptr #if (__hcc_workweek__ > 17312) , From 498fe92734396730a66f3bfd4e6ad4aa5a2a4b27 Mon Sep 17 00:00:00 2001 From: Michael LIAO Date: Fri, 24 May 2019 10:22:42 -0400 Subject: [PATCH 15/19] [hip] Minor fix to silence compilation warnings. - Add parenthese to silence repeative compilation warnings across projects built against hip. --- include/hip/hcc_detail/device_functions.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index 808ed216e4..d275bb1820 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -1007,7 +1007,7 @@ void __syncthreads() SIZE 15:11 Range: 1..32 */ -#define GETREG_IMMED(SZ,OFF,REG) (SZ << 11) | (OFF << 6) | REG +#define GETREG_IMMED(SZ,OFF,REG) (((SZ) << 11) | ((OFF) << 6) | (REG)) /* __smid returns the wave's assigned Compute Unit and Shader Engine. From 754d745d84525d5dafdcce6f1f82dba55dc8a449 Mon Sep 17 00:00:00 2001 From: Derek Stinson Date: Tue, 4 Jun 2019 14:19:39 -0400 Subject: [PATCH 16/19] fixed targets flag and TARGET GPU --target-isa= didn't work any longer. --- docs/markdown/hip_kernel_language.md | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/docs/markdown/hip_kernel_language.md b/docs/markdown/hip_kernel_language.md index 5479813675..9e2d0409f7 100644 --- a/docs/markdown/hip_kernel_language.md +++ b/docs/markdown/hip_kernel_language.md @@ -793,10 +793,13 @@ hipcc now supports compiling C++/HIP kernels to binary code objects. The user can specify the target for which the binary can be generated. HIP/HCC does not yet support fat binaries so only a single target may be specified. The file format for binary is `.co` which means Code Object. The following command builds the code object using `hipcc`. -`hipcc --genco --target-isa=[TARGET GPU] [INPUT FILE] -o [OUTPUT FILE]` -```[TARGET GPU] = gfx803/gfx701 +`hipcc --genco --targets [TARGET GPU] [INPUT FILE] -o [OUTPUT FILE]` + +``` +[TARGET GPU] = gfx900 gfx803 gfx701 [INPUT FILE] = Name of the file containing kernels -[OUTPUT FILE] = Name of the generated code object file``` +[OUTPUT FILE] = Name of the generated code object file +``` Note that one important fact to remember when using binary code objects is that the number of arguments to the kernel are different on HCC and NVCC path. Refer to the sample in samples/0_Intro/module_api for differences in the arguments to be passed to the kernel. From 9bb4ecfcfe95c9cddf2bb1700b53a2c9832c805a Mon Sep 17 00:00:00 2001 From: Michael LIAO Date: Tue, 4 Jun 2019 15:15:26 -0400 Subject: [PATCH 17/19] [hip] Make vector type's scalar conversion explicit. --- include/hip/hcc_detail/hip_vector_types.h | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/include/hip/hcc_detail/hip_vector_types.h b/include/hip/hcc_detail/hip_vector_types.h index d19a0e0e3f..8f693427d9 100644 --- a/include/hip/hcc_detail/hip_vector_types.h +++ b/include/hip/hcc_detail/hip_vector_types.h @@ -114,7 +114,7 @@ THE SOFTWARE. typename U, typename std::enable_if< std::is_convertible{}>::type* = nullptr> - inline __host__ __device__ + explicit inline __host__ __device__ HIP_vector_type(U x) noexcept { for (auto i = 0u; i != rank; ++i) data[i] = x; @@ -668,20 +668,20 @@ __MAKE_VECTOR_TYPE__(double, double); #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \ static inline __device__ __host__ \ - type make_##type(comp x) { type r = {x}; return r; } + type make_##type(comp x) { type r{x}; return r; } #define DECLOP_MAKE_TWO_COMPONENT(comp, type) \ static inline __device__ __host__ \ - type make_##type(comp x, comp y) { type r = {x, y}; return r; } + type make_##type(comp x, comp y) { type r{x, y}; return r; } #define DECLOP_MAKE_THREE_COMPONENT(comp, type) \ static inline __device__ __host__ \ - type make_##type(comp x, comp y, comp z) { type r = {x, y, z}; return r; } + type make_##type(comp x, comp y, comp z) { type r{x, y, z}; return r; } #define DECLOP_MAKE_FOUR_COMPONENT(comp, type) \ static inline __device__ __host__ \ type make_##type(comp x, comp y, comp z, comp w) { \ - type r = {x, y, z, w}; \ + type r{x, y, z, w}; \ return r; \ } From 52a8f2fed4c4bda5ad96ec09cf82f3bbd3b960f1 Mon Sep 17 00:00:00 2001 From: Michael LIAO Date: Wed, 5 Jun 2019 08:57:09 -0400 Subject: [PATCH 18/19] [hip] Replace implicit conversions with explicit ones. --- include/hip/hcc_detail/hip_vector_types.h | 24 +++++++++++------------ 1 file changed, 12 insertions(+), 12 deletions(-) diff --git a/include/hip/hcc_detail/hip_vector_types.h b/include/hip/hcc_detail/hip_vector_types.h index 8f693427d9..49df41f5a7 100644 --- a/include/hip/hcc_detail/hip_vector_types.h +++ b/include/hip/hcc_detail/hip_vector_types.h @@ -286,14 +286,14 @@ THE SOFTWARE. HIP_vector_type operator+( const HIP_vector_type& x, U y) noexcept { - return HIP_vector_type{x} += y; + return HIP_vector_type{x} += HIP_vector_type{y}; } template inline __host__ __device__ HIP_vector_type operator+( U x, const HIP_vector_type& y) noexcept { - return y + x; + return HIP_vector_type{x} += y; } template @@ -308,7 +308,7 @@ THE SOFTWARE. HIP_vector_type operator-( const HIP_vector_type& x, U y) noexcept { - return HIP_vector_type{x} -= y; + return HIP_vector_type{x} -= HIP_vector_type{y}; } template inline __host__ __device__ @@ -330,14 +330,14 @@ THE SOFTWARE. HIP_vector_type operator*( const HIP_vector_type& x, U y) noexcept { - return HIP_vector_type{x} *= y; + return HIP_vector_type{x} *= HIP_vector_type{y}; } template inline __host__ __device__ HIP_vector_type operator*( U x, const HIP_vector_type& y) noexcept { - return y * x; + return HIP_vector_type{x} *= y; } template @@ -352,7 +352,7 @@ THE SOFTWARE. HIP_vector_type operator/( const HIP_vector_type& x, U y) noexcept { - return HIP_vector_type{x} /= y; + return HIP_vector_type{x} /= HIP_vector_type{y}; } template inline __host__ __device__ @@ -423,7 +423,7 @@ THE SOFTWARE. HIP_vector_type operator%( const HIP_vector_type& x, U y) noexcept { - return HIP_vector_type{x} %= y; + return HIP_vector_type{x} %= HIP_vector_type{y}; } template< typename T, @@ -456,7 +456,7 @@ THE SOFTWARE. HIP_vector_type operator^( const HIP_vector_type& x, U y) noexcept { - return HIP_vector_type{x} ^= y; + return HIP_vector_type{x} ^= HIP_vector_type{y}; } template< typename T, @@ -489,7 +489,7 @@ THE SOFTWARE. HIP_vector_type operator|( const HIP_vector_type& x, U y) noexcept { - return HIP_vector_type{x} |= y; + return HIP_vector_type{x} |= HIP_vector_type{y}; } template< typename T, @@ -522,7 +522,7 @@ THE SOFTWARE. HIP_vector_type operator&( const HIP_vector_type& x, U y) noexcept { - return HIP_vector_type{x} &= y; + return HIP_vector_type{x} &= HIP_vector_type{y}; } template< typename T, @@ -555,7 +555,7 @@ THE SOFTWARE. HIP_vector_type operator>>( const HIP_vector_type& x, U y) noexcept { - return HIP_vector_type{x} >>= y; + return HIP_vector_type{x} >>= HIP_vector_type{y}; } template< typename T, @@ -588,7 +588,7 @@ THE SOFTWARE. HIP_vector_type operator<<( const HIP_vector_type& x, U y) noexcept { - return HIP_vector_type{x} <<= y; + return HIP_vector_type{x} <<= HIP_vector_type{y}; } template< typename T, From bc528b1e8bffe797ff9dba27f8a514b7adbfaee2 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 13 Jun 2019 23:32:44 +0530 Subject: [PATCH 19/19] HACK for SWDEV-173477/SWDEV-190701 --- src/hip_module.cpp | 16 ++++++++++++++-- 1 file changed, 14 insertions(+), 2 deletions(-) diff --git a/src/hip_module.cpp b/src/hip_module.cpp index 3f6585e059..d69be7f689 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -428,8 +428,20 @@ namespace hip_impl { std::tie(*dptr, *bytes) = read_global_description(it0->second.cbegin(), it0->second.cend(), name); - - return *dptr ? hipSuccess : hipErrorNotFound; + // HACK for SWDEV-173477 + // + // For code objects with global symbols of length 0, ROCR runtime's fix + // may not be working correctly. Therefore the + // result from read_agent_globals() can't be trusted entirely. + // + // As a workaround to tame applications which depend on the existence of + // global symbols with length 0, always return hipSuccess here. + // + // This behavior shall be reverted once ROCR runtime has been fixed to + // address SWDEV-173477 and SWDEV-190701 + + //return *dptr ? hipSuccess : hipErrorNotFound; + return hipSuccess; } hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,