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

Merge branch 'master' into amd-master-next

Browse Source 
Conflicts:
	CMakeLists.txt
	tests/src/texture/simpleTexture2DLayered.cpp
	tests/src/texture/simpleTexture3D.cpp

Change-Id: I4aa4754d391b5f37ddf15fa0bcfc84d9da020119
This commit is contained in:
Lad, Aditya
2020-03-06 14:10:04 -05:00
parent 4b312caca4 5c036520b1
commit d80edf9541
42 changed files with 1410 additions and 640 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/hip_hcc.cpp
+14 -43
View File
@@ -39,6 +39,7 @@ THE SOFTWARE.
#include <algorithm>
#include <atomic>
#include <mutex>
#include <unordered_set>
#include <hc.hpp>
#include <hc_am.hpp>
@@ -71,7 +72,6 @@ int HIP_API_BLOCKING = 0;
int HIP_PRINT_ENV = 0;
int HIP_TRACE_API = 0;
std::string HIP_TRACE_API_COLOR("green");
int HIP_PROFILE_API = 0;
// TODO - DB_START/STOP need more testing.
std::string HIP_DB_START_API;
@@ -149,12 +149,10 @@ uint64_t recordApiTrace(TlsData *tls, std::string* fullStr, const std::string& a
if ((tid < g_dbStartTriggers.size()) && (apiSeqNum >= g_dbStartTriggers[tid].nextTrigger())) {
printf("info: resume profiling at %lu\n", apiSeqNum);
RESUME_PROFILING;
g_dbStartTriggers.pop_back();
};
if ((tid < g_dbStopTriggers.size()) && (apiSeqNum >= g_dbStopTriggers[tid].nextTrigger())) {
printf("info: stop profiling at %lu\n", apiSeqNum);
STOP_PROFILING;
g_dbStopTriggers.pop_back();
};
@@ -1294,9 +1292,6 @@ void HipReadEnv() {
"executes.");
READ_ENV_S(release, HIP_TRACE_API_COLOR, 0,
"Color to use for HIP_API. None/Red/Green/Yellow/Blue/Magenta/Cyan/White");
READ_ENV_I(release, HIP_PROFILE_API, 0,
"Add HIP API markers to ATP file generated with CodeXL. 0x1=short API name, "
"0x2=full API name including args.");
READ_ENV_S(release, HIP_DB_START_API, 0,
"Comma-separated list of tid.api_seq_num for when to start debug and profiling.");
READ_ENV_S(release, HIP_DB_STOP_API, 0,
@@ -1372,14 +1367,6 @@ void HipReadEnv() {
HIP_DB |= 0x1;
}
if (HIP_PROFILE_API && !COMPILE_HIP_ATP_MARKER) {
fprintf(stderr,
"warning: env var HIP_PROFILE_API=0x%x but COMPILE_HIP_ATP_MARKER=0. (perhaps "
"enable COMPILE_HIP_ATP_MARKER in src code before compiling?)\n",
HIP_PROFILE_API);
HIP_PROFILE_API = 0;
}
if (HIP_DB) {
fprintf(stderr, "HIP_DB=0x%x [%s]\n", HIP_DB, HIP_DB_string(HIP_DB).c_str());
}
@@ -1423,11 +1410,6 @@ void HipReadEnv() {
// This function creates a vector with only the GPU accelerators.
// It is called with C++11 call_once, which provided thread-safety.
void ihipInit() {
#if COMPILE_HIP_ATP_MARKER
amdtInitializeActivityLogger();
amdtScopedMarker("ihipInit", "HIP", NULL);
#endif
HipReadEnv();
@@ -1617,7 +1599,7 @@ hipStream_t ihipSyncAndResolveStream(hipStream_t stream, bool lockAcquired) {
void ihipPrintKernelLaunch(const char* kernelName, const grid_launch_parm* lp,
const hipStream_t stream) {
if ((HIP_TRACE_API & (1 << TRACE_KCMD)) || HIP_PROFILE_API ||
if ((HIP_TRACE_API & (1 << TRACE_KCMD)) ||
(COMPILE_HIP_DB & HIP_TRACE_API)) {
GET_TLS();
std::stringstream os;
@@ -1630,14 +1612,6 @@ void ihipPrintKernelLaunch(const char* kernelName, const grid_launch_parm* lp,
std::string fullStr;
recordApiTrace(tls, &fullStr, os.str());
}
if (HIP_PROFILE_API == 0x1) {
std::string shortAtpString("hipLaunchKernel:");
shortAtpString += kernelName;
MARKER_BEGIN(shortAtpString.c_str(), "HIP");
} else if (HIP_PROFILE_API == 0x2) {
MARKER_BEGIN(os.str().c_str(), "HIP");
}
}
}
@@ -1696,9 +1670,6 @@ void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, grid_launc
tprintf(DB_SYNC, "ihipPostLaunchKernel, unlocking stream\n");
stream->lockclose_postKernelCommand(kernelName, lp.av, unlockPostponed);
if (HIP_PROFILE_API) {
MARKER_END();
}
}
//=================================================================================================
@@ -1796,6 +1767,8 @@ const char* ihipErrorString(hipError_t hip_error) {
return "hipErrorMissingConfiguration";
case hipErrorLaunchFailure:
return "hipErrorLaunchFailure";
case hipErrorCooperativeLaunchTooLarge:
return "hipErrorCooperativeLaunchTooLarge";
case hipErrorPriorLaunchFailure:
return "hipErrorPriorLaunchFailure";
case hipErrorLaunchTimeOut:
@@ -2478,29 +2451,17 @@ bool ihipStream_t::locked_copy2DAsync(void* dst, const void* src, size_t width,
return retStatus;
}
//-------------------------------------------------------------------------------------------------
//-------------------------------------------------------------------------------------------------
// Profiler, really these should live elsewhere:
hipError_t hipProfilerStart() {
HIP_INIT_API(hipProfilerStart);
#if COMPILE_HIP_ATP_MARKER
amdtResumeProfiling(AMDT_ALL_PROFILING);
#endif
return ihipLogStatus(hipSuccess);
};
hipError_t hipProfilerStop() {
HIP_INIT_API(hipProfilerStop);
#if COMPILE_HIP_ATP_MARKER
amdtStopProfiling(AMDT_ALL_PROFILING);
#endif
return ihipLogStatus(hipSuccess);
};
//-------------------------------------------------------------------------------------------------
//-------------------------------------------------------------------------------------------------
// HCC-specific accessor functions:
@@ -2540,6 +2501,16 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** a
// TODO - add a contect sequence number for debug. Print operator<< ctx:0.1 (device.ctx)
namespace hip_impl {
std::unordered_set<std::string>& get_all_gpuarch() {
static std::unordered_set<std::string> r{};
static std::once_flag init;
std::call_once(init, []() {
for (int i=0; i < g_deviceCnt; i++){
r.insert("hcc-amdgcn-amd-amdhsa--gfx"+std::to_string(g_deviceArray[i]->_props.gcnArch));
}});
return r;
}
std::vector<hsa_agent_t> all_hsa_agents() {
std::vector<hsa_agent_t> r{};
std::vector<hc::accelerator> visible_accelerators;
src/hip_hcc_internal.h
+2 -39
View File
@@ -63,7 +63,6 @@ extern int HIP_LAUNCH_BLOCKING;
extern int HIP_API_BLOCKING;
extern int HIP_PRINT_ENV;
extern int HIP_PROFILE_API;
// extern int HIP_TRACE_API;
extern int HIP_ATP;
extern int HIP_DB;
@@ -250,34 +249,6 @@ extern const char* API_COLOR_END;
// Must be enabled at runtime with HIP_TRACE_API
#define COMPILE_HIP_TRACE_API 0x3
// Compile code that generates trace markers for CodeXL ATP at HIP function begin/end.
// ATP is standard CodeXL format that includes timestamps for kernels, HSA RT APIs, and HIP APIs.
#ifndef COMPILE_HIP_ATP_MARKER
#define COMPILE_HIP_ATP_MARKER 0
#endif
// Compile support for trace markers that are displayed on CodeXL GUI at start/stop of each function
// boundary.
// TODO - currently we print the trace message at the beginning. if we waited, we could also
// tls->tidInfo return codes, and any values returned through ptr-to-args (ie the pointers allocated
// by hipMalloc).
#if COMPILE_HIP_ATP_MARKER
#include "CXLActivityLogger.h"
#define MARKER_BEGIN(markerName, group) amdtBeginMarker(markerName, group, nullptr);
#define MARKER_END() amdtEndMarker();
#define RESUME_PROFILING amdtResumeProfiling(AMDT_ALL_PROFILING);
#define STOP_PROFILING amdtStopProfiling(AMDT_ALL_PROFILING);
#else
// Swallow scoped markers:
#define MARKER_BEGIN(markerName, group)
#define MARKER_END()
#define RESUME_PROFILING
#define STOP_PROFILING
#endif
//---
// HIP Trace modes - use with HIP_TRACE_API=...
#define TRACE_ALL 0 // 0x01
@@ -336,22 +307,17 @@ static inline uint64_t getTicks() { return hc::get_system_ticks(); }
//---
extern uint64_t recordApiTrace(TlsData *tls, std::string* fullStr, const std::string& apiStr);
#if COMPILE_HIP_ATP_MARKER || (COMPILE_HIP_TRACE_API & 0x1)
#if (COMPILE_HIP_TRACE_API & 0x1)
#define API_TRACE(forceTrace, ...) \
GET_TLS(); \
uint64_t hipApiStartTick = 0; \
{ \
tls->tidInfo.incApiSeqNum(); \
if (forceTrace || \
(HIP_PROFILE_API || (COMPILE_HIP_DB && (HIP_TRACE_API & (1 << TRACE_ALL))))) { \
(COMPILE_HIP_DB && (HIP_TRACE_API & (1 << TRACE_ALL)))) { \
std::string apiStr = std::string(__func__) + " (" + ToString(__VA_ARGS__) + ')'; \
std::string fullStr; \
hipApiStartTick = recordApiTrace(tls, &fullStr, apiStr); \
if (HIP_PROFILE_API == 0x1) { \
MARKER_BEGIN(__func__, "HIP") \
} else if (HIP_PROFILE_API == 0x2) { \
MARKER_BEGIN(fullStr.c_str(), "HIP"); \
} \
} \
}
@@ -398,9 +364,6 @@ extern uint64_t recordApiTrace(TlsData *tls, std::string* fullStr, const std::st
tls->tidInfo.apiSeqNum(), __func__, localHipStatus, \
ihipErrorString(localHipStatus), ticks, API_COLOR_END); \
} \
if (HIP_PROFILE_API) { \
MARKER_END(); \
} \
localHipStatus; \
})
src/hip_memory.cpp
+175 -160
View File
@@ -37,7 +37,6 @@ __device__ uint32_t __hip_device_page_flag[__HIP_NUM_PAGES];
namespace hip_internal {
namespace {
inline
const char* hsa_to_string(hsa_status_t err) noexcept
{
@@ -140,6 +139,8 @@ namespace {
return r;
}()};
constexpr std::uint32_t is_cpu_owned{UINT32_MAX};
inline
hsa_amd_pointer_info_t info(const void* p)
{
@@ -149,13 +150,14 @@ namespace {
const_cast<void*>(p), &r, nullptr, nullptr, nullptr),
__FILE__, __func__, __LINE__);
r.size = is_large_BAR || (type(r.agentOwner) == HSA_DEVICE_TYPE_CPU) ?
UINT32_MAX : sizeof(hsa_amd_pointer_info_t);
if (type(r.agentOwner) == HSA_DEVICE_TYPE_CPU) r.size = is_cpu_owned;
return r;
}
constexpr size_t staging_sz{4 * 1024 * 1024}; // 2 Pages.
constexpr size_t staging_sz{4 * 1024 * 1024}; // 2 Pages.
constexpr size_t max_h2d_std_memcpy_sz{8 * 1024}; // 8 KiB.
constexpr size_t max_d2h_std_memcpy_sz{64}; // 1 cacheline.
thread_local const std::unique_ptr<void, void (*)(void *)> staging_buffer{
[]() {
@@ -202,7 +204,7 @@ namespace {
} // Unnamed namespace.
inline
void do_copy(void* __restrict dst, const void* __restrict src, std::size_t n,
void do_copy(void* __restrict dst, const void* __restrict src, size_t n,
hsa_agent_t da, hsa_agent_t sa) {
hsa_signal_silent_store_relaxed(copy_signal, 1);
throwing_result_check(
@@ -224,14 +226,20 @@ void do_std_memcpy(
inline
void d2h_copy(void* __restrict dst, const void* __restrict src, size_t n,
hsa_amd_pointer_info_t si) {
// TODO: characterise direct largeBAR reads from agent-allocated memory.
// if (si.size == UINT32_MAX) {
// return do_std_memcpy(dst, src, n);
// }
const auto di{info(dst)};
const auto is_locked{di.type == HSA_EXT_POINTER_TYPE_LOCKED};
if (di.type == HSA_EXT_POINTER_TYPE_LOCKED) {
if (!is_locked && si.size == is_cpu_owned) {
return do_std_memcpy(dst, src, n);
}
if (!is_locked && is_large_BAR && n <= max_d2h_std_memcpy_sz) {
return do_std_memcpy(dst, src, n);
}
if (di.type == HSA_EXT_POINTER_TYPE_HSA) {
return do_copy(dst, src, n, si.agentOwner, si.agentOwner);
}
if (is_locked) {
dst = static_cast<char*>(di.agentBaseAddress) +
(static_cast<char*>(dst) -
static_cast<char*>(di.hostBaseAddress));
@@ -247,7 +255,7 @@ void d2h_copy(void* __restrict dst, const void* __restrict src, size_t n,
throwing_result_check(hsa_amd_memory_lock(dst, n, &si.agentOwner, 1,
const_cast<void**>(&dst)),
__FILE__, __func__, __LINE__);
__FILE__, __func__, __LINE__);
do_copy(dst, src, n, si.agentOwner, si.agentOwner);
}
@@ -256,16 +264,23 @@ void d2h_copy(void* __restrict dst, const void* __restrict src, size_t n,
inline
void h2d_copy(void* __restrict dst, const void* __restrict src, size_t n,
hsa_amd_pointer_info_t di) {
if (di.size == UINT32_MAX) {
const auto si{info(const_cast<void*>(src))};
const auto is_locked{si.type == HSA_EXT_POINTER_TYPE_LOCKED};
if (!is_locked && di.size == is_cpu_owned) {
return do_std_memcpy(dst, src, n);
}
if (!is_locked && is_large_BAR && n <= max_h2d_std_memcpy_sz) {
return do_std_memcpy(dst, src, n);
}
if (si.type == HSA_EXT_POINTER_TYPE_HSA) {
return do_copy(dst, src, n, di.agentOwner, di.agentOwner);
}
const auto si{info(const_cast<void*>(src))};
if (si.type == HSA_EXT_POINTER_TYPE_LOCKED) {
if (is_locked) {
src = static_cast<char*>(si.agentBaseAddress) +
(static_cast<const char*>(src) -
static_cast<char*>(si.hostBaseAddress));
(static_cast<const char*>(src) -
static_cast<char*>(si.hostBaseAddress));
do_copy(dst, src, n, di.agentOwner, di.agentOwner);
}
else if (n <= staging_sz) {
@@ -279,7 +294,7 @@ void h2d_copy(void* __restrict dst, const void* __restrict src, size_t n,
throwing_result_check(hsa_amd_memory_lock(const_cast<void*>(src), n,
&di.agentOwner, 1,
const_cast<void**>(&src)),
__FILE__, __func__, __LINE__);
__FILE__, __func__, __LINE__);
do_copy(dst, src, n, di.agentOwner, di.agentOwner);
}
@@ -288,69 +303,36 @@ void h2d_copy(void* __restrict dst, const void* __restrict src, size_t n,
inline
void generic_copy(void* __restrict dst, const void* __restrict src, size_t n,
hsa_amd_pointer_info_t di, hsa_amd_pointer_info_t si) {
if (di.size == UINT32_MAX && si.size == UINT32_MAX) {
if (di.size == is_cpu_owned && si.size == is_cpu_owned) {
return do_std_memcpy(dst, src, n);
}
if (di.size == is_cpu_owned) return d2h_copy(dst, src, n, si);
if (si.size == is_cpu_owned) return h2d_copy(dst, src, n, di);
std::unique_ptr<void, void (*)(void*)> lck0{
nullptr, [](void* p) { hsa_amd_memory_unlock(p); }};
std::unique_ptr<void, void (*)(void*)> lck1{nullptr, lck0.get_deleter()};
throwing_result_check(hsa_amd_agents_allow_access(1u, &si.agentOwner,
nullptr,
di.agentBaseAddress),
__FILE__, __func__, __LINE__);
switch (si.type) {
case HSA_EXT_POINTER_TYPE_HSA:
if (di.type == HSA_EXT_POINTER_TYPE_HSA) {
hsa_memory_copy(dst, src, n);
return; // TODO: do_copy(dst, src, n, di.agentOwner, si.agentOwner);
}
if (di.type == HSA_EXT_POINTER_TYPE_UNKNOWN ||
di.type == HSA_EXT_POINTER_TYPE_LOCKED) {
return d2h_copy(dst, src, n, si);
}
break;
case HSA_EXT_POINTER_TYPE_LOCKED:
if (di.type == HSA_EXT_POINTER_TYPE_UNKNOWN) {
std::memcpy(dst, si.hostBaseAddress, n);
return;
}
if (di.type == HSA_EXT_POINTER_TYPE_LOCKED) {
std::memcpy(di.hostBaseAddress, si.hostBaseAddress, n);
return;
}
src = si.agentBaseAddress;
si.agentOwner = di.agentOwner;
break;
case HSA_EXT_POINTER_TYPE_UNKNOWN:
if (di.type == HSA_EXT_POINTER_TYPE_UNKNOWN) {
std::memcpy(dst, src, n);
return;
}
if (di.type == HSA_EXT_POINTER_TYPE_LOCKED) {
std::memcpy(di.hostBaseAddress, src, n);
return;
}
return h2d_copy(dst, src, n, di);
default: do_copy(dst, src, n, di.agentOwner, si.agentOwner); break;
}
return do_copy(dst, src, n, di.agentOwner, si.agentOwner);
}
inline
void memcpy_impl(void* __restrict dst, const void* __restrict src, size_t n,
hipMemcpyKind k) noexcept {
hipMemcpyKind k) {
switch (k) {
case hipMemcpyHostToHost: std::memcpy(dst, src, n); break;
case hipMemcpyHostToDevice:
return is_large_BAR ? do_std_memcpy(dst, src, n)
: h2d_copy(dst, src, n, info(dst));
case hipMemcpyDeviceToHost:
// TODO: characterise direct largeBAR reads from agent-allocated memory.
return /*is_large_BAR ? do_std_memcpy(dst, src, n)
: */d2h_copy(dst, src, n, info(src));
case hipMemcpyDeviceToDevice: hsa_memory_copy(dst, src, n); break;
case hipMemcpyHostToDevice: return h2d_copy(dst, src, n, info(dst));
case hipMemcpyDeviceToHost: return d2h_copy(dst, src, n, info(src));
case hipMemcpyDeviceToDevice: {
const auto di{info(dst)};
const auto si{info(src)};
throwing_result_check(hsa_amd_agents_allow_access(1u, &si.agentOwner,
nullptr,
di.agentBaseAddress),
__FILE__, __func__, __LINE__);
return do_copy(dst, src, n, di.agentOwner, si.agentOwner);
}
default: return generic_copy(dst, src, n, info(dst), info(src));
}
}
@@ -1290,7 +1272,7 @@ hipError_t hipMemcpyToSymbolAsync(void* dst, const void* src, size_t count,
if (dst == nullptr) {
return ihipLogStatus(hipErrorInvalidSymbol);
}
if (kind == hipMemcpyDeviceToHost || kind == hipMemcpyHostToHost) {
return ihipLogStatus(hipErrorInvalidMemcpyDirection);
} else if (kind == hipMemcpyDeviceToDevice) {
@@ -1322,7 +1304,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* src, size_t count,
if (src == nullptr || dst == nullptr) {
return ihipLogStatus(hipErrorInvalidSymbol);
}
if (kind == hipMemcpyHostToDevice || kind == hipMemcpyHostToHost) {
return ihipLogStatus(hipErrorInvalidMemcpyDirection);
} else if (kind == hipMemcpyDeviceToDevice) {
@@ -1540,111 +1522,144 @@ hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t
return ihipLogStatus(e);
}
int getByteSizeFromFormat(const hipChannelFormatDesc& desc){
int byteSize =0;
switch (desc.f) {
case hipChannelFormatKindUnsigned:
switch (desc.x) {
case 32:
byteSize = sizeof(uint32_t);
break;
case 16:
byteSize = sizeof(uint16_t);
break;
case 8:
byteSize = sizeof(uint8_t);
break;
default:
byteSize = sizeof(uint32_t);
}
break;
case hipChannelFormatKindSigned:
switch (desc.x) {
case 32:
byteSize = sizeof(int32_t);
break;
case 16:
byteSize = sizeof(int16_t);
break;
case 8:
byteSize = sizeof(int8_t);
break;
default:
byteSize = sizeof(int32_t);
}
break;
case hipChannelFormatKindFloat:
switch (desc.x) {
case 32:
byteSize = sizeof(float);
break;
case 16:
byteSize = sizeof(_Float16);
break;
default:
byteSize = sizeof(float);
}
break;
case hipChannelFormatKindNone:
default:
break;
}
return byteSize;
}
hipError_t ihipMemcpy3D(const struct hipMemcpy3DParms* p, hipStream_t stream, bool isAsync) {
hipError_t e = hipSuccess;
if(p) {
size_t byteSize, width, height, depth, widthInBytes, srcPitch, dstPitch, ySize;
hipChannelFormatDesc desc;
void* srcPtr;void* dstPtr;
size_t dstByteSize, srcByteSize, copyWidth, copyHeight, copyDepth, widthInBytes, srcPitch, dstPitch, srcYsize, dstYsize;
size_t srcXoffset, srcYoffset, srcZoffset, dstXoffset, dstYoffset, dstZoffset;
size_t srcWidth, srcHeight, srcDepth, dstWidth, dstHeight, dstDepth;
void* srcPtr, *dstPtr;
bool copyWidthUpdate= false;
copyDepth = p->extent.depth;
copyHeight = p->extent.height;
copyWidth = p->extent.width; // in bytes ?
dstXoffset = p->dstPos.x;
dstYoffset = p->dstPos.y;
dstZoffset = p->dstPos.z;
srcXoffset = p->srcPos.x;
srcYoffset = p->srcPos.y;
srcZoffset = p->srcPos.z;
if (p->dstArray != nullptr) {
if (p->dstArray->isDrv == false) {
switch (p->dstArray->desc.f) {
case hipChannelFormatKindSigned:
byteSize = sizeof(int);
break;
case hipChannelFormatKindUnsigned:
byteSize = sizeof(unsigned int);
break;
case hipChannelFormatKindFloat:
byteSize = sizeof(float);
break;
case hipChannelFormatKindNone:
byteSize = sizeof(size_t);
break;
default:
byteSize = 0;
break;
}
depth = p->extent.depth;
height = p->extent.height;
width = p->extent.width;
widthInBytes = p->extent.width * byteSize;
srcPitch = p->srcPtr.pitch;
srcPtr = p->srcPtr.ptr;
ySize = p->srcPtr.ysize;
desc = p->dstArray->desc;
dstPtr = p->dstArray->data;
hsa_ext_image_data_info_t imageInfo;
if(hipTextureType2DLayered == p->dstArray->textureType)
GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_2DA, width, height, 0, desc, imageInfo, depth);
else
GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_3D, width, height, depth, desc, imageInfo);
dstPitch = imageInfo.size/(height == 0 ? 1 : height)/(depth == 0 ? 1 : depth);
} else {
depth = p->Depth;
height = p->Height;
widthInBytes = p->WidthInBytes;
width = p->dstArray->width;
hsa_ext_image_channel_order_t channelOrder;
switch(p->dstArray->NumChannels) {
case 2:
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG;
break;
case 3:
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGB;
break;
case 4:
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA;
break;
case 1:
default:
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R;
break;
}
hsa_ext_image_channel_type_t channelType;
e = ihipArrayToImageFormat(p->dstArray->Format,channelType);
srcPitch = p->srcPitch;
srcPtr = (void*)p->srcHost;
ySize = p->srcHeight;
dstPtr = p->dstArray->data;
hsa_ext_image_data_info_t imageInfo;
if(hipTextureType2DLayered == p->dstArray->textureType)
GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_2DA, width, height, 0, channelOrder, channelType, imageInfo, depth);
else
GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_3D, width, height, depth, channelOrder, channelType, imageInfo);
dstPitch = imageInfo.size/(height == 0 ? 1 : height)/(depth == 0 ? 1 : depth);
if ((p->dstArray->isDrv == true) ||( p->dstPtr.ptr!= nullptr)){
return hipErrorInvalidValue;
}
// Array destination
dstByteSize = getByteSizeFromFormat(p->dstArray->desc);
hipChannelFormatDesc desc;
desc = p->dstArray->desc;
dstPtr = p->dstArray->data;
dstWidth = p->dstArray->width;
dstHeight = p->dstArray->height;
dstDepth = p->dstArray->depth;
dstPitch = dstByteSize * alignUp(dstWidth, IMAGE_PITCH_ALIGNMENT);
if(!copyWidthUpdate) {
copyWidth = copyWidth * dstByteSize;
copyWidthUpdate = true;
}
} else {
// Non array destination
depth = p->extent.depth;
height = p->extent.height;
widthInBytes = p->extent.width;
srcPitch = p->srcPtr.pitch;
srcPtr = p->srcPtr.ptr;
//Non Array destination
dstPtr = p->dstPtr.ptr;
ySize = p->srcPtr.ysize;
dstWidth = p->dstPtr.xsize;
dstHeight = p->dstPtr.ysize;
dstPitch = p->dstPtr.pitch;
}
if (p->srcArray != nullptr) {
if ((p->srcArray->isDrv == true) ||( p->srcPtr.ptr!= nullptr)){
return hipErrorInvalidValue;
}
// Array source
srcByteSize = getByteSizeFromFormat(p->srcArray->desc);
hipChannelFormatDesc desc;
desc = p->srcArray->desc;
srcPtr = p->srcArray->data;
srcWidth = p->srcArray->width;
srcHeight = p->srcArray->height;
srcDepth = p->srcArray->depth;
srcPitch = srcByteSize * alignUp(srcWidth, IMAGE_PITCH_ALIGNMENT);
if(!copyWidthUpdate) {
copyWidth = copyWidth * srcByteSize;
copyWidthUpdate = true;
}
} else {
//Non Array source
srcPtr = p->srcPtr.ptr;
srcWidth = p->srcPtr.xsize;
srcHeight = p->srcPtr.ysize;
srcPitch = p->srcPtr.pitch;
}
stream = ihipSyncAndResolveStream(stream);
try {
if((widthInBytes == dstPitch) && (widthInBytes == srcPitch)) {
if((copyWidth == dstPitch) && (copyWidth == srcPitch)&& (copyHeight == dstHeight) &&(copyHeight == srcHeight)) {
if(isAsync)
stream->locked_copyAsync((void*)dstPtr, (void*)srcPtr, widthInBytes*height*depth, p->kind);
stream->locked_copyAsync((void*)dstPtr, (void*)srcPtr, copyWidth*copyHeight*copyDepth, p->kind);
else
stream->locked_copySync((void*)dstPtr, (void*)srcPtr, widthInBytes*height*depth, p->kind, false);
stream->locked_copySync((void*)dstPtr, (void*)srcPtr, copyWidth*copyHeight*copyDepth, p->kind, false);
} else {
for (int i = 0; i < depth; i++) {
for (int j = 0; j < height; j++) {
// TODO: p->srcPos or p->dstPos are not 0.
for (int i = 0; i < copyDepth; i++) {
for (int j = 0; j < copyHeight; j++) {
unsigned char* src =
(unsigned char*)srcPtr + i * ySize * srcPitch + j * srcPitch;
(unsigned char*)srcPtr + (i + srcZoffset) * srcHeight * srcPitch + (j + srcYoffset) * srcPitch + srcXoffset;
unsigned char* dst =
(unsigned char*)dstPtr + i * height * dstPitch + j * dstPitch;
(unsigned char*)dstPtr + (i + dstZoffset) * dstHeight * dstPitch + (j + dstYoffset) * dstPitch + dstXoffset;
if(isAsync)
stream->locked_copyAsync(dst, src, widthInBytes, p->kind);
stream->locked_copyAsync(dst, src, copyWidth, p->kind);
else
stream->locked_copySync(dst, src, widthInBytes, p->kind);
stream->locked_copySync(dst, src, copyWidth, p->kind);
}
}
}
src/hip_module.cpp
+78 -35
View File
@@ -134,10 +134,10 @@ extern hipError_t ihipGetDeviceProperties(hipDeviceProp_t* props, int device);
return ihipLogStatus(hipStatus); \
}
hipError_t ihipModuleLaunchKernel(TlsData *tls, hipFunction_t f, uint32_t gridSizeX,
uint32_t gridSizeY, uint32_t gridSizeZ,
uint32_t blockSizeX, uint32_t blockSizeY,
uint32_t blockSizeZ, size_t sharedMemBytes,
hipError_t ihipModuleLaunchKernel(TlsData *tls, hipFunction_t f, uint32_t globalWorkSizeX,
uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
uint32_t localWorkSizeX, uint32_t localWorkSizeY,
uint32_t localWorkSizeZ, size_t sharedMemBytes,
hipStream_t hStream, void** kernelParams, void** extra,
hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags, bool isStreamLocked = 0,
void** impCoopParams = 0) {
@@ -146,14 +146,6 @@ hipError_t ihipModuleLaunchKernel(TlsData *tls, hipFunction_t f, uint32_t gridSi
auto ctx = ihipGetTlsDefaultCtx();
hipError_t ret = hipSuccess;
size_t globalWorkSizeX = (size_t)gridSizeX * (size_t)blockSizeX;
size_t globalWorkSizeY = (size_t)gridSizeY * (size_t)blockSizeY;
size_t globalWorkSizeZ = (size_t)gridSizeZ * (size_t)blockSizeZ;
if(globalWorkSizeX > UINT32_MAX || globalWorkSizeY > UINT32_MAX || globalWorkSizeZ > UINT32_MAX)
{
return hipErrorInvalidConfiguration;
}
if (ctx == nullptr) {
ret = hipErrorInvalidDevice;
@@ -211,8 +203,8 @@ hipError_t ihipModuleLaunchKernel(TlsData *tls, hipFunction_t f, uint32_t gridSi
lp.dynamic_group_mem_bytes =
sharedMemBytes; // TODO - this should be part of preLaunchKernel.
hStream = ihipPreLaunchKernel(
hStream, dim3(globalWorkSizeX/blockSizeX, globalWorkSizeY/blockSizeY, globalWorkSizeZ/blockSizeZ),
dim3(blockSizeX, blockSizeY, blockSizeZ), &lp, f->_name.c_str(), isStreamLocked);
hStream, dim3(globalWorkSizeX/localWorkSizeX, globalWorkSizeY/localWorkSizeY, globalWorkSizeZ/localWorkSizeZ),
dim3(localWorkSizeX, localWorkSizeY, localWorkSizeZ), &lp, f->_name.c_str(), isStreamLocked);
hsa_kernel_dispatch_packet_t aql;
@@ -221,9 +213,9 @@ hipError_t ihipModuleLaunchKernel(TlsData *tls, hipFunction_t f, uint32_t gridSi
// aql.completion_signal._handle = 0;
// aql.kernarg_address = 0;
aql.workgroup_size_x = blockSizeX;
aql.workgroup_size_y = blockSizeY;
aql.workgroup_size_z = blockSizeZ;
aql.workgroup_size_x = localWorkSizeX;
aql.workgroup_size_y = localWorkSizeY;
aql.workgroup_size_z = localWorkSizeZ;
aql.grid_size_x = globalWorkSizeX;
aql.grid_size_y = globalWorkSizeY;
aql.grid_size_z = globalWorkSizeZ;
@@ -283,8 +275,17 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, uint32_t gridDimX, uint32_t gr
void** kernelParams, void** extra) {
HIP_INIT_API(hipModuleLaunchKernel, f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes,
hStream, kernelParams, extra);
size_t globalWorkSizeX = (size_t)gridDimX * (size_t)blockDimX;
size_t globalWorkSizeY = (size_t)gridDimY * (size_t)blockDimY;
size_t globalWorkSizeZ = (size_t)gridDimZ * (size_t)blockDimZ;
if(globalWorkSizeX > UINT32_MAX || globalWorkSizeY > UINT32_MAX || globalWorkSizeZ > UINT32_MAX)
{
return hipErrorInvalidConfiguration;
}
return ihipLogStatus(ihipModuleLaunchKernel(tls,
f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY,
f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, blockDimX, blockDimY,
blockDimZ, sharedMemBytes, hStream, kernelParams, extra, nullptr, nullptr, 0));
}
@@ -297,11 +298,8 @@ hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
HIP_INIT_API(hipExtModuleLaunchKernel, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX,
localWorkSizeY, localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra);
if(localWorkSizeX == 0 || localWorkSizeY == 0 || localWorkSizeZ == 0)
return hipErrorInvalidValue;
return ihipLogStatus(ihipModuleLaunchKernel(tls,
f, globalWorkSizeX/localWorkSizeX, globalWorkSizeY/localWorkSizeY, globalWorkSizeZ/localWorkSizeZ, localWorkSizeX, localWorkSizeY,
f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY,
localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, flags));
}
@@ -314,11 +312,8 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
HIP_INIT_API(hipHccModuleLaunchKernel, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX,
localWorkSizeY, localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra);
if(localWorkSizeX == 0 || localWorkSizeY == 0 || localWorkSizeZ == 0)
return hipErrorInvalidValue;
return ihipLogStatus(ihipModuleLaunchKernel(tls,
f, globalWorkSizeX/localWorkSizeX, globalWorkSizeY/localWorkSizeY, globalWorkSizeZ/localWorkSizeZ, localWorkSizeX, localWorkSizeY,
f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY,
localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, 0));
}
@@ -364,14 +359,26 @@ hipError_t ihipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList
}
GET_TLS();
size_t globalWorkSizeX = 0, globalWorkSizeY = 0, globalWorkSizeZ = 0;
// launch kernels for each device
for (int i = 0; i < numDevices; ++i) {
const hipLaunchParams& lp = launchParamsList[i];
globalWorkSizeX = (size_t)lp.gridDim.x * (size_t)lp.blockDim.x;
globalWorkSizeY = (size_t)lp.gridDim.y * (size_t)lp.blockDim.y;
globalWorkSizeZ = (size_t)lp.gridDim.z * (size_t)lp.blockDim.z;
if(globalWorkSizeX > UINT32_MAX || globalWorkSizeY > UINT32_MAX || globalWorkSizeZ > UINT32_MAX)
{
return hipErrorInvalidConfiguration;
}
result = ihipModuleLaunchKernel(tls, kds[i],
lp.gridDim.x,
lp.gridDim.y,
lp.gridDim.z,
lp.gridDim.x * lp.blockDim.x,
lp.gridDim.y * lp.blockDim.y,
lp.gridDim.z * lp.blockDim.z,
lp.blockDim.x, lp.blockDim.y,
lp.blockDim.z, lp.sharedMem,
lp.stream, lp.args, nullptr, nullptr, nullptr, 0,
@@ -424,6 +431,14 @@ hipError_t ihipLaunchCooperativeKernel(const void* f, dim3 gridDim,
return hipErrorInvalidConfiguration;
}
size_t globalWorkSizeX = (size_t)gridDim.x * (size_t)blockDimX.x;
size_t globalWorkSizeY = (size_t)gridDim.y * (size_t)blockDimX.y;
size_t globalWorkSizeZ = (size_t)gridDim.z * (size_t)blockDimX.z;
if(globalWorkSizeX > UINT32_MAX || globalWorkSizeY > UINT32_MAX || globalWorkSizeZ > UINT32_MAX)
{
return hipErrorInvalidConfiguration;
}
// Prepare the kernel descriptor for initializing the GWS
hipFunction_t gwsKD = ps.kernel_descriptor(
reinterpret_cast<std::uintptr_t>(&init_gws),
@@ -483,9 +498,9 @@ hipError_t ihipLaunchCooperativeKernel(const void* f, dim3 gridDim,
// launch the main kernel
result = ihipModuleLaunchKernel(tls, kd,
gridDim.x,
gridDim.y,
gridDim.z,
gridDim.x * blockDimX.x,
gridDim.y * blockDimX.y,
gridDim.z * blockDimX.z,
blockDimX.x, blockDimX.y, blockDimX.z,
sharedMemBytes, stream, kernelParams, nullptr, nullptr,
nullptr, 0, true, impCoopParams);
@@ -620,6 +635,8 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL
void* impCoopParams[1];
ulong prev_sum = 0;
size_t globalWorkSizeX = 0, globalWorkSizeY = 0, globalWorkSizeZ = 0;
// launch the main kernels for each device
for (int i = 0; i < numDevices; ++i) {
const hipLaunchParams& lp = launchParamsList[i];
@@ -636,10 +653,18 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL
impCoopParams[0] = &mg_info_ptr[i];
globalWorkSizeX = (size_t)lp.gridDim.x * (size_t)lp.blockDim.x;
globalWorkSizeY = (size_t)lp.gridDim.y * (size_t)lp.blockDim.y;
globalWorkSizeZ = (size_t)lp.gridDim.z * (size_t)lp.blockDim.z;
if(globalWorkSizeX > UINT32_MAX || globalWorkSizeY > UINT32_MAX || globalWorkSizeZ > UINT32_MAX)
{
return hipErrorInvalidConfiguration;
}
result = ihipModuleLaunchKernel(tls, kds[i],
lp.gridDim.x,
lp.gridDim.y,
lp.gridDim.z,
lp.gridDim.x * lp.blockDim.x,
lp.gridDim.y * lp.blockDim.y,
lp.gridDim.z * lp.blockDim.z,
lp.blockDim.x, lp.blockDim.y,
lp.blockDim.z, lp.sharedMem,
lp.stream, lp.args, nullptr, nullptr, nullptr, 0,
@@ -1469,6 +1494,15 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(
tls, numBlocks, f, blockSize, dynSharedMemPerBlk));
}
hipError_t hipDrvOccupancyMaxActiveBlocksPerMultiprocessor(
int* numBlocks, hipFunction_t f, int blockSize, size_t dynSharedMemPerBlk)
{
HIP_INIT_API(hipDrvOccupancyMaxActiveBlocksPerMultiprocessor, numBlocks, f, blockSize, dynSharedMemPerBlk);
return ihipLogStatus(ihipOccupancyMaxActiveBlocksPerMultiprocessor(
tls, (uint32_t*) numBlocks, f, blockSize, dynSharedMemPerBlk));
}
hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk,
unsigned int flags)
@@ -1479,6 +1513,15 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
tls, numBlocks, f, blockSize, dynSharedMemPerBlk));
}
hipError_t hipDrvOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
int* numBlocks, hipFunction_t f, int blockSize, size_t dynSharedMemPerBlk,
unsigned int flags)
{
HIP_INIT_API(hipDrvOccupancyMaxActiveBlocksPerMultiprocessorWithFlags, numBlocks, f, blockSize, dynSharedMemPerBlk, flags);
return ihipLogStatus(ihipOccupancyMaxActiveBlocksPerMultiprocessor(
tls, (uint32_t*) numBlocks, f, blockSize, dynSharedMemPerBlk));
}
hipError_t hipLaunchKernel(
const void* func_addr, dim3 numBlocks, dim3 dimBlocks, void** args,
size_t sharedMemBytes, hipStream_t stream)
src/hip_peer.cpp
+17 -10
View File
@@ -128,17 +128,24 @@ hipError_t ihipEnablePeerAccess(TlsData* tls, hipCtx_t peerCtx, unsigned int fla
if (thisCtx == peerCtx) {
err = hipErrorInvalidDevice; // Can't enable peer access to self.
} else if ((thisCtx != NULL) && (peerCtx != NULL)) {
LockedAccessor_CtxCrit_t peerCrit(peerCtx->criticalData());
// Add thisCtx to peerCtx's access list so that new allocations on peer will be made
// visible to this device:
bool isNewPeer = peerCrit->addPeerWatcher(peerCtx, thisCtx);
if (isNewPeer) {
tprintf(DB_MEM, "device=%s can now see all memory allocated on peer=%s\n",
thisCtx->toString().c_str(), peerCtx->toString().c_str());
am_memtracker_update_peers(peerCtx->getDevice()->_acc, peerCrit->peerCnt(),
peerCrit->peerAgents());
int canAccess = 0;
if ((hipSuccess != ihipDeviceCanAccessPeer(&canAccess,thisCtx,peerCtx)) || (canAccess == 0)){
tprintf(DB_MEM, "device=%s can't access peer=%s\n",thisCtx->toString().c_str(), peerCtx->toString().c_str());
err = hipErrorInvalidDevice;
} else {
err = hipErrorPeerAccessAlreadyEnabled;
LockedAccessor_CtxCrit_t peerCrit(peerCtx->criticalData());
// Add thisCtx to peerCtx's access list so that new allocations on peer will be made
// visible to this device:
bool isNewPeer = peerCrit->addPeerWatcher(peerCtx, thisCtx);
if (isNewPeer) {
tprintf(DB_MEM, "device=%s can now see all memory allocated on peer=%s\n",
thisCtx->toString().c_str(), peerCtx->toString().c_str());
am_memtracker_update_peers(peerCtx->getDevice()->_acc, peerCrit->peerCnt(),
peerCrit->peerAgents());
} else {
err = hipErrorPeerAccessAlreadyEnabled;
}
}
} else {
err = hipErrorInvalidDevice;
src/hiprtc.cpp
+38 -26
View File
@@ -50,7 +50,7 @@ THE SOFTWARE.
#include <iostream>
#include <sys/stat.h>
const char* hiprtcGetErrorString(hiprtcResult x)
extern "C" const char* hiprtcGetErrorString(hiprtcResult x)
{
switch (x) {
case HIPRTC_SUCCESS:
@@ -95,6 +95,21 @@ inline bool fileExists (const std::string& name) {
}
} // namespace hip_impl
namespace
{
char* demangle(const char* x)
{
if (!x) return nullptr;
int s{};
char* tmp = abi::__cxa_demangle(x, nullptr, nullptr, &s);
if (s != 0) return nullptr;
return tmp;
}
} // Unnamed namespace.
namespace
{
struct Symbol {
@@ -158,7 +173,7 @@ struct _hiprtcProgram {
{
using namespace std;
char* demangled = hip_impl::demangle(name.c_str());
char* demangled = demangle(name.c_str());
name.assign(demangled == nullptr ? "" : demangled);
free(demangled);
@@ -352,7 +367,7 @@ namespace
}
} // Unnamed namespace.
hiprtcResult hiprtcAddNameExpression(hiprtcProgram p, const char* n)
extern "C" hiprtcResult hiprtcAddNameExpression(hiprtcProgram p, const char* n)
{
if (!n) return HIPRTC_ERROR_INVALID_INPUT;
if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM;
@@ -413,21 +428,6 @@ namespace
};
} // Unnamed namespace.
namespace hip_impl
{
char* demangle(const char* x)
{
if (!x) return nullptr;
int s{};
char* tmp = abi::__cxa_demangle(x, nullptr, nullptr, &s);
if (s != 0) return nullptr;
return tmp;
}
} // Namespace hip_impl.
namespace
{
const std::string& defaultTarget()
@@ -492,7 +492,7 @@ namespace
}
} // Unnamed namespace.
hiprtcResult hiprtcCompileProgram(hiprtcProgram p, int n, const char** o)
extern "C" hiprtcResult hiprtcCompileProgram(hiprtcProgram p, int n, const char** o)
{
using namespace std;
@@ -530,7 +530,7 @@ hiprtcResult hiprtcCompileProgram(hiprtcProgram p, int n, const char** o)
return HIPRTC_SUCCESS;
}
hiprtcResult hiprtcCreateProgram(hiprtcProgram* p, const char* src,
extern "C" hiprtcResult hiprtcCreateProgram(hiprtcProgram* p, const char* src,
const char* name, int n, const char** hdrs,
const char** incs)
{
@@ -548,14 +548,14 @@ hiprtcResult hiprtcCreateProgram(hiprtcProgram* p, const char* src,
return HIPRTC_SUCCESS;
}
hiprtcResult hiprtcDestroyProgram(hiprtcProgram* p)
extern "C" hiprtcResult hiprtcDestroyProgram(hiprtcProgram* p)
{
if (!p) return HIPRTC_SUCCESS;
return _hiprtcProgram::destroy(*p);
}
hiprtcResult hiprtcGetLoweredName(hiprtcProgram p, const char* n,
extern "C" hiprtcResult hiprtcGetLoweredName(hiprtcProgram p, const char* n,
const char** ln)
{
using namespace std;
@@ -576,7 +576,7 @@ hiprtcResult hiprtcGetLoweredName(hiprtcProgram p, const char* n,
return HIPRTC_SUCCESS;
}
hiprtcResult hiprtcGetProgramLog(hiprtcProgram p, char* l)
extern "C" hiprtcResult hiprtcGetProgramLog(hiprtcProgram p, char* l)
{
if (!l) return HIPRTC_ERROR_INVALID_INPUT;
if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM;
@@ -588,7 +588,7 @@ hiprtcResult hiprtcGetProgramLog(hiprtcProgram p, char* l)
return HIPRTC_SUCCESS;
}
hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram p, std::size_t* sz)
extern "C" hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram p, std::size_t* sz)
{
if (!sz) return HIPRTC_ERROR_INVALID_INPUT;
if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM;
@@ -599,7 +599,7 @@ hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram p, std::size_t* sz)
return HIPRTC_SUCCESS;
}
hiprtcResult hiprtcGetCode(hiprtcProgram p, char* c)
extern "C" hiprtcResult hiprtcGetCode(hiprtcProgram p, char* c)
{
if (!c) return HIPRTC_ERROR_INVALID_INPUT;
if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM;
@@ -610,7 +610,7 @@ hiprtcResult hiprtcGetCode(hiprtcProgram p, char* c)
return HIPRTC_SUCCESS;
}
hiprtcResult hiprtcGetCodeSize(hiprtcProgram p, std::size_t* sz)
extern "C" hiprtcResult hiprtcGetCodeSize(hiprtcProgram p, std::size_t* sz)
{
if (!sz) return HIPRTC_ERROR_INVALID_INPUT;
if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM;
@@ -620,3 +620,15 @@ hiprtcResult hiprtcGetCodeSize(hiprtcProgram p, std::size_t* sz)
return HIPRTC_SUCCESS;
}
extern "C" hiprtcResult hiprtcVersion(int* major, int* minor)
{
if (major == nullptr || minor == nullptr) {
return HIPRTC_ERROR_INVALID_INPUT;
}
*major = 9;
*minor = 0;
return HIPRTC_SUCCESS;
}
src/program_state.inl
+11 -5
View File
@@ -26,6 +26,7 @@
#include <cstddef>
#include <cstdint>
#include <cstdio>
#include <deque>
#include <memory>
#include <mutex>
#include <stdexcept>
@@ -202,7 +203,7 @@ public:
std::function<void(hsa_code_object_reader_t*)>>;
std::pair<
std::mutex,
std::vector<std::pair<std::string, RAII_code_reader>>> code_readers;
std::deque<std::pair<std::string, RAII_code_reader>>> code_readers;
program_state_impl() {
// Create placeholder for each agent for the per-agent members.
@@ -244,7 +245,8 @@ public:
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);
if(bundle.blob.size())
impl.code_object_blobs.second[elf][triple_to_hsa_isa(bundle.triple)].push_back(bundle.blob);
}
blob_it += tmp.bundled_code_size;
@@ -418,13 +420,17 @@ public:
decltype(code_readers.second)::iterator it;
{
std::lock_guard<std::mutex> lck{code_readers.first};
it = code_readers.second.emplace(code_readers.second.end(),
move(file), move(tmp));
code_readers.second.emplace_back(move(file), move(tmp));
it = std::prev(code_readers.second.end());
}
auto check_hsa_error = [](hsa_status_t s) {
if (s != HSA_STATUS_SUCCESS) {
hip_throw(std::runtime_error{"error when loading code object"});
const char* hsa_err_msg;
hsa_status_string(s, &hsa_err_msg);
hip_throw(std::runtime_error{
std::string("error when loading code object: ") +
hsa_err_msg});
}
};