AI/rocm-systems
2
0
Forkuj 0
Kod Zgłoszenia Oczekujące zmiany Actions Packages Projekty Wydania Wiki Aktywność

Merge 'master' into 'amd-master'

Przeglądaj źródła 
Change-Id: Ibca6232ef00edf0c11a5db5f6f10a5bcbb963625


[ROCm/hip commit: 43c928aaa2]
This commit is contained in:
Jenkins
2019-05-13 05:46:17 -04:00
rodzic 6a0eccb154 c4567ad01a
commit 7dbaf7cbd8
144 zmienionych plików z 1376 dodań i 1060 usunięć
Ściągnij plik aktualizacji Ściągnij plik porównania Expand all files Collapse all files
projects/hip/CMakeLists.txt
+1
Wyświetl plik
@@ -232,6 +232,7 @@ if(HIP_PLATFORM STREQUAL "hcc")
set(SOURCE_FILES_RUNTIME
src/code_object_bundle.cpp
src/program_state.cpp
src/hip_clang.cpp
src/hip_hcc.cpp
src/hip_context.cpp
projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp
+14 -51
Wyświetl plik
@@ -59,7 +59,7 @@ template <
typename std::enable_if<n == sizeof...(Ts)>::type* = nullptr>
inline std::vector<std::uint8_t> make_kernarg(
const std::tuple<Ts...>&,
const std::vector<std::pair<std::size_t, std::size_t>>&,
const kernargs_size_align&,
std::vector<std::uint8_t> kernarg) {
return kernarg;
}
@@ -70,7 +70,7 @@ template <
typename std::enable_if<n != sizeof...(Ts)>::type* = nullptr>
inline std::vector<std::uint8_t> make_kernarg(
const std::tuple<Ts...>& formals,
const std::vector<std::pair<std::size_t, std::size_t>>& size_align,
const kernargs_size_align& size_align,
std::vector<std::uint8_t> kernarg) {
using T = typename std::tuple_element<n, std::tuple<Ts...>>::type;
@@ -86,13 +86,12 @@ inline std::vector<std::uint8_t> make_kernarg(
#endif
kernarg.resize(round_up_to_next_multiple_nonnegative(
kernarg.size(), size_align[n].second) + size_align[n].first);
kernarg.size(), size_align.alignment(n)) + size_align.size(n));
std::memcpy(
kernarg.data() + kernarg.size() - size_align[n].first,
kernarg.data() + kernarg.size() - size_align.size(n),
&std::get<n>(formals),
size_align[n].first);
size_align.size(n));
return make_kernarg<n + 1>(formals, size_align, std::move(kernarg));
}
@@ -104,45 +103,17 @@ inline std::vector<std::uint8_t> make_kernarg(
if (sizeof...(Formals) == 0) return {};
auto it = function_names().find(reinterpret_cast<std::uintptr_t>(kernel));
if (it == function_names().cend()) {
hip_throw(std::runtime_error{"Undefined __global__ function."});
}
auto it1 = kernargs().find(it->second);
if (it1 == kernargs().end()) {
hip_throw(std::runtime_error{
"Missing metadata for __global__ function: " + it->second});
}
std::tuple<Formals...> to_formals{std::move(actuals)};
std::vector<std::uint8_t> kernarg;
kernarg.reserve(sizeof(to_formals));
return make_kernarg<0>(to_formals, it1->second, std::move(kernarg));
auto& ps = hip_impl::get_program_state();
return make_kernarg<0>(to_formals,
ps.get_kernargs_size_align(
reinterpret_cast<std::uintptr_t>(kernel)),
std::move(kernarg));
}
inline
std::string name(std::uintptr_t function_address)
{
const auto it = function_names().find(function_address);
if (it == function_names().cend()) {
hip_throw(std::runtime_error{
"Invalid function passed to hipLaunchKernelGGL."});
}
return it->second;
}
inline
std::string name(hsa_agent_t agent)
{
char n[64]{};
hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, n);
return std::string{n};
}
hsa_agent_t target_agent(hipStream_t stream);
@@ -156,17 +127,10 @@ void hipLaunchKernelGGLImpl(
hipStream_t stream,
void** kernarg) {
auto agent = target_agent(stream);
auto it = functions(agent).find(function_address);
const auto& kd = hip_impl::get_program_state().kernel_descriptor(function_address,
target_agent(stream));
if (it == functions(agent).cend()) {
hip_throw(std::runtime_error{
"No device code available for function: " +
name(function_address) +
", for agent: " + name(agent)});
}
hipModuleLaunchKernel(it->second, numBlocks.x, numBlocks.y, numBlocks.z,
hipModuleLaunchKernel(kd, numBlocks.x, numBlocks.y, numBlocks.z,
dimBlocks.x, dimBlocks.y, dimBlocks.z, sharedMemBytes,
stream, nullptr, kernarg);
}
@@ -178,8 +142,7 @@ void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
std::uint32_t sharedMemBytes, hipStream_t stream,
Args... args) {
hip_impl::hip_init();
auto kernarg = hip_impl::make_kernarg(
kernel, std::tuple<Args...>{std::move(args)...});
auto kernarg = hip_impl::make_kernarg(kernel, std::tuple<Args...>{std::move(args)...});
std::size_t kernarg_size = kernarg.size();
void* config[]{
projects/hip/include/hip/hcc_detail/hip_runtime_api.h
+117 -76
Wyświetl plik
@@ -67,6 +67,7 @@ THE SOFTWARE.
#define HIP_LAUNCH_PARAM_END ((void*)0x03)
#ifdef __cplusplus
#include <algorithm>
#include <mutex>
#include <string>
#include <unordered_map>
@@ -190,7 +191,7 @@ enum hipLimit_t {
0x2 ///< Map the allocation into the address space for the current device. The device pointer
///< can be obtained with #hipHostGetDevicePointer.
#define hipHostRegisterIoMemory 0x4 ///< Not supported.
#define hipExtHostRegisterCoarseGrained 0x8 ///< Coarse Grained host memory lock
#define hipDeviceScheduleAuto 0x0 ///< Automatically select between Spin and Yield
#define hipDeviceScheduleSpin \
@@ -2643,90 +2644,129 @@ std::vector<Agent_global> 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<Agent_global>>> globals_from_module;
std::unordered_map<
hsa_agent_t,
std::pair<
std::once_flag,
std::vector<Agent_global>>> 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<std::mutex> 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);
// HACK for SWDEV-173477
//
// For code objects with global symbols of length 0, ROCR runtime would
// ignore them even though they exist in the symbol table. 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
//return *dptr ? hipSuccess : hipErrorNotFound;
return hipSuccess;
}
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<Agent_global> 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;
}
};
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")))
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
static std::unordered_map<
std::string, std::vector<Agent_global>> agent_globals;
std::string key(hash_for(hmod));
if (agent_globals.count(key) == 0) {
static std::mutex mtx;
std::lock_guard<std::mutex> lck{mtx};
if (agent_globals.count(key) == 0) {
agent_globals.emplace(
key, read_agent_globals(this_agent(), executable_for(hmod)));
}
}
const auto it0 = agent_globals.find(key);
if (it0 == agent_globals.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);
// HACK for SWDEV-173477
//
// For code objects with global symbols of length 0, ROCR runtime would
// ignore them even though they exist in the symbol table. 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
//return *dptr ? hipSuccess : hipErrorNotFound;
return hipSuccess;
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) {
static std::unordered_map<hsa_agent_t, std::pair<std::once_flag,
std::vector<Agent_global>>> globals;
static std::once_flag f;
auto agent = this_agent();
// Create placeholder for each agent in the map.
std::call_once(f, []() {
for (auto&& x : hip_impl::all_hsa_agents()) {
(void)globals[x];
}
});
if (globals.find(agent) == globals.cend()) {
hip_throw(std::runtime_error{"invalid agent"});
}
std::call_once(globals[agent].first, [](hsa_agent_t aa) {
std::vector<Agent_global> tmp0;
for (auto&& executable : executables(aa)) {
auto tmp1 = read_agent_globals(aa, executable);
tmp0.insert(tmp0.end(), make_move_iterator(tmp1.begin()),
make_move_iterator(tmp1.end()));
}
globals[aa].second = move(tmp0);
}, agent);
const auto it = globals.find(agent);
if (it == globals.cend()) return hipErrorNotInitialized;
std::tie(*dptr, *bytes) = read_global_description(it->second.second.cbegin(),
it->second.second.cend(), name);
return *dptr ? hipSuccess : hipErrorNotFound;
return get_agent_globals().read_agent_global_from_process(dptr, bytes, name);
}
} // Namespace hip_impl.
#if defined(__cplusplus)
@@ -2748,6 +2788,7 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes,
#endif // __HIP_VDI__
hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const char* name);
/**
* @brief builds module from code object which resides in host memory. Image is pointer to that
* location.
projects/hip/include/hip/hcc_detail/program_state.hpp
+31 -594
Wyświetl plik
@@ -22,38 +22,17 @@ THE SOFTWARE.
#pragma once
#include "code_object_bundle.hpp"
#include "hsa_helpers.hpp"
#if !defined(__cpp_exceptions)
#define try if (true)
#define catch(...) if (false)
#endif
#include "elfio/elfio.hpp"
#if !defined(__cpp_exceptions)
#undef try
#undef catch
#endif
#include <hsa/amd_hsa_kernel_code.h>
#include <hsa/hsa.h>
#include <hsa/hsa_ext_amd.h>
#include <hsa/hsa_ven_amd_loader.h>
#include <link.h>
#include <cassert>
#include <cstddef>
#include <cstdint>
#include <cstdlib>
#include <istream>
#include <memory>
#include <mutex>
#include <sstream>
#include <stdexcept>
#include <string>
#include <unordered_map>
#include <utility>
#include <vector>
struct ihipModuleSymbol_t;
@@ -84,590 +63,48 @@ inline constexpr bool operator==(hsa_isa_t x, hsa_isa_t y) {
namespace hip_impl {
std::vector<hsa_agent_t> all_hsa_agents();
class Kernel_descriptor {
std::uint64_t kernel_object_{};
amd_kernel_code_t const* kernel_header_{nullptr};
std::string name_{};
public:
Kernel_descriptor() = default;
Kernel_descriptor(std::uint64_t kernel_object, const std::string& name)
: kernel_object_{kernel_object}, name_{name}
{
bool supported{false};
std::uint16_t min_v{UINT16_MAX};
auto r = hsa_system_major_extension_supported(
HSA_EXTENSION_AMD_LOADER, 1, &min_v, &supported);
if (r != HSA_STATUS_SUCCESS || !supported) return;
hsa_ven_amd_loader_1_01_pfn_t tbl{};
r = hsa_system_get_major_extension_table(
HSA_EXTENSION_AMD_LOADER,
1,
sizeof(tbl),
reinterpret_cast<void*>(&tbl));
if (r != HSA_STATUS_SUCCESS) return;
if (!tbl.hsa_ven_amd_loader_query_host_address) return;
r = tbl.hsa_ven_amd_loader_query_host_address(
reinterpret_cast<void*>(kernel_object_),
reinterpret_cast<const void**>(&kernel_header_));
if (r != HSA_STATUS_SUCCESS) return;
}
Kernel_descriptor(const Kernel_descriptor&) = default;
Kernel_descriptor(Kernel_descriptor&&) = default;
~Kernel_descriptor() = default;
Kernel_descriptor& operator=(const Kernel_descriptor&) = default;
Kernel_descriptor& operator=(Kernel_descriptor&&) = default;
operator hipFunction_t() const { // TODO: this is awful and only meant for illustration.
return reinterpret_cast<hipFunction_t>(const_cast<Kernel_descriptor*>(this));
}
};
template<typename P>
inline
ELFIO::section* find_section_if(ELFIO::elfio& reader, P p) {
const auto it = std::find_if(
reader.sections.begin(), reader.sections.end(), std::move(p));
return it != reader.sections.end() ? *it : nullptr;
}
inline
__attribute__((visibility("hidden")))
const std::unordered_map<
hsa_isa_t, std::vector<std::vector<char>>>& code_object_blobs() {
static std::unordered_map<hsa_isa_t, std::vector<std::vector<char>>> r;
static std::once_flag f;
std::call_once(f, []() {
static std::vector<std::vector<char>> blobs{};
dl_iterate_phdr([](dl_phdr_info* info, std::size_t, void*) {
ELFIO::elfio tmp;
const auto elf =
info->dlpi_addr ? info->dlpi_name : "/proc/self/exe";
if (!tmp.load(elf)) return 0;
const auto it = find_section_if(tmp, [](const ELFIO::section* x) {
return x->get_name() == ".kernel";
});
if (!it) return 0;
blobs.emplace_back(it->get_data(), it->get_data() + it->get_size());
return 0;
}, nullptr);
for (auto&& multi_arch_blob : blobs) {
auto it = multi_arch_blob.begin();
while (it != multi_arch_blob.end()) {
Bundled_code_header tmp{it, multi_arch_blob.end()};
if (!valid(tmp)) break;
for (auto&& bundle : bundles(tmp)) {
r[triple_to_hsa_isa(bundle.triple)].push_back(bundle.blob);
}
it += tmp.bundled_code_size;
};
}
});
return r;
}
struct Symbol {
std::string name;
ELFIO::Elf64_Addr value = 0;
ELFIO::Elf_Xword size = 0;
ELFIO::Elf_Half sect_idx = 0;
std::uint8_t bind = 0;
std::uint8_t type = 0;
std::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;
}
inline
__attribute__((visibility("hidden")))
const std::unordered_map<
std::string,
std::pair<ELFIO::Elf64_Addr, ELFIO::Elf_Xword>>& symbol_addresses() {
static std::unordered_map<
std::string, std::pair<ELFIO::Elf64_Addr, ELFIO::Elf_Xword>> r;
static std::once_flag f;
std::call_once(f, []() {
dl_iterate_phdr([](dl_phdr_info* info, std::size_t, void*) {
ELFIO::elfio tmp;
const auto elf =
info->dlpi_addr ? info->dlpi_name : "/proc/self/exe";
if (!tmp.load(elf)) return 0;
auto it = find_section_if(tmp, [](const ELFIO::section* x) {
return x->get_type() == SHT_SYMTAB;
});
if (!it) return 0;
const ELFIO::symbol_section_accessor symtab{tmp, it};
for (auto i = 0u; i != symtab.get_symbols_num(); ++i) {
auto s = read_symbol(symtab, i);
if (s.type != STT_OBJECT || s.sect_idx == SHN_UNDEF) continue;
const auto addr = s.value + info->dlpi_addr;
r.emplace(std::move(s.name), std::make_pair(addr, s.size));
}
return 0;
}, nullptr);
});
return r;
}
inline
__attribute__((visibility("hidden")))
std::unordered_map<std::string, void*>& globals() {
static std::unordered_map<std::string, void*> r;
static std::once_flag f;
std::call_once(f, []() { r.reserve(symbol_addresses().size()); });
return r;
}
inline
std::vector<std::string> copy_names_of_undefined_symbols(
const ELFIO::symbol_section_accessor& section) {
std::vector<std::string> r;
for (auto i = 0u; i != section.get_symbols_num(); ++i) {
// TODO: this is boyscout code, caching the temporaries
// may be of worth.
auto tmp = read_symbol(section, i);
if (tmp.sect_idx != SHN_UNDEF || tmp.name.empty()) continue;
r.push_back(std::move(tmp.name));
}
return r;
}
[[noreturn]]
void hip_throw(const std::exception&);
inline
void associate_code_object_symbols_with_host_allocation(
const ELFIO::elfio& reader,
ELFIO::section* code_object_dynsym,
hsa_agent_t agent,
hsa_executable_t executable) {
if (!code_object_dynsym) return;
class kernargs_size_align;
class program_state_impl;
class program_state {
public:
program_state();
~program_state();
const auto undefined_symbols = copy_names_of_undefined_symbols(
ELFIO::symbol_section_accessor{reader, code_object_dynsym});
hipFunction_t kernel_descriptor(std::uintptr_t,
hsa_agent_t);
kernargs_size_align get_kernargs_size_align(std::uintptr_t);
hsa_executable_t load_executable(const char*, const size_t,
hsa_executable_t,
hsa_agent_t);
for (auto&& x : undefined_symbols) {
if (globals().find(x) != globals().cend()) return;
void* global_addr_by_name(const char* name);
const auto it1 = symbol_addresses().find(x);
// to fix later
const std::vector<hsa_executable_t>& executables(hsa_agent_t agent);
if (it1 == symbol_addresses().cend()) {
hip_throw(std::runtime_error{
"Global symbol: " + x + " is undefined."});
}
program_state(const program_state&) = delete;
static std::mutex mtx;
std::lock_guard<std::mutex> lck{mtx};
private:
program_state_impl* impl;
};
if (globals().find(x) != globals().cend()) return;
globals().emplace(x, (void*)(it1->second.first));
void* p = nullptr;
hsa_amd_memory_lock(
reinterpret_cast<void*>(it1->second.first),
it1->second.second,
nullptr, // All agents.
0,
&p);
hsa_executable_agent_global_variable_define(
executable, agent, x.c_str(), p);
}
}
inline
void load_code_object_and_freeze_executable(
const std::string& file, hsa_agent_t agent, hsa_executable_t executable) {
// TODO: the following sequence is inefficient, should be refactored
// into a single load of the file and subsequent ELFIO
// processing.
static const auto cor_deleter = [](hsa_code_object_reader_t* p) {
if (!p) return;
hsa_code_object_reader_destroy(*p);
delete p;
};
using RAII_code_reader =
std::unique_ptr<hsa_code_object_reader_t, decltype(cor_deleter)>;
if (file.empty()) return;
RAII_code_reader tmp{new hsa_code_object_reader_t, cor_deleter};
hsa_code_object_reader_create_from_memory(
file.data(), file.size(), tmp.get());
hsa_executable_load_agent_code_object(
executable, agent, *tmp, nullptr, nullptr);
hsa_executable_freeze(executable, nullptr);
static std::vector<RAII_code_reader> code_readers;
static std::mutex mtx;
std::lock_guard<std::mutex> lck{mtx};
code_readers.push_back(move(tmp));
}
inline
hsa_executable_t load_executable(const std::string& file,
hsa_executable_t executable,
hsa_agent_t agent) {
ELFIO::elfio reader;
std::stringstream tmp{file};
if (!reader.load(tmp)) return hsa_executable_t{};
const auto code_object_dynsym = find_section_if(
reader, [](const ELFIO::section* x) {
return x->get_type() == SHT_DYNSYM;
});
associate_code_object_symbols_with_host_allocation(reader,
code_object_dynsym,
agent, executable);
load_code_object_and_freeze_executable(file, agent, executable);
return executable;
}
class kernargs_size_align {
public:
std::size_t size(std::size_t n) const;
std::size_t alignment(std::size_t n) const;
private:
const void* handle;
friend kernargs_size_align program_state::get_kernargs_size_align(std::uintptr_t);
};
inline
__attribute__((visibility("hidden")))
const std::vector<hsa_executable_t>& executables(hsa_agent_t agent) {
static std::unordered_map<hsa_agent_t, std::pair<std::once_flag,
std::vector<hsa_executable_t>>> r;
static std::once_flag f;
// Create placeholder for each agent in the map.
std::call_once(f, []() {
for (auto&& x : hip_impl::all_hsa_agents()) {
(void)r[x];
}
});
if (r.find(agent) == r.cend()) {
hip_throw(std::runtime_error{"invalid agent"});
}
std::call_once(r[agent].first, [](hsa_agent_t aa) {
hsa_agent_iterate_isas(aa, [](hsa_isa_t x, void* pa) {
const auto it = code_object_blobs().find(x);
if (it == code_object_blobs().cend()) return HSA_STATUS_SUCCESS;
hsa_agent_t a = *static_cast<hsa_agent_t*>(pa);
for (auto&& blob : it->second) {
hsa_executable_t tmp = {};
hsa_executable_create_alt(
HSA_PROFILE_FULL,
HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT,
nullptr,
&tmp);
// TODO: this is massively inefficient and only meant for
// illustration.
std::string blob_to_str{blob.cbegin(), blob.cend()};
tmp = load_executable(blob_to_str, tmp, a);
if (tmp.handle) r[a].second.push_back(tmp);
}
return HSA_STATUS_SUCCESS;
}, &aa);
}, agent);
return r[agent].second;
}
inline
std::vector<std::pair<std::uintptr_t, std::string>> function_names_for(
const ELFIO::elfio& reader, ELFIO::section* symtab) {
std::vector<std::pair<std::uintptr_t, std::string>> r;
ELFIO::symbol_section_accessor symbols{reader, symtab};
for (auto i = 0u; i != symbols.get_symbols_num(); ++i) {
// TODO: this is boyscout code, caching the temporaries
// may be of worth.
auto tmp = read_symbol(symbols, i);
if (tmp.type != STT_FUNC) continue;
if (tmp.type == SHN_UNDEF) continue;
if (tmp.name.empty()) continue;
r.emplace_back(tmp.value, tmp.name);
}
return r;
}
inline
__attribute__((visibility("hidden")))
const std::unordered_map<std::uintptr_t, std::string>& function_names() {
static std::unordered_map<std::uintptr_t, std::string> r;
static std::once_flag f;
std::call_once(f, []() {
dl_iterate_phdr([](dl_phdr_info* info, std::size_t, void*) {
ELFIO::elfio tmp;
const auto elf =
info->dlpi_addr ? info->dlpi_name : "/proc/self/exe";
if (!tmp.load(elf)) return 0;
const auto it = find_section_if(tmp, [](const ELFIO::section* x) {
return x->get_type() == SHT_SYMTAB;
});
if (!it) return 0;
auto names = function_names_for(tmp, it);
for (auto&& x : names) x.first += info->dlpi_addr;
r.insert(
std::make_move_iterator(names.begin()),
std::make_move_iterator(names.end()));
return 0;
}, nullptr);
});
return r;
}
inline
__attribute__((visibility("hidden")))
const std::unordered_map<
std::string, std::vector<hsa_executable_symbol_t>>& kernels(hsa_agent_t agent) {
static std::unordered_map<hsa_agent_t, std::pair<std::once_flag,
std::unordered_map<std::string, std::vector<hsa_executable_symbol_t>>>> r;
static std::once_flag f;
// Create placeholder for each agent in the map.
std::call_once(f, []() {
for (auto&& x : hip_impl::all_hsa_agents()) {
(void)r[x];
}
});
if (r.find(agent) == r.cend()) {
hip_throw(std::runtime_error{"invalid agent"});
}
std::call_once(r[agent].first, [](hsa_agent_t aa) {
static const auto copy_kernels = [](
hsa_executable_t, hsa_agent_t a, hsa_executable_symbol_t x, void*) {
if (type(x) == HSA_SYMBOL_KIND_KERNEL) r[a].second[name(x)].push_back(x);
return HSA_STATUS_SUCCESS;
};
for (auto&& executable : executables(aa)) {
hsa_executable_iterate_agent_symbols(
executable, aa, copy_kernels, nullptr);
}
}, agent);
return r[agent].second;
}
inline
__attribute__((visibility("hidden")))
const std::unordered_map<
std::uintptr_t,
Kernel_descriptor>& functions(hsa_agent_t agent) {
static std::unordered_map<hsa_agent_t, std::pair<std::once_flag,
std::unordered_map<std::uintptr_t, Kernel_descriptor>>> r;
static std::once_flag f;
// Create placeholder for each agent in the map.
std::call_once(f, []() {
for (auto&& x : hip_impl::all_hsa_agents()) {
(void)r[x];
}
});
if (r.find(agent) == r.cend()) {
hip_throw(std::runtime_error{"invalid agent"});
}
std::call_once(r[agent].first, [](hsa_agent_t aa) {
for (auto&& function : function_names()) {
const auto it = kernels(aa).find(function.second);
if (it == kernels(aa).cend()) continue;
for (auto&& kernel_symbol : it->second) {
r[aa].second.emplace(
function.first,
Kernel_descriptor{kernel_object(kernel_symbol), it->first});
}
}
}, agent);
return r[agent].second;
}
inline
std::size_t parse_args(
const std::string& metadata,
std::size_t f,
std::size_t l,
std::vector<std::pair<std::size_t, std::size_t>>& size_align) {
if (f == l) return f;
if (!size_align.empty()) return l;
do {
static constexpr size_t size_sz{5};
f = metadata.find("Size:", f) + size_sz;
if (l <= f) return f;
auto size = std::strtoul(&metadata[f], nullptr, 10);
static constexpr size_t align_sz{6};
f = metadata.find("Align:", f) + align_sz;
char* l{};
auto align = std::strtoul(&metadata[f], &l, 10);
f += (l - &metadata[f]) + 1;
size_align.emplace_back(size, align);
} while (true);
}
inline
void read_kernarg_metadata(
ELFIO::elfio& reader,
std::unordered_map<
std::string,
std::vector<std::pair<std::size_t, std::size_t>>>& kernargs) {
// TODO: this is inefficient.
auto it = find_section_if(reader, [](const ELFIO::section* x) {
return x->get_type() == SHT_NOTE;
});
if (!it) 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{};
acc.get_note(i, type, name, desc, desc_size);
if (name != "AMD") continue; // TODO: switch to using NT_AMD_AMDGPU_HSA_METADATA.
std::string tmp{
static_cast<char*>(desc), static_cast<char*>(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);
}
}
inline
__attribute__((visibility("hidden")))
const std::unordered_map<
std::string, std::vector<std::pair<std::size_t, std::size_t>>>& kernargs() {
static std::unordered_map<
std::string, std::vector<std::pair<std::size_t, std::size_t>>> r;
static std::once_flag f;
std::call_once(f, []() {
for (auto&& isa_blobs : code_object_blobs()) {
for (auto&& blob : isa_blobs.second) {
std::stringstream tmp{std::string{blob.cbegin(), blob.cend()}};
ELFIO::elfio reader;
if (!reader.load(tmp)) continue;
read_kernarg_metadata(reader, r);
}
}
});
return r;
program_state& get_program_state() {
static program_state ps;
return ps;
}
} // Namespace hip_impl.
projects/hip/samples/2_Cookbook/0_MatrixTranspose/Readme.md
+26 -22
Wyświetl plik
@@ -7,30 +7,32 @@ This tutorial shows how to get write simple HIP application. We will write the s
HIP is a C++ runtime API and kernel language that allows developers to create portable applications that can run on AMD and other GPU’s. Our goal was to rise above the lowest-common-denominator paths and deliver a solution that allows you, the developer, to use essential hardware features and maximize your application’s performance on GPU hardware.
## Requirement:
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
## prerequiste knowledge:
Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
## Simple Matrix Transpose
## Simple Matrix Transpose
Here is simple example showing how to write your first program in HIP.
In order to use the HIP framework, we need to add the "hip_runtime.h" header file. SInce its c++ api you can add any header file you have been using earlier while writing your c/c++ program. For gpgpu programming, we have host(microprocessor) and the device(gpu).
In order to use the HIP framework, we need to add the "hip_runtime.h" header file. SInce its c++ api you can add any header file you have been using earlier while writing your c/c++ program. For gpgpu programming, we have host(microprocessor) and the device(gpu).
## Device-side code
We will work on device side code first, Here is simple example showing a snippet of HIP device side code:
`__global__ void matrixTranspose(float *out, `
` float *in, `
` const int width, `
` const int height) `
`{ `
` int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; `
` int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; `
` `
` out[y * width + x] = in[x * height + y]; `
`} `
```
__global__ void matrixTranspose(float *out,
float *in,
const int width,
const int height)
{
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
out[y * width + x] = in[x * height + y];
}
```
`__global__` keyword is the Function-Type Qualifiers, it is used with functions that are executed on device and are called/launched from the hosts.
other function-type qualifiers are:
@@ -40,11 +42,11 @@ other function-type qualifiers are:
`__host__` can combine with `__device__`, in which case the function compiles for both the host and device. These functions cannot use the HIP grid coordinate functions (for example, "hipThreadIdx_x", will talk about it latter). A possible workaround is to pass the necessary coordinate info as an argument to the function.
`__host__` cannot combine with `__global__`.
`__global__` functions are often referred to as *kernels*, and calling one is termed *launching the kernel*.
`__global__` functions are often referred to as *kernels*, and calling one is termed *launching the kernel*.
Next keyword is `void`. HIP `__global__` functions must have a `void` return type. Global functions require the caller to specify an "execution configuration" that includes the grid and block dimensions. The execution configuration can also include other information for the launch, such as the amount of additional shared memory to allocate and the stream where the kernel should execute.
The kernel function begins with
The kernel function begins with
` int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;`
` int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;`
here the keyword hipBlockIdx_x, hipBlockIdx_y and hipBlockIdx_z(not used here) are the built-in functions to identify the threads in a block. The keyword hipBlockDim_x, hipBlockDim_y and hipBlockDim_z(not used here) are to identify the dimensions of the block.
@@ -60,18 +62,20 @@ We allocated memory to the Matrix on host side by using malloc and initiallized
here the first parameter is the destination pointer, second is the source pointer, third is the size of memory copy and the last specify the direction on memory copy(which is in this case froom host to device). While in order to transfer memory from device to host, use `hipMemcpyDeviceToHost` and for device to device memory copy use `hipMemcpyDeviceToDevice`.
Now, we'll see how to launch the kernel.
` hipLaunchKernelGGL(matrixTranspose, `
` dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y), `
` dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), `
` 0, 0, `
` gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT); `
```
hipLaunchKernelGGL(matrixTranspose,
dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
0, 0,
gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT);
```
HIP introduces a standard C++ calling convention to pass the execution configuration to the kernel (this convention replaces the `Cuda <<< >>>` syntax). In HIP,
- Kernels launch with the `"hipLaunchKernelGGL"` function
- The first five parameters to hipLaunchKernelGGL are the following:
- **symbol kernelName**: the name of the kernel to launch. To support template kernels which contains "," use the HIP_KERNEL_NAME macro. In current application it's "matrixTranspose".
- **dim3 gridDim**: 3D-grid dimensions specifying the number of blocks to launch. In MatrixTranspose sample, it's "dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y)".
- **dim3 blockDim**: 3D-block dimensions specifying the number of threads in each block.In MatrixTranspose sample, it's "dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y)".
- **dim3 gridDim**: 3D-grid dimensions specifying the number of blocks to launch. In MatrixTranspose sample, it's "dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y)".
- **dim3 blockDim**: 3D-block dimensions specifying the number of threads in each block.In MatrixTranspose sample, it's "dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y)".
- **size_t dynamicShared**: amount of additional shared memory to allocate when launching the kernel. In MatrixTranspose sample, it's '0'.
- **hipStream_t**: stream where the kernel should execute. A value of 0 corresponds to the NULL stream.In MatrixTranspose sample, it's '0'.
- Kernel arguments follow these first five parameters. Here, these are "gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT".
projects/hip/samples/2_Cookbook/10_inline_asm/Readme.md
+3 -3
Wyświetl plik
@@ -6,7 +6,7 @@ This tutorial is about how to use inline GCN asm in kernel. In this tutorial, we
If you want to take advantage of the extra performance benefits of writing in assembly as well as take advantage of special GPU hardware features that were only available through assemby, then this tutorial is for you. In this tutorial we'll be explaining how to start writing inline asm in kernel.
For more insight Please read the following blogs by Ben Sander
For more insight Please read the following blogs by Ben Sander
[The Art of AMDGCN Assembly: How to Bend the Machine to Your Will](gpuopen.com/amdgcn-assembly)
[AMD GCN Assembly: Cross-Lane Operations](http://gpuopen.com/amd-gcn-assembly-cross-lane-operations/)
@@ -15,13 +15,13 @@ For more information:
[User Guide for AMDGPU Back-end](llvm.org/docs/AMDGPUUsage.html)
## Requirement:
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
## prerequiste knowledge:
Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
## Simple Matrix Transpose
## Simple Matrix Transpose
We will be using the Simple Matrix Transpose application from the our very first tutorial.
projects/hip/samples/2_Cookbook/1_hipEvent/Readme.md
+16 -10
Wyświetl plik
@@ -7,13 +7,13 @@ This tutorial is follow-up of the previous one where we learn how to write our f
Memory transfer and kernel execution are the most important parameter in parallel computing (specially HPC and machine learning). Memory bottlenecks is the main problem why we are not able to get the highest performance, therefore obtaining the memory transfer timing and kernel execution timing plays key role in application optimization.
## Requirement:
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
## prerequiste knowledge:
Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
## Simple Matrix Transpose
## Simple Matrix Transpose
We will be using the Simple Matrix Transpose application from the previous tutorial and modify it to learn how to get the performance score for memory transfer and kernel execution time.
@@ -21,12 +21,16 @@ We will be using the Simple Matrix Transpose application from the previous tutor
We'll learn how to use the event management functionality of HIP runtime api. In the same sourcecode, we used for MatrixTranspose we will declare the following events as follows:
` hipEvent_t start, stop;`
```
hipEvent_t start, stop;
```
We'll create the event with the help of following code:
` hipEventCreate(&start);`
` hipEventCreate(&stop);`
```
hipEventCreate(&start);
hipEventCreate(&stop);
```
We'll use the "eventMs" variable to store the time taken value:
` float eventMs = 1.0f;`
@@ -41,11 +45,13 @@ Now, we'll have the operation for which we need to compute the time taken. For t
` hipMemcpy(gpuMatrix, Matrix, NUM*sizeof(float), hipMemcpyHostToDevice);`
and for kernel execution time we'll use `hipKernelLaunch`:
` hipLaunchKernelGGL(matrixTranspose, `
` dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y), `
` dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), `
` 0, 0, `
` gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT); `
```
hipLaunchKernelGGL(matrixTranspose,
dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
0, 0,
gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT);
```
Now to mark the end of the eventRecord, we will again use the hipEventRecord by passing the stop event:
` hipEventRecord(stop, NULL);`
projects/hip/samples/2_Cookbook/3_shared_memory/Readme.md
+2 -2
Wyświetl plik
@@ -7,13 +7,13 @@ Earlier we learned how to write our first hip program, in which we compute Matri
As we mentioned earlier that Memory bottlenecks is the main problem why we are not able to get the highest performance, therefore minimizing the latency for memory access plays prominent role in application optimization. In this tutorial, we'll learn how to use static shared memory and will explain the dynamic one latter.
## Requirement:
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
## prerequiste knowledge:
Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
## Simple Matrix Transpose
## Simple Matrix Transpose
We will be using the Simple Matrix Transpose application from the previous tutorial and modify it to learn how to use shared memory.
projects/hip/samples/2_Cookbook/4_shfl/Readme.md
+12 -10
Wyświetl plik
@@ -5,23 +5,25 @@ In this tutorial, we'll explain how to use the warp shfl operations to improve t
## Introduction:
Let's talk about Warp first. The kernel code is executed in groups of fixed number of threads known as Warp. For nvidia WarpSize is 32 while for AMD, 32 for Polaris architecture and 64 for rest. Threads in a warp are referred to as lanes and are numbered from 0 to warpSize -1. With the help of shfl ops, we can directly exchange values of variable between threads without using any memory ops within a warp. There are four types of shfl ops:
` int __shfl (int var, int srcLane, int width=warpSize); `
` float __shfl (float var, int srcLane, int width=warpSize); `
` int __shfl_up (int var, unsigned int delta, int width=warpSize); `
` float __shfl_up (float var, unsigned int delta, int width=warpSize); `
` int __shfl_down (int var, unsigned int delta, int width=warpSize); `
` float __shfl_down (float var, unsigned int delta, int width=warpSize); `
` int __shfl_xor (int var, int laneMask, int width=warpSize) `
` float __shfl_xor (float var, int laneMask, int width=warpSize); `
```
int __shfl (int var, int srcLane, int width=warpSize);
float __shfl (float var, int srcLane, int width=warpSize);
int __shfl_up (int var, unsigned int delta, int width=warpSize);
float __shfl_up (float var, unsigned int delta, int width=warpSize);
int __shfl_down (int var, unsigned int delta, int width=warpSize);
float __shfl_down (float var, unsigned int delta, int width=warpSize);
int __shfl_xor (int var, int laneMask, int width=warpSize)
float __shfl_xor (float var, int laneMask, int width=warpSize);
```
## Requirement:
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
## prerequiste knowledge:
Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
## Simple Matrix Transpose
## Simple Matrix Transpose
We will be using the Simple Matrix Transpose application from the previous tutorial and modify it to learn how to use shared memory.
projects/hip/samples/2_Cookbook/5_2dshfl/Readme.md
+16 -12
Wyświetl plik
@@ -5,31 +5,35 @@ This tutorial is follow-up of the previous tutorial, where we learned how to use
## Introduction:
Let's talk about Warp first. The kernel code is executed in groups of fixed number of threads known as Warp. For nvidia WarpSize is 32 while for AMD, 32 for Polaris architecture and 64 for rest. Threads in a warp are referred to as lanes and are numbered from 0 to warpSize -1. With the help of shfl ops, we can directly exchange values of variable between threads without using any memory ops within a warp. There are four types of shfl ops:
` int __shfl (int var, int srcLane, int width=warpSize); `
` float __shfl (float var, int srcLane, int width=warpSize); `
` int __shfl_up (int var, unsigned int delta, int width=warpSize); `
` float __shfl_up (float var, unsigned int delta, int width=warpSize); `
` int __shfl_down (int var, unsigned int delta, int width=warpSize); `
` float __shfl_down (float var, unsigned int delta, int width=warpSize); `
` int __shfl_xor (int var, int laneMask, int width=warpSize) `
` float __shfl_xor (float var, int laneMask, int width=warpSize); `
```
int __shfl (int var, int srcLane, int width=warpSize);
float __shfl (float var, int srcLane, int width=warpSize);
int __shfl_up (int var, unsigned int delta, int width=warpSize);
float __shfl_up (float var, unsigned int delta, int width=warpSize);
int __shfl_down (int var, unsigned int delta, int width=warpSize);
float __shfl_down (float var, unsigned int delta, int width=warpSize);
int __shfl_xor (int var, int laneMask, int width=warpSize);
float __shfl_xor (float var, int laneMask, int width=warpSize);
```
## Requirement:
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
## prerequiste knowledge:
Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
## Simple Matrix Transpose
## Simple Matrix Transpose
We will be using the Simple Matrix Transpose application from the previous tutorial and modify it to learn how to use shared memory.
## __shfl ops in 2D
In the same sourcecode, we used for MatrixTranspose. We'll add the following:
` int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; `
` out[x*width + y] = __shfl(val,y*width + x); `
```
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
out[x*width + y] = __shfl(val,y*width + x);
```
With the help of this application, we can say that kernel code can be converted into multi-dimensional threads with ease.
projects/hip/samples/2_Cookbook/6_dynamic_shared/Readme.md
+5 -3
Wyświetl plik
@@ -7,13 +7,13 @@ Earlier we learned how to use static shared memory. In this tutorial, we'll expl
As we mentioned earlier that Memory bottlenecks is the main problem why we are not able to get the highest performance, therefore minimizing the latency for memory access plays prominent role in application optimization. In this tutorial, we'll learn how to use dynamic shared memory.
## Requirement:
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
## prerequiste knowledge:
Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
## Simple Matrix Transpose
## Simple Matrix Transpose
We will be using the Simple Matrix Transpose application from the previous tutorial and modify it to learn how to use shared memory.
@@ -25,11 +25,13 @@ Shared memory is way more faster than that of global and constant memory and acc
here the first parameter is the data type while the second one is the variable name.
The other important change is:
` hipLaunchKernelGGL(matrixTranspose, `
```
hipLaunchKernelGGL(matrixTranspose,
dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
sizeof(float)*WIDTH*WIDTH, 0,
gpuTransposeMatrix , gpuMatrix, WIDTH);
```
here we replaced 4th parameter with amount of additional shared memory to allocate when launching the kernel.
## How to build and run:
projects/hip/samples/2_Cookbook/7_streams/Readme.md
+12 -6
Wyświetl plik
@@ -7,13 +7,13 @@ In all Earlier tutorial we used single stream, In this tutorial, we'll explain h
The various instances of kernel to be executed on device in exact launch order defined by Host are called streams. We can launch multiple streams on a single device. We will learn how to learn two streams which can we scaled with ease.
## Requirement:
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
## prerequiste knowledge:
Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
## Simple Matrix Transpose
## Simple Matrix Transpose
We will be using the Simple Matrix Transpose application from the previous tutorial and modify it to learn how to launch multiple streams.
@@ -23,22 +23,28 @@ In this tutorial, we'll use both instances of shared memory (i.e., static and dy
` hipStream_t streams[num_streams]; `
and create stream using `hipStreamCreate` as follows:
` for(int i=0;i<num_streams;i++) `
` hipStreamCreate(&streams[i]); `
```
for(int i=0;i<num_streams;i++)
hipStreamCreate(&streams[i]);
```
and while kernel launch, we make the following changes in 5th parameter to hipLaunchKernelGGL(having 0 as the default stream value):
` hipLaunchKernelGGL(matrixTranspose_static_shared, `
```
hipLaunchKernelGGL(matrixTranspose_static_shared,
dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
0, streams[0],
gpuTransposeMatrix[0], data[0], width);
```
` hipLaunchKernelGGL(matrixTranspose_dynamic_shared, `
```
hipLaunchKernelGGL(matrixTranspose_dynamic_shared,
dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
sizeof(float)*WIDTH*WIDTH, streams[1],
gpuTransposeMatrix[1], data[1], width);
```
here we replaced 4th parameter with amount of additional shared memory to allocate when launching the kernel.
projects/hip/samples/2_Cookbook/9_unroll/Readme.md
+11 -9
Wyświetl plik
@@ -1,6 +1,6 @@
## Using Pragma unroll ###
In this tutorial, we'll explain how to use #pragma unroll to improve the performance.
In this tutorial, we'll explain how to use #pragma unroll to improve the performance.
## Introduction:
@@ -8,24 +8,26 @@ Loop unrolling optimization hints can be specified with #pragma unroll and #prag
Specifying #pragma unroll without a parameter directs the loop unroller to attempt to fully unroll the loop if the trip count is known at compile time and attempt to partially unroll the loop if the trip count is not known at compile time.
## Requirement:
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
## prerequiste knowledge:
Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
## Simple Matrix Transpose
## Simple Matrix Transpose
For this tutorial we will be using MatrixTranspose with shfl operation i.e., our 4_shfl tutorial since it is the only examples where we used loops inside the kernel.
In this tutorial, we'll use `#pragma unroll`. In the same sourcecode, we used for MatrixTranspose. We'll add it just before the for loop as following:
`#pragma unroll `
` for(int i=0;i<width;i++) `
` { `
` for(int j=0;j<width;j++) `
` out[i*width + j] = __shfl(val,j*width + i); `
` } `
```
#pragma unroll
for(int i=0;i<width;i++)
{
for(int j=0;j<width;j++)
out[i*width + j] = __shfl(val,j*width + i);
}
```
Specifying the optional parameter, #pragma unroll value, directs the unroller to unroll the loop value times. Be careful while using it.
Specifying #pragma nounroll indicates that the loop should not be unroll. #pragma unroll 1 will show the same behaviour.
projects/hip/src/hip_clang.cpp
+3 -1
Wyświetl plik
@@ -85,7 +85,9 @@ __hipRegisterFatBinary(const void* data)
reinterpret_cast<uintptr_t>(header) + desc->offset), desc->size};
if (HIP_DUMP_CODE_OBJECT)
__hipDumpCodeObject(image);
module->executable = hip_impl::load_executable(image, module->executable, agent);
module->executable = hip_impl::get_program_state().load_executable(image.data(), image.size(),
module->executable,
agent);
if (module->executable.handle) {
modules->at(deviceId) = module;
projects/hip/src/hip_hcc.cpp
+10
Wyświetl plik
@@ -2489,4 +2489,14 @@ namespace hip_impl {
std::terminate();
#endif
}
std::mutex executables_cache_mutex;
std::vector<hsa_executable_t>& executables_cache(
std::string elf, hsa_isa_t isa, hsa_agent_t agent) {
static std::unordered_map<std::string,
std::unordered_map<hsa_isa_t,
std::unordered_map<hsa_agent_t, std::vector<hsa_executable_t>>>> cache;
return cache[elf][isa][agent];
}
} // Namespace hip_impl.
projects/hip/src/hip_memory.cpp
+12 -2
Wyświetl plik
@@ -957,15 +957,25 @@ hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags)
}
// TODO-test : multi-gpu access to registered host memory.
if (ctx) {
if (flags == hipHostRegisterDefault || flags == hipHostRegisterPortable ||
flags == hipHostRegisterMapped) {
if ((flags == hipHostRegisterDefault) || (flags & hipHostRegisterPortable) ||
(flags & hipHostRegisterMapped) || (flags == hipExtHostRegisterCoarseGrained)) {
auto device = ctx->getWriteableDevice();
std::vector<hc::accelerator> vecAcc;
for (int i = 0; i < g_deviceCnt; i++) {
vecAcc.push_back(ihipGetDevice(i)->_acc);
}
#if (__hcc_workweek__ >= 19183)
if(flags & hipExtHostRegisterCoarseGrained) {
am_status = hc::am_memory_host_lock(device->_acc, hostPtr, sizeBytes, &vecAcc[0],
vecAcc.size());
} else {
am_status = hc::am_memory_host_lock_with_flag(device->_acc, hostPtr, sizeBytes, &vecAcc[0],
vecAcc.size());
}
#else
am_status = hc::am_memory_host_lock(device->_acc, hostPtr, sizeBytes, &vecAcc[0],
vecAcc.size());
#endif
if ( am_status == AM_SUCCESS ) {
am_status = hc::am_memtracker_getinfo(&amPointerInfo, hostPtr);
projects/hip/src/hip_module.cpp
+10 -11
Wyświetl plik
@@ -25,6 +25,7 @@ THE SOFTWARE.
#include "hip/hcc_detail/hsa_helpers.hpp"
#include "hip/hcc_detail/program_state.hpp"
#include "hip_hcc_internal.h"
#include "program_state.inl"
#include "trace_helper.h"
#include <hsa/amd_hsa_kernel_code.h>
@@ -289,7 +290,7 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes,
if (!name) return hipErrorNotInitialized;
return hip_impl::read_agent_global_from_module(dptr, bytes, hmod, name);
return hip_impl::get_agent_globals().read_agent_global_from_module(dptr, bytes, hmod, name);
}
namespace hip_impl {
@@ -512,11 +513,8 @@ hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func)
if (!func) return hipErrorInvalidDeviceFunction;
auto agent = this_agent();
const auto it = functions(agent).find(reinterpret_cast<uintptr_t>(func));
if (it == functions(agent).cend()) return hipErrorInvalidDeviceFunction;
const auto header = static_cast<hipFunction_t>(it->second)->_header;
auto kd = get_program_state().kernel_descriptor(reinterpret_cast<uintptr_t>(func), agent);
const auto header = kd->_header;
if (!header) throw runtime_error{"Ill-formed Kernel_descriptor."};
@@ -548,7 +546,8 @@ hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) {
auto content = tmp.empty() ? read_elf_file_as_string(image) : tmp;
(*module)->executable = load_executable(content, (*module)->executable,
(*module)->executable = get_program_state().load_executable(
content.data(), content.size(), (*module)->executable,
this_agent());
// compute the hash of the code object
@@ -591,10 +590,10 @@ hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const
if (!texRef) return ihipLogStatus(hipErrorInvalidValue);
if (!hmod || !name) return ihipLogStatus(hipErrorNotInitialized);
auto addr = get_program_state().global_addr_by_name(name);
if (addr == nullptr) return ihipLogStatus(hipErrorInvalidValue);
const auto it = globals().find(name);
if (it == globals().end()) return ihipLogStatus(hipErrorInvalidValue);
*texRef = reinterpret_cast<textureReference*>(it->second);
*texRef = reinterpret_cast<textureReference*>(addr);
return ihipLogStatus(hipSuccess);
}
projects/hip/src/program_state.cpp
+63
Wyświetl plik
@@ -0,0 +1,63 @@
#include "../include/hip/hcc_detail/program_state.hpp"
#include <hsa/hsa.h>
#include <cstdint>
#include <stdexcept>
#include <unordered_map>
#include <vector>
// 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<const std::vector<std::pair<std::size_t, std::size_t>>*>(handle))[n].first;
}
std::size_t kernargs_size_align::alignment(std::size_t n) const{
return (*reinterpret_cast<const std::vector<std::pair<std::size_t, std::size_t>>*>(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() {
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;
}
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);
}
const std::vector<hsa_executable_t>& 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);
return kd;
}
kernargs_size_align program_state::get_kernargs_size_align(std::uintptr_t kernel) {
kernargs_size_align t;
t.handle = reinterpret_cast<const void*>(&impl->kernargs_size_align(kernel));
return t;
}
};
projects/hip/src/program_state.inl
+713
Wyświetl plik
@@ -0,0 +1,713 @@
#include "../include/hip/hcc_detail/program_state.hpp"
#include "../include/hip/hcc_detail/code_object_bundle.hpp"
#include "../include/hip/hcc_detail/hsa_helpers.hpp"
#if !defined(__cpp_exceptions)
#define try if (true)
#define catch(...) if (false)
#endif
#include "../include/hip/hcc_detail/elfio/elfio.hpp"
#if !defined(__cpp_exceptions)
#undef try
#undef catch
#endif
#include <hsa/amd_hsa_kernel_code.h>
#include <hsa/hsa.h>
#include <hsa/hsa_ext_amd.h>
#include <hsa/hsa_ven_amd_loader.h>
#include <link.h>
#include <cassert>
#include <cstddef>
#include <cstdint>
#include <cstdio>
#include <memory>
#include <mutex>
#include <string>
#include <sstream>
#include <unordered_map>
#include <utility>
#include <vector>
namespace hip_impl {
[[noreturn]]
void hip_throw(const std::exception&);
std::vector<hsa_agent_t> all_hsa_agents();
extern std::mutex executables_cache_mutex;
std::vector<hsa_executable_t>& executables_cache(std::string, hsa_isa_t, hsa_agent_t);
template<typename P>
inline
ELFIO::section* find_section_if(ELFIO::elfio& reader, P p) {
const auto it = std::find_if(
reader.sections.begin(), reader.sections.end(), std::move(p));
return it != reader.sections.end() ? *it : nullptr;
}
struct Symbol {
std::string name;
ELFIO::Elf64_Addr value = 0;
ELFIO::Elf_Xword size = 0;
ELFIO::Elf_Half sect_idx = 0;
std::uint8_t bind = 0;
std::uint8_t type = 0;
std::uint8_t other = 0;
};
class Kernel_descriptor {
std::uint64_t kernel_object_{};
amd_kernel_code_t const* kernel_header_{nullptr};
std::string name_{};
public:
Kernel_descriptor() = default;
Kernel_descriptor(std::uint64_t kernel_object, const std::string& name)
: kernel_object_{kernel_object}, name_{name}
{
bool supported{false};
std::uint16_t min_v{UINT16_MAX};
auto r = hsa_system_major_extension_supported(
HSA_EXTENSION_AMD_LOADER, 1, &min_v, &supported);
if (r != HSA_STATUS_SUCCESS || !supported) return;
hsa_ven_amd_loader_1_01_pfn_t tbl{};
r = hsa_system_get_major_extension_table(
HSA_EXTENSION_AMD_LOADER,
1,
sizeof(tbl),
reinterpret_cast<void*>(&tbl));
if (r != HSA_STATUS_SUCCESS) return;
if (!tbl.hsa_ven_amd_loader_query_host_address) return;
r = tbl.hsa_ven_amd_loader_query_host_address(
reinterpret_cast<void*>(kernel_object_),
reinterpret_cast<const void**>(&kernel_header_));
if (r != HSA_STATUS_SUCCESS) return;
}
Kernel_descriptor(const Kernel_descriptor&) = default;
Kernel_descriptor(Kernel_descriptor&&) = default;
~Kernel_descriptor() = default;
Kernel_descriptor& operator=(const Kernel_descriptor&) = default;
Kernel_descriptor& operator=(Kernel_descriptor&&) = default;
operator hipFunction_t() const { // TODO: this is awful and only meant for illustration.
return reinterpret_cast<hipFunction_t>(const_cast<Kernel_descriptor*>(this));
}
};
class program_state_impl {
public:
std::pair<
std::once_flag,
std::unordered_map<
std::string,
std::unordered_map<
hsa_isa_t,
std::vector<std::vector<char>>>>> code_object_blobs;
std::pair<
std::once_flag,
std::unordered_map<
std::string,
std::pair<ELFIO::Elf64_Addr, ELFIO::Elf_Xword>>> symbol_addresses;
std::unordered_map<
hsa_agent_t,
std::pair<
std::once_flag,
std::vector<hsa_executable_t>>> executables;
std::unordered_map<
hsa_agent_t,
std::pair<
std::once_flag,
std::unordered_map<
std::string,
std::vector<hsa_executable_symbol_t>>>> kernels;
std::pair<
std::once_flag,
std::unordered_map<
std::string, std::vector<std::pair<std::size_t, std::size_t>>>> kernargs;
std::pair<
std::once_flag,
std::unordered_map<std::uintptr_t, std::string>> function_names;
std::unordered_map<
hsa_agent_t,
std::pair<
std::once_flag,
std::unordered_map<
std::uintptr_t,
Kernel_descriptor>>> functions;
std::tuple<
std::once_flag,
std::mutex,
std::unordered_map<std::string, void*>> globals;
using RAII_code_reader =
std::unique_ptr<hsa_code_object_reader_t,
std::function<void(hsa_code_object_reader_t*)>>;
std::pair<
std::mutex,
std::vector<RAII_code_reader>> code_readers;
program_state_impl() {
// Create placeholder for each agent for the per-agent members.
for (auto&& x : hip_impl::all_hsa_agents()) {
(void)executables[x];
(void)kernels[x];
(void)functions[x];
}
}
const std::unordered_map<
std::string,
std::unordered_map<
hsa_isa_t,
std::vector<std::vector<char>>>>& get_code_object_blobs() {
std::call_once(code_object_blobs.first, [this]() {
dl_iterate_phdr([](dl_phdr_info* info, std::size_t, void* p) {
ELFIO::elfio tmp;
const auto elf =
info->dlpi_addr ? info->dlpi_name : "/proc/self/exe";
if (!tmp.load(elf)) return 0;
const auto it = find_section_if(tmp, [](const ELFIO::section* x) {
return x->get_name() == ".kernel";
});
if (!it) return 0;
auto& impl = *static_cast<program_state_impl*>(p);
std::vector<char> multi_arch_blob(it->get_data(), it->get_data() + it->get_size());
auto blob_it = multi_arch_blob.begin();
while (blob_it != multi_arch_blob.end()) {
Bundled_code_header tmp{blob_it, multi_arch_blob.end()};
if (!valid(tmp)) break;
for (auto&& bundle : bundles(tmp)) {
impl.code_object_blobs.second[elf][triple_to_hsa_isa(bundle.triple)].push_back(bundle.blob);
}
blob_it += tmp.bundled_code_size;
};
return 0;
}, this);
});
return code_object_blobs.second;
}
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;
}
const std::unordered_map<
std::string,
std::pair<ELFIO::Elf64_Addr, ELFIO::Elf_Xword>>& get_symbol_addresses() {
std::call_once(symbol_addresses.first, [this]() {
dl_iterate_phdr([](dl_phdr_info* info, std::size_t, void* psi_ptr) {
if (!psi_ptr)
return 0;
program_state_impl* t = static_cast<program_state_impl*>(psi_ptr);
ELFIO::elfio tmp;
const auto elf =
info->dlpi_addr ? info->dlpi_name : "/proc/self/exe";
if (!tmp.load(elf)) return 0;
auto it = find_section_if(tmp, [](const ELFIO::section* x) {
return x->get_type() == SHT_SYMTAB;
});
if (!it) return 0;
const ELFIO::symbol_section_accessor symtab{tmp, it};
for (auto i = 0u; i != symtab.get_symbols_num(); ++i) {
auto s = t->read_symbol(symtab, i);
if (s.type != STT_OBJECT || s.sect_idx == SHN_UNDEF) continue;
const auto addr = s.value + info->dlpi_addr;
t->symbol_addresses.second.emplace(std::move(s.name), std::make_pair(addr, s.size));
}
return 0;
}, this);
});
return symbol_addresses.second;
}
std::unordered_map<std::string, void*>& get_globals() {
std::call_once(std::get<0>(globals), [this]() {
std::get<2>(globals).reserve(get_symbol_addresses().size());
});
return std::get<2>(globals);
}
std::mutex& get_globals_mutex() {
return std::get<1>(globals);
}
std::vector<std::string> copy_names_of_undefined_symbols(
const ELFIO::symbol_section_accessor& section) {
std::vector<std::string> r;
for (auto i = 0u; i != section.get_symbols_num(); ++i) {
// TODO: this is boyscout code, caching the temporaries
// may be of worth.
auto tmp = read_symbol(section, i);
if (tmp.sect_idx != SHN_UNDEF || tmp.name.empty()) continue;
r.push_back(std::move(tmp.name));
}
return r;
}
void associate_code_object_symbols_with_host_allocation(
const ELFIO::elfio& reader,
ELFIO::section* code_object_dynsym,
hsa_agent_t agent,
hsa_executable_t executable) {
if (!code_object_dynsym) return;
const auto undefined_symbols = copy_names_of_undefined_symbols(
ELFIO::symbol_section_accessor{reader, code_object_dynsym});
auto& g = get_globals();
auto& g_mutex = get_globals_mutex();
for (auto&& x : undefined_symbols) {
if (g.find(x) != g.cend()) return;
const auto it1 = get_symbol_addresses().find(x);
if (it1 == get_symbol_addresses().cend()) {
hip_throw(std::runtime_error{
"Global symbol: " + x + " is undefined."});
}
std::lock_guard<std::mutex> lck{g_mutex};
if (g.find(x) != g.cend()) return;
g.emplace(x, (void*)(it1->second.first));
void* p = nullptr;
hsa_amd_memory_lock(
reinterpret_cast<void*>(it1->second.first),
it1->second.second,
nullptr, // All agents.
0,
&p);
hsa_executable_agent_global_variable_define(
executable, agent, x.c_str(), p);
}
}
void load_code_object_and_freeze_executable(
const std::string& file, hsa_agent_t agent, hsa_executable_t executable) {
// TODO: the following sequence is inefficient, should be refactored
// into a single load of the file and subsequent ELFIO
// processing.
if (file.empty()) return;
static const auto cor_deleter = [] (hsa_code_object_reader_t* p) {
if (!p) return;
hsa_code_object_reader_destroy(*p);
delete p;
};
RAII_code_reader tmp{new hsa_code_object_reader_t, cor_deleter};
hsa_code_object_reader_create_from_memory(
file.data(), file.size(), tmp.get());
hsa_executable_load_agent_code_object(
executable, agent, *tmp, nullptr, nullptr);
hsa_executable_freeze(executable, nullptr);
std::lock_guard<std::mutex> lck{code_readers.first};
code_readers.second.push_back(move(tmp));
}
const std::vector<hsa_executable_t>& get_executables(hsa_agent_t agent) {
if (executables.find(agent) == executables.cend()) {
hip_throw(std::runtime_error{"invalid agent"});
}
std::call_once(executables[agent].first, [this](hsa_agent_t aa) {
auto data = std::make_pair(this, &aa);
hsa_agent_iterate_isas(aa, [](hsa_isa_t x, void* d) {
auto& p = *static_cast<decltype(data)*>(d);
auto& impl = *(p.first);
for (const auto code_object_it : impl.get_code_object_blobs()) {
const auto elf = code_object_it.first;
const auto code_object_blobs = code_object_it.second;
const auto it = code_object_blobs.find(x);
if (it == code_object_blobs.cend()) continue;
hsa_agent_t a = *static_cast<hsa_agent_t*>(p.second);
std::lock_guard<std::mutex> lck{executables_cache_mutex};
std::vector<hsa_executable_t>& current_exes =
hip_impl::executables_cache(elf, x, a);
// check the cache for already loaded executables
if (current_exes.empty()) {
// executables do not yet exist for this elf+isa+agent, create and cache them
for (auto&& blob : it->second) {
hsa_executable_t tmp = {};
hsa_executable_create_alt(
HSA_PROFILE_FULL,
HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT,
nullptr,
&tmp);
// TODO: this is massively inefficient and only meant for
// illustration.
tmp = impl.load_executable(blob.data(), blob.size(), tmp, a);
if (tmp.handle) current_exes.push_back(tmp);
}
}
// append cached executables to our agent's vector of executables
impl.executables[a].second.insert(impl.executables[a].second.end(),
current_exes.begin(), current_exes.end());
}
return HSA_STATUS_SUCCESS;
}, &data);
}, agent);
return executables[agent].second;
}
hsa_executable_t load_executable(const char* data,
const size_t data_size,
hsa_executable_t executable,
hsa_agent_t agent) {
ELFIO::elfio reader;
std::string ts = std::string(data, data_size);
std::stringstream tmp{ts};
if (!reader.load(tmp)) return hsa_executable_t{};
const auto code_object_dynsym = find_section_if(
reader, [](const ELFIO::section* x) {
return x->get_type() == SHT_DYNSYM;
});
associate_code_object_symbols_with_host_allocation(reader,
code_object_dynsym,
agent, executable);
load_code_object_and_freeze_executable(ts, agent, executable);
return executable;
}
std::vector<std::pair<std::uintptr_t, std::string>> function_names_for(
const ELFIO::elfio& reader, ELFIO::section* symtab) {
std::vector<std::pair<std::uintptr_t, std::string>> r;
ELFIO::symbol_section_accessor symbols{reader, symtab};
for (auto i = 0u; i != symbols.get_symbols_num(); ++i) {
// TODO: this is boyscout code, caching the temporaries
// may be of worth.
auto tmp = read_symbol(symbols, i);
if (tmp.type != STT_FUNC) continue;
if (tmp.type == SHN_UNDEF) continue;
if (tmp.name.empty()) continue;
r.emplace_back(tmp.value, tmp.name);
}
return r;
}
const std::unordered_map<std::uintptr_t, std::string>& get_function_names() {
std::call_once(function_names.first, [this]() {
dl_iterate_phdr([](dl_phdr_info* info, std::size_t, void* p) {
ELFIO::elfio tmp;
const auto elf =
info->dlpi_addr ? info->dlpi_name : "/proc/self/exe";
if (!tmp.load(elf)) return 0;
const auto it = find_section_if(tmp, [](const ELFIO::section* x) {
return x->get_type() == SHT_SYMTAB;
});
if (!it) return 0;
auto& impl = *static_cast<program_state_impl*>(p);
auto names = impl.function_names_for(tmp, it);
for (auto&& x : names) x.first += info->dlpi_addr;
impl.function_names.second.insert(
std::make_move_iterator(names.begin()),
std::make_move_iterator(names.end()));
return 0;
}, this);
});
return function_names.second;
}
const std::unordered_map<
std::string, std::vector<hsa_executable_symbol_t>>& get_kernels(hsa_agent_t agent) {
if (kernels.find(agent) == kernels.cend()) {
hip_throw(std::runtime_error{"invalid agent"});
}
std::call_once(kernels[agent].first, [this](hsa_agent_t aa) {
static const auto copy_kernels = [](
hsa_executable_t, hsa_agent_t a, hsa_executable_symbol_t x, void* p) {
auto& impl = *static_cast<program_state_impl*>(p);
if (type(x) == HSA_SYMBOL_KIND_KERNEL) impl.kernels[a].second[hip_impl::name(x)].push_back(x);
return HSA_STATUS_SUCCESS;
};
for (auto&& executable : get_executables(aa)) {
hsa_executable_iterate_agent_symbols(
executable, aa, copy_kernels, this);
}
}, agent);
return kernels[agent].second;
}
const std::unordered_map<
std::uintptr_t,
Kernel_descriptor>& get_functions(hsa_agent_t agent) {
if (functions.find(agent) == functions.cend()) {
hip_throw(std::runtime_error{"invalid agent"});
}
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);
if (it == get_kernels(aa).cend()) continue;
for (auto&& kernel_symbol : it->second) {
functions[aa].second.emplace(
function.first,
Kernel_descriptor{kernel_object(kernel_symbol), it->first});
}
}
}, agent);
return functions[agent].second;
}
std::size_t parse_args(
const std::string& metadata,
std::size_t f,
std::size_t l,
std::vector<std::pair<std::size_t, std::size_t>>& size_align) {
if (f == l) return f;
if (!size_align.empty()) return l;
do {
static constexpr size_t size_sz{5};
f = metadata.find("Size:", f) + size_sz;
if (l <= f) return f;
auto size = std::strtoul(&metadata[f], nullptr, 10);
static constexpr size_t align_sz{6};
f = metadata.find("Align:", f) + align_sz;
char* l{};
auto align = std::strtoul(&metadata[f], &l, 10);
f += (l - &metadata[f]) + 1;
size_align.emplace_back(size, align);
} while (true);
}
void read_kernarg_metadata(
ELFIO::elfio& reader,
std::unordered_map<
std::string,
std::vector<std::pair<std::size_t, std::size_t>>>& kernargs) {
// TODO: this is inefficient.
auto it = find_section_if(reader, [](const ELFIO::section* x) {
return x->get_type() == SHT_NOTE;
});
if (!it) 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{};
acc.get_note(i, type, name, desc, desc_size);
if (name != "AMD") continue; // TODO: switch to using NT_AMD_AMDGPU_HSA_METADATA.
std::string tmp{
static_cast<char*>(desc), static_cast<char*>(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);
}
}
const std::unordered_map<std::string,
std::vector<std::pair<std::size_t, std::size_t>>>& get_kernargs() {
std::call_once(kernargs.first, [this]() {
for (auto&& name_and_isa_blobs : get_code_object_blobs()) {
for (auto&& isa_blobs : name_and_isa_blobs.second) {
for (auto&& blob : isa_blobs.second) {
std::stringstream tmp{std::string{blob.cbegin(), blob.cend()}};
ELFIO::elfio reader;
if (!reader.load(tmp)) continue;
read_kernarg_metadata(reader, kernargs.second);
}
}
}
});
return kernargs.second;
}
std::string name(std::uintptr_t function_address)
{
const auto it = get_function_names().find(function_address);
if (it == get_function_names().cend()) {
hip_throw(std::runtime_error{
"Invalid function passed to hipLaunchKernelGGL."});
}
return it->second;
}
std::string name(hsa_agent_t agent)
{
char n[64]{};
hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, n);
return std::string{n};
}
const Kernel_descriptor& kernel_descriptor(std::uintptr_t function_address,
hsa_agent_t agent) {
auto it0 = get_functions(agent).find(function_address);
if (it0 == get_functions(agent).cend()) {
hip_throw(std::runtime_error{
"No device code available for function: " +
std::string(name(function_address)) +
", for agent: " + name(agent)});
}
return it0->second;
}
const std::vector<std::pair<std::size_t, std::size_t>>&
kernargs_size_align(std::uintptr_t kernel) {
auto it = get_function_names().find(kernel);
if (it == get_function_names().cend()) {
hip_throw(std::runtime_error{"Undefined __global__ function."});
}
auto it1 = get_kernargs().find(it->second);
if (it1 == get_kernargs().end()) {
hip_throw(std::runtime_error{
"Missing metadata for __global__ function: " + it->second});
}
return it1->second;
}
}; // class program_state_impl
};
projects/hip/tests/README.md
+27 -25
Wyświetl plik
@@ -31,17 +31,17 @@ The parser looks for a code block similar to the one below.
```
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* RUN: %t
* TEST: %t
* //Small copy
* RUN: %t -N 10 --memsetval 0x42
* TEST: %t -N 10 --memsetval 0x42
* // Oddball size
* RUN: %t -N 10013 --memsetval 0x5a
* TEST: %t -N 10013 --memsetval 0x5a
* // Big copy
* RUN: %t -N 256M --memsetval 0xa6
* TEST: %t -N 256M --memsetval 0xa6
* HIT_END
*/
```
In the above, BUILD commands provide instructions on how to build the test case while RUN commands provide instructions on how to execute the test case.
In the above, BUILD commands provide instructions on how to build the test case while TEST commands provide instructions on how to execute the test case.
#### BUILD command
@@ -57,36 +57,38 @@ NVCC_OPTIONS: All options specified after this delimiter are passed to hipcc on
EXCLUDE_HIP_PLATFORM: This can be used to exclude a test case from HCC, NVCC or both platforms.
#### RUN command
#### BUILD_CMD command
The supported syntax for the RUN command is:
The supported syntax for the BUILD_CMD command is:
```
RUN: %t <arguments_to_test_executable> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all>
BUILD_CMD: <targetname> <build_command> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all>
```
%s: refers to current source file name. Additional source files needed for the test can be specified by name (including relative path).
%t: refers to target executable named derived by removing the extension from the current source file. Alternatively a target executable name can be specified.
%hc: refers to hipcc pointed to by $CMAKE_INSTALL_PREFIX/bin/hipcc.
%cc: refers to system c compiler pointed to by $CC.
%S: refers to path to current source file.
%T: refers to path to current build target.
#### TEST command
The supported syntax for the TEST command is:
```
TEST: %t <arguments_to_test_executable> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all>
```
%t: refers to target executable named derived by removing the extension from the current source file. Alternatively a target executable name can be specified.
EXCLUDE_HIP_PLATFORM: This can be used to exclude a test case from HCC, NVCC or both platforms. Note that if the test has been excluded for a specific platform in the BUILD command, it is automatically excluded from the RUN command as well for the same platform.
EXCLUDE_HIP_PLATFORM: This can be used to exclude a test case from HCC, NVCC or both platforms. Note that if the test has been excluded for a specific platform in the BUILD command, it is automatically excluded from the TEST command as well for the same platform.
#### RUN_NAMED command
#### TEST_NAMED command
When using the RUN command, HIT will squash and append the arguments specified to the test executable name to generate the CMAKE test name. Sometimes we might want to specify a more descriptive name. The RUN_NAMED command is used for that. The supported syntax for the RUN_NAMED command is:
When using the TEST command, HIT will squash and append the arguments specified to the test executable name to generate the CMAKE test name. Sometimes we might want to specify a more descriptive name. The TEST_NAMED command is used for that. The supported syntax for the TEST_NAMED command is:
```
RUN: %t CMAKE_TEST_NAME <arguments_to_test_executable> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all>
TEST: %t CMAKE_TEST_NAME <arguments_to_test_executable> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all>
```
#### CMAKECMD command
The supported syntax for the CMAKECMD command is:
```
CMAKECMD: <cmake_command> <options_to_cmake_command>
```
cmake_command: refers to any of the commands supported by ```cmake -E``` as specified in the [cmake documentation](https://cmake.org/cmake/help/latest/manual/cmake.1.html#command-line-tool-mode). Note that the commands are limited by the version of cmake the user is running.
options_to_cmake_command: refers to the arguments supported by the specific cmake_command. The arguments are parsed by HIT to replace special markers. The markers supported by HIT are:
%S: Refers to the source directory containing the current source file.
%B: Refers to the build directory for the current cmake project i.e. CMAKE_CURRENT_BINARY_DIR.
### Running tests:
```
ctest
@@ -111,7 +113,7 @@ Find the test and commandline that fail:
(From the build directory, perhaps hip/build)
grep -IR hipMemcpy-modes -IR ../tests/
../tests/src/runtimeApi/memory/hipMemcpy.cpp: * RUN_NAMED: %t hipMemcpy-modes --tests 0x1
../tests/src/runtimeApi/memory/hipMemcpy.cpp: * TEST_NAMED: %t hipMemcpy-modes --tests 0x1
# Guidelines for adding new tests
projects/hip/tests/hit/HIT.cmake
+48 -12
Wyświetl plik
@@ -55,8 +55,31 @@ macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_opt
endforeach()
endmacro()
# Helper macro to parse RUN instructions
macro(PARSE_RUN_COMMAND _target _arguments _exclude_platforms)
# Helper macro to parse CUSTOM BUILD instructions
macro(PARSE_CUSTOMBUILD_COMMAND _target _buildcmd _exclude_platforms)
set(${_target})
set(${_buildcmd} " ")
set(${_exclude_platforms})
set(_target_found FALSE)
set(_exclude_platforms_found FALSE)
foreach(arg ${ARGN})
if(NOT _target_found)
set(_target_found TRUE)
set(${_target} ${arg})
elseif("x${arg}" STREQUAL "xEXCLUDE_HIP_PLATFORM")
set(_exclude_platforms_found TRUE)
else()
if(_exclude_platforms_found)
set(${_exclude_platforms} ${arg})
else()
list(APPEND ${_buildcmd} ${arg})
endif()
endif()
endforeach()
endmacro()
# Helper macro to parse TEST instructions
macro(PARSE_TEST_COMMAND _target _arguments _exclude_platforms)
set(${_target})
set(${_arguments} " ")
set(${_exclude_platforms})
@@ -78,8 +101,8 @@ macro(PARSE_RUN_COMMAND _target _arguments _exclude_platforms)
endforeach()
endmacro()
# Helper macro to parse RUN_NAMED instructions
macro(PARSE_RUN_NAMED_COMMAND _target _testname _arguments _exclude_platforms)
# Helper macro to parse TEST_NAMED instructions
macro(PARSE_TEST_NAMED_COMMAND _target _testname _arguments _exclude_platforms)
set(${_target})
set(${_arguments} " ")
set(${_exclude_platforms})
@@ -155,22 +178,35 @@ macro(HIT_ADD_FILES _dir _label _parent)
endif()
endforeach()
# Run cmake commands
execute_process(COMMAND ${HIP_SRC_PATH}/tests/hit/parser --cmakeCMDs ${file}
# Custom build commands
execute_process(COMMAND ${HIP_SRC_PATH}/tests/hit/parser --customBuildCMDs ${file}
OUTPUT_VARIABLE _contents
ERROR_QUIET
WORKING_DIRECTORY ${_dir}
OUTPUT_STRIP_TRAILING_WHITESPACE)
string(REGEX REPLACE "\n" ";" _contents "${_contents}")
string(REGEX REPLACE "%hc" "${HIP_HIPCC_EXECUTABLE}" _contents "${_contents}")
string(REGEX REPLACE "%cc" "${CC}" _contents "${_contents}")
string(REGEX REPLACE "%S" ${_dir} _contents "${_contents}")
string(REGEX REPLACE "%B" ${CMAKE_CURRENT_BINARY_DIR} _contents "${_contents}")
string(REGEX REPLACE "%T" ${_label} _contents "${_contents}")
foreach(_cmd ${_contents})
string(REGEX REPLACE " " ";" _cmd "${_cmd}")
execute_process(COMMAND ${CMAKE_COMMAND} -E ${_cmd})
parse_custombuild_command(_target _buildcmd _exclude_platforms ${_cmd})
string(REGEX REPLACE "/" "." target ${_label}/${_target})
insert_into_map("_exclude" "${target}" "${_exclude_platforms}")
if(_exclude_platforms STREQUAL "all" OR _exclude_platforms STREQUAL ${HIP_PLATFORM})
else()
string(REGEX REPLACE ";" " " _buildcmd "${_buildcmd}")
#string(CONCAT buildscript ${CMAKE_CURRENT_BINARY_DIR}/${target} ".sh")
#file(WRITE ${buildscript} ${_buildcmd})
#add_custom_target(${target} COMMAND ${buildscript})
add_custom_target(${target} COMMAND sh -c "${_buildcmd}")
add_dependencies(${_parent} ${target})
endif()
endforeach()
# Add tests
execute_process(COMMAND ${HIP_SRC_PATH}/tests/hit/parser --runCMDs ${file}
execute_process(COMMAND ${HIP_SRC_PATH}/tests/hit/parser --testCMDs ${file}
OUTPUT_VARIABLE _contents
ERROR_QUIET
WORKING_DIRECTORY ${_dir}
@@ -178,7 +214,7 @@ macro(HIT_ADD_FILES _dir _label _parent)
string(REGEX REPLACE "\n" ";" _contents "${_contents}")
foreach(_cmd ${_contents})
string(REGEX REPLACE " " ";" _cmd "${_cmd}")
parse_run_command(_target _arguments _exclude_platforms ${_cmd})
parse_test_command(_target _arguments _exclude_platforms ${_cmd})
string(REGEX REPLACE "/" "." target ${_label}/${_target})
read_from_map("_exclude" "${target}" _exclude_platforms_from_build)
if(_exclude_platforms STREQUAL "all" OR _exclude_platforms STREQUAL ${HIP_PLATFORM} OR
@@ -189,7 +225,7 @@ macro(HIT_ADD_FILES _dir _label _parent)
endforeach()
# Add named tests
execute_process(COMMAND ${HIP_SRC_PATH}/tests/hit/parser --runNamedCMDs ${file}
execute_process(COMMAND ${HIP_SRC_PATH}/tests/hit/parser --testNamedCMDs ${file}
OUTPUT_VARIABLE _contents
ERROR_QUIET
WORKING_DIRECTORY ${_dir}
@@ -197,7 +233,7 @@ macro(HIT_ADD_FILES _dir _label _parent)
string(REGEX REPLACE "\n" ";" _contents "${_contents}")
foreach(_cmd ${_contents})
string(REGEX REPLACE " " ";" _cmd "${_cmd}")
parse_run_named_command(_target _testname _arguments _exclude_platforms ${_cmd})
parse_test_named_command(_target _testname _arguments _exclude_platforms ${_cmd})
string(REGEX REPLACE "/" "." target ${_label}/${_target})
read_from_map("_exclude" "${target}" _exclude_platforms_from_build)
if(_exclude_platforms STREQUAL "all" OR _exclude_platforms STREQUAL ${HIP_PLATFORM} OR
projects/hip/tests/hit/parser
+46 -38
Wyświetl plik
@@ -4,49 +4,57 @@ use 5.006; use v5.10.1;
use File::Basename;
use File::Spec;
my $patBUILD = "^".quotemeta(" * BUILD:");
my $patTEST = "^".quotemeta(" * TEST:");
my $patTEST_NAMED = "^".quotemeta(" * TEST_NAMED:");
my $patBUILD_CMD = "^".quotemeta(" * BUILD_CMD:");
# Scan input file for HIT information
sub parse_file {
my $file = shift;
(my $exe = $file) =~ s/\.[^.]+$//g;
my (@buildCMDs, @runCMDs, @runNamedCMDs, @cmakeCMDs);
my (@buildCMDs, @testCMDs, @testNamedCMDs, @customBuildCMDs);
if (open (SOURCE, '<:encoding(UTF-8)', "$file")) {
while (<SOURCE>) {
my $line=$_;
# Look for BUILD instructions
if ($line =~ /^ \* BUILD:/) {
if ($line =~ /$patBUILD/) {
$line =~ s/^ \* BUILD: //g; # Remove " * BUILD: "
$line =~ s/%s/$file/g; # Substitute %s -> filename
$line =~ s/%t/$exe/g; # Substitute %t -> targetname
$line =~ s/\R//g; # Remove line endings
push @buildCMDs, $line;
}
# Look for RUN instructions
if ($line =~ /^ \* RUN:/) {
$line =~ s/^ \* RUN: //g; # Remove " * RUN: "
# Look for TEST instructions
if ($line =~ /$patTEST/) {
$line =~ s/^ \* TEST: //g; # Remove " * TEST: "
$line =~ s/%s/$file/g; # Substitute %s -> filename
$line =~ s/%t/$exe/g; # Subsitute %t -> targetname
$line =~ s/\R//g; # Remove line endings
push @runCMDs, $line;
push @testCMDs, $line;
}
# Look for RUN_NAMED instructions
if ($line =~ /^ \* RUN_NAMED:/) {
$line =~ s/^ \* RUN_NAMED: //g; # Remove " * RUN_NAMED: "
# Look for TEST_NAMED instructions
if ($line =~ /$patTEST_NAMED/) {
$line =~ s/^ \* TEST_NAMED: //g;# Remove " * TEST_NAMED: "
$line =~ s/%s/$file/g; # Substitute %s -> filename
$line =~ s/%t/$exe/g; # Subsitute %t -> targetname
$line =~ s/\R//g; # Remove line endings
push @runNamedCMDs, $line;
push @testNamedCMDs, $line;
}
# Look for CMAKECMD instructions
if ($line =~ /^ \* CMAKECMD:/) {
$line =~ s/^ \* CMAKECMD: //g; # Remove " * CMAKECMD: "
# Substitute %S -> srcdir and %B -> builddir happens in cmake
# Look for BUILD_CMD instructions
if ($line =~ /$patBUILD_CMD/) {
$line =~ s/^ \* BUILD_CMD: //g; # Remove " * BUILD_CMD: "
$line =~ s/%s/$file/g; # Substitute %s -> filename
$line =~ s/%t/$exe/g; # Substitute %t -> targetname
# Substitute %hc -> /path/to/hipcc and %cc -> /path/to/cc happens in cmake
# Substitute %S -> src dir and %T -> target build dir happens in cmake
$line =~ s/\R//g; # Remove line endings
push @cmakeCMDs, $line;
push @customBuildCMDs, $line;
}
}
close(SOURCE);
}
return (\@buildCMDs, \@runCMDs, \@runNamedCMDs, \@cmakeCMDs);
return (\@buildCMDs, \@testCMDs, \@testNamedCMDs, \@customBuildCMDs);
}
# Exit if no arguments specified
@@ -58,52 +66,52 @@ if(scalar @ARGV == 0){
# Parse command
my @options = ();
my $retBuildCMDs = 0;
my $retRunCMDs = 0;
my $retRunNamedCMDs = 0;
my $retCmakeCMDs = 0;
my $retTestCMDs = 0;
my $retTestNamedCMDs = 0;
my $retCustomBuildCMDs = 0;
foreach $arg (@ARGV) {
if ($retBuildCMDs or $retRunCMDs or $retRunNamedCMDs or $retCmakeCMDs) {
if ($retBuildCMDs or $retTestCMDs or $retTestNamedCMDs or $retCustomBuildCMDs) {
push (@options, $arg);
}
if ($arg eq '--buildCMDs') {
$retBuildCMDs = 1;
}
if ($arg eq '--runCMDs') {
$retRunCMDs = 1;
if ($arg eq '--testCMDs') {
$retTestCMDs = 1;
}
if ($arg eq '--runNamedCMDs') {
$retRunNamedCMDs = 1;
if ($arg eq '--testNamedCMDs') {
$retTestNamedCMDs = 1;
}
if ($arg eq '--cmakeCMDs') {
$retCmakeCMDs = 1;
if ($arg eq '--customBuildCMDs') {
$retCustomBuildCMDs = 1;
}
}
# Atleast one command needs to be specified
if (($retBuildCMDs eq 0) and ($retRunCMDs eq 0) and ($retRunNamedCMDs eq 0) and ($retCmakeCMDs eq 0)) {
die "Usage: $0 <--buildCMDs|--runCMDs|--runNamedCMDs|--cmakeCMDs> FILENAMEs\n";
if (($retBuildCMDs eq 0) and ($retTestCMDs eq 0) and ($retTestNamedCMDs eq 0) and($retCustomBuildCMDs eq 0)) {
die "Usage: $0 <--buildCMDs|--testCMDs|--testNamedCMDs|--customBuildCMDs> FILENAMEs\n";
}
# Iterate over input files
foreach $file (@options) {
# Convert absolute path to path relative to working directory
my $relfile = File::Spec->abs2rel($file);
my ($buildCMDs, $runCMDs, $runNamedCMDs, $cmakeCMDs) = parse_file("$relfile");
my ($buildCMDs, $testCMDs, $testNamedCMDs, $customBuildCMDs) = parse_file("$relfile");
if ($retBuildCMDs) {
# print "BuildCMDs:\n";
print "$_\n" for @$buildCMDs;
}
if ($retRunCMDs) {
# print "RunCMDs:\n";
print "$_\n" for @$runCMDs;
if ($retTestCMDs) {
# print "TestCMDs:\n";
print "$_\n" for @$testCMDs;
}
if ($retRunNamedCMDs) {
# print "RunNamedCMDs:\n";
print "$_\n" for @$runNamedCMDs;
if ($retTestNamedCMDs) {
# print "TestNamedCMDs:\n";
print "$_\n" for @$testNamedCMDs;
}
if ($retCmakeCMDs) {
# print "CmakeCMDs:\n";
print "$_\n" for @$cmakeCMDs;
if ($retCustomBuildCMDs) {
# print "CustomBuildCMDs:\n";
print "$_\n" for @$customBuildCMDs;
}
}
projects/hip/tests/src/Functional/device/hipFuncDeviceSynchronize.cpp
+1 -1
Wyświetl plik
@@ -24,7 +24,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/Functional/device/hipFuncGetDevice.cpp
+1 -1
Wyświetl plik
@@ -24,7 +24,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/Functional/device/hipFuncSetDevice.cpp
+1 -1
Wyświetl plik
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* RUN: %t EXCLUDE_HIP_PLATFORM
* TEST: %t EXCLUDE_HIP_PLATFORM
* HIT_END
*/
projects/hip/tests/src/Functional/device/hipFuncSetDeviceFlags.cpp
+1 -1
Wyświetl plik
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/compiler/hipClassKernel.cpp
+1 -1
Wyświetl plik
@@ -21,7 +21,7 @@ THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
#include "hipClassKernel.h"
projects/hip/tests/src/context/hipCtx_simple.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/context/hipMemsetD8.cpp
+4 -4
Wyświetl plik
@@ -24,13 +24,13 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* //Small copy
* RUN: %t -N 10 --memsetval 0x42
* TEST: %t -N 10 --memsetval 0x42
* // Oddball size
* RUN: %t -N 10013 --memsetval 0x5a
* TEST: %t -N 10013 --memsetval 0x5a
* // Big copy
* RUN: %t -N 256M --memsetval 0xa6
* TEST: %t -N 256M --memsetval 0xa6
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipAsynchronousStreams.cpp
+1 -1
Wyświetl plik
@@ -20,7 +20,7 @@ THE SOFTWARE.
/*
* HIT_START
* BUILD: %t %s ../test_common.cpp HIPCC_OPTIONS -std=c++11 EXCLUDE_HIP_PLATFORM nvcc
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipDeviceMalloc.cpp
+1 -1
Wyświetl plik
@@ -18,7 +18,7 @@ THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s NVCC_OPTIONS -std=c++11
* RUN: %t EXCLUDE_HIP_PLATFORM nvcc
* TEST: %t EXCLUDE_HIP_PLATFORM nvcc
* HIT_END
*/
#include "test_common.h"
projects/hip/tests/src/deviceLib/hipDeviceMemcpy.cpp
+1 -1
Wyświetl plik
@@ -9,7 +9,7 @@
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipDoublePrecisionMathDevice.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipDoublePrecisionMathHost.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipFloatMath.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipFloatMathPrecise.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipIntegerIntrinsics.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipMathFunctions.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipSimpleAtomicsTest.cpp
+1 -1
Wyświetl plik
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -std=c++11 --gpu-architecture=sm_60
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp
+1 -1
Wyświetl plik
@@ -21,7 +21,7 @@ THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipSinglePrecisionMathDevice.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipSinglePrecisionMathHost.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipStdComplex.cpp
+1 -1
Wyświetl plik
@@ -18,7 +18,7 @@ THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s NVCC_OPTIONS -std=c++11
* RUN: %t EXCLUDE_HIP_PLATFORM nvcc
* TEST: %t EXCLUDE_HIP_PLATFORM nvcc
* HIT_END
*/
#include "test_common.h"
projects/hip/tests/src/deviceLib/hipTestClock.cpp
+1 -1
Wyświetl plik
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipTestDevice.cpp
+1 -1
Wyświetl plik
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipTestDeviceDouble.cpp
+1 -1
Wyświetl plik
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* XXBUILD: %t %s ../test_common.cpp
* XXRUN: %t
* XXTEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipTestDeviceSymbol.cpp
+1 -1
Wyświetl plik
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipTestDotFunctions.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -std=c++11
* RUN: %t EXCLUDE_HIP_PLATFORM nvcc
* TEST: %t EXCLUDE_HIP_PLATFORM nvcc
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipTestFMA.cpp
+1 -1
Wyświetl plik
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipTestHalf.cpp
+1 -1
Wyświetl plik
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipTestIncludeMath.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipTestNativeHalf.cpp
+1 -1
Wyświetl plik
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipTestNew.cpp
+1 -1
Wyświetl plik
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipThreadFence.cpp
+1 -1
Wyświetl plik
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipVectorTypes.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hipVectorTypesDevice.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hip_anyall.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp NVCC_OPTIONS --Wno-deprecated-declarations
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hip_ballot.cpp
+1 -1
Wyświetl plik
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp NVCC_OPTIONS --Wno-deprecated-declarations
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hip_bitextract.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hip_bitinsert.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hip_brev.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hip_clz.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hip_ffs.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hip_mbcnt.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hip_popc.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hip_test_ldg.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp NVCC_OPTIONS --gpu-architecture=sm_35
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hip_threadfence_system.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -std=c++11
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hip_trig.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/dynamicLoading/complex_loading_behavior.cpp
+6 -9
Wyświetl plik
@@ -20,13 +20,11 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/**
* Build the test:
* hipcc complex_loading_behavior.cpp -o libfoo.so -fPIC -lpthread -shared -DTEST_SHARED_LIBRARY
* hipcc complex_loading_behavior.cpp -o complex_loading_behavior -ldl
*
* Run the test:
* ./complex_loading_behavior
/* HIT_START
* BUILD_CMD: libfoo %hc %S/%s -o libfoo.so -fPIC -lpthread -shared -DTEST_SHARED_LIBRARY
* BUILD_CMD: %t %hc %S/%s -o %T/%t -ldl
* TEST: %t
* HIT_END
*/
#if !defined(TEST_SHARED_LIBRARY)
@@ -93,8 +91,7 @@ int launch_local_kernel() {
free(B_h);
free(C_h);
std::cout << "local launch succedded\n";
std::cout << "PASSED!\n";
return 0;
}
projects/hip/tests/src/hipC.c
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s EXCLUDE_HIP_PLATFORM all
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/hipEnvVarDriver.cpp
+1 -1
Wyświetl plik
@@ -17,7 +17,7 @@ OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWA
/* HIT_START
* BUILD: %t %s test_common.cpp NVCC_OPTIONS -std=c++11
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/hipHcc.cpp
+1 -1
Wyświetl plik
@@ -23,7 +23,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s
* RUN: %t EXCLUDE_HIP_PLATFORM all
* TEST: %t EXCLUDE_HIP_PLATFORM all
* HIT_END
*/
projects/hip/tests/src/kernel/hipDynamicShared.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM hcc
* RUN: %t EXCLUDE_HIP_PLATFORM nvcc
* TEST: %t EXCLUDE_HIP_PLATFORM nvcc
* HIT_END
*/
projects/hip/tests/src/kernel/hipDynamicShared2.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM hcc
* RUN: %t EXCLUDE_HIP_PLATFORM nvcc
* TEST: %t EXCLUDE_HIP_PLATFORM nvcc
* HIT_END
*/
projects/hip/tests/src/kernel/hipEmptyKernel.cpp
+1 -1
Wyświetl plik
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/kernel/hipGridLaunch.cpp
+1 -1
Wyświetl plik
@@ -23,7 +23,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/kernel/hipLanguageExtensions.cpp
+1 -1
Wyświetl plik
@@ -23,7 +23,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/kernel/hipLaunchParm.cpp
+1 -1
Wyświetl plik
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -std=c++11
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/kernel/hipLaunchParmFunctor.cpp
+1 -1
Wyświetl plik
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp HIPCC_OPTIONS -O3
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/kernel/hipPrintfKernel.cpp
+1 -1
Wyświetl plik
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/kernel/hipTestConstant.cpp
+1 -1
Wyświetl plik
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/kernel/hipTestGlobalVariable.cpp
+1 -1
Wyświetl plik
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/kernel/hipTestMallocKernel.cpp
+1 -1
Wyświetl plik
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM all
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/kernel/hipTestMemKernel.cpp
+1 -1
Wyświetl plik
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/kernel/inline_asm_vadd.cpp
+1 -1
Wyświetl plik
@@ -17,7 +17,7 @@ OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWA
/* HIT_START
* BUILD: %t %s
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp
+3 -3
Wyświetl plik
@@ -24,9 +24,9 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t EXCLUDE_HIP_PLATFORM hcc
* RUN: %t --memcpyWithPeer EXCLUDE_HIP_PLATFORM hcc
* RUN: %t --mirrorPeers EXCLUDE_HIP_PLATFORM hcc
* TEST: %t EXCLUDE_HIP_PLATFORM hcc
* TEST: %t --memcpyWithPeer EXCLUDE_HIP_PLATFORM hcc
* TEST: %t --mirrorPeers EXCLUDE_HIP_PLATFORM hcc
* HIT_END
*/
projects/hip/tests/src/runtimeApi/device/hipChooseDevice.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/runtimeApi/device/hipDeviceComputeCapability.cpp
+1 -1
Wyświetl plik
@@ -24,7 +24,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/runtimeApi/device/hipDeviceGetByPCIBusId.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/runtimeApi/device/hipDeviceGetName.cpp
+1 -1
Wyświetl plik
@@ -24,7 +24,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/runtimeApi/device/hipDeviceGetPCIBusId.cpp
+1 -1
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/runtimeApi/device/hipDeviceSynchronize.cpp
+1 -1
Wyświetl plik
@@ -24,7 +24,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/runtimeApi/device/hipDeviceTotalMem.cpp
+1 -1
Wyświetl plik
@@ -24,7 +24,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/runtimeApi/device/hipGetDevice.cpp
+1 -1
Wyświetl plik
@@ -24,7 +24,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp
+1 -1
Wyświetl plik
@@ -23,7 +23,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/runtimeApi/device/hipRuntimeGetVersion.cpp
+1 -1
Wyświetl plik
@@ -26,7 +26,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/runtimeApi/device/hipSetCachceConfig.cpp
+1 -1
Wyświetl plik
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/
projects/hip/tests/src/runtimeApi/device/hipSetDevice.cpp
+1 -1
Wyświetl plik
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* RUN: %t EXCLUDE_HIP_PLATFORM
* TEST: %t EXCLUDE_HIP_PLATFORM
* HIT_END
*/
projects/hip/tests/src/runtimeApi/device/hipSetDeviceFlags.cpp
+1 -1
Wyświetl plik
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* RUN: %t
* TEST: %t
* HIT_END
*/

Some files were not shown because too many files have changed in this diff Show More