AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

Merge 'master' into 'amd-master'

Просмотр исходного кода 
Change-Id: Icc364a87efd81ceb59122b1a33af8b3837ae1ab3


[ROCm/hip commit: a865f24d1c]
Этот коммит содержится в:
Jenkins
2018-11-23 04:09:47 -06:00
родитель 3572acefd7 e317a9a1fa
Коммит 5f41a4aef3
18 изменённых файлов: 2997 добавлений и 194 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
projects/hip/CMakeLists.txt
+18
Просмотреть файл
@@ -150,6 +150,23 @@ if(NOT DEFINED COMPILE_HIP_ATP_MARKER)
endif()
add_to_config(_buildInfo COMPILE_HIP_ATP_MARKER)
################
# Detect profiling API
################
if(USE_PROF_API EQUAL 1)
if(NOT DEFINED PROF_API_HEADER_PATH)
set(PROF_API_HEADER_PATH /opt/rocm/roctracer/include/ext)
endif ()
find_path(PROF_API_HEADER_DIR NAMES prof_protocol.h PATHS ${PROF_API_HEADER_PATH} NO_DEFAULT_PATH)
if(NOT PROF_API_HEADER_DIR)
MESSAGE("PROF_API_HEADER_PATH = ${PROF_API_HEADER_PATH}")
MESSAGE(FATAL_ERROR "Profiling API header not found, use -DPROF_API_HEADER_PATH=<path to prof_protocol.h header>")
else()
add_definitions(-DUSE_PROF_API=1)
include_directories ( ${PROF_API_HEADER_DIR} )
MESSAGE(STATUS "Profiling API: ${PROF_API_HEADER_DIR}")
endif()
endif()
#############################
# Build steps
@@ -204,6 +221,7 @@ if(HIP_PLATFORM STREQUAL "hcc")
src/grid_launch.cpp
src/hip_texture.cpp
src/hip_surface.cpp
src/hip_intercept.cpp
src/env.cpp
src/program_state.cpp)
projects/hip/include/hip/hcc_detail/hip_prof_api.h
+204
Просмотреть файл
@@ -0,0 +1,204 @@
// automatically generated sources
#ifndef _HIP_PROF_API_H
#define _HIP_PROF_API_H
#include <atomic>
#include <iostream>
#include <mutex>
#include "hip/hcc_detail/hip_prof_str.h"
template <typename Record, typename Fun, typename Act>
class api_callbacks_table_templ {
public:
typedef std::recursive_mutex mutex_t;
typedef Record record_t;
typedef Fun fun_t;
typedef Act act_t;
// HIP API callbacks table
struct hip_cb_table_entry_t {
volatile std::atomic<bool> sync;
volatile std::atomic<uint32_t> sem;
act_t act;
void* a_arg;
fun_t fun;
void* arg;
};
struct hip_cb_table_t {
hip_cb_table_entry_t arr[HIP_API_ID_NUMBER];
};
api_callbacks_table_templ() {
memset(&callbacks_table_, 0, sizeof(callbacks_table_));
}
bool set_activity(uint32_t id, act_t fun, void* arg) {
std::lock_guard<mutex_t> lock(mutex_);
bool ret = true;
if (id == HIP_API_ID_ANY) {
for (unsigned i = 0; i < HIP_API_ID_NUMBER; ++i) set_activity(i, fun, arg);
} else if (id < HIP_API_ID_NUMBER) {
cb_sync(id);
callbacks_table_.arr[id].act = fun;
callbacks_table_.arr[id].a_arg = arg;
cb_release(id);
} else {
ret = false;
}
return ret;
}
bool set_callback(uint32_t id, fun_t fun, void* arg) {
std::lock_guard<mutex_t> lock(mutex_);
bool ret = true;
if (id == HIP_API_ID_ANY) {
for (unsigned i = 0; i < HIP_API_ID_NUMBER; ++i) set_callback(i, fun, arg);
} else if (id < HIP_API_ID_NUMBER) {
cb_sync(id);
callbacks_table_.arr[id].fun = fun;
callbacks_table_.arr[id].arg = arg;
cb_release(id);
} else {
ret = false;
}
return ret;
}
inline hip_cb_table_entry_t& entry(const uint32_t& id) {
return callbacks_table_.arr[id];
}
inline void sem_sync(const uint32_t& id) {
sem_increment(id);
if (entry(id).sync.load() == true) sync_wait(id);
}
inline void sem_release(const uint32_t& id) {
sem_decrement(id);
}
private:
inline void cb_sync(const uint32_t& id) {
entry(id).sync.store(true);
while (entry(id).sem.load() != 0) {}
}
inline void cb_release(const uint32_t& id) {
entry(id).sync.store(false);
}
inline void sem_increment(const uint32_t& id) {
const uint32_t prev = entry(id).sem.fetch_add(1);
if (prev == UINT32_MAX) {
std::cerr << "sem overflow id = " << id << std::endl << std::flush;
abort();
}
}
inline void sem_decrement(const uint32_t& id) {
const uint32_t prev = entry(id).sem.fetch_sub(1);
if (prev == 0) {
std::cerr << "sem corrupted id = " << id << std::endl << std::flush;
abort();
}
}
void sync_wait(const uint32_t& id) {
sem_decrement(id);
while (entry(id).sync.load() == true) {}
sem_increment(id);
}
mutex_t mutex_;
hip_cb_table_t callbacks_table_;
};
#if USE_PROF_API
#include <prof_protocol.h>
static const uint32_t HIP_DOMAIN_ID = ACTIVITY_DOMAIN_HIP_API;
typedef activity_record_t hip_api_record_t;
typedef activity_rtapi_callback_t hip_api_callback_t;
typedef activity_sync_callback_t hip_act_callback_t;
// HIP API callbacks spawner object macro
#define HIP_CB_SPAWNER_OBJECT(CB_ID) \
hip_api_data_t api_data{}; \
INIT_CB_ARGS_DATA(CB_ID, api_data); \
api_callbacks_spawner_t<HIP_API_ID_##CB_ID> __api_tracer(HIP_API_ID_##CB_ID, api_data);
typedef api_callbacks_table_templ<hip_api_record_t,
hip_api_callback_t,
hip_act_callback_t> api_callbacks_table_t;
extern api_callbacks_table_t callbacks_table;
template <int cid_>
class api_callbacks_spawner_t {
public:
api_callbacks_spawner_t(const hip_api_id_t& cid, hip_api_data_t& api_data) :
api_data_(api_data),
record_({})
{
if (cid_ >= HIP_API_ID_NUMBER) {
fprintf(stderr, "HIP %s bad id %d\n", __FUNCTION__, cid_);
abort();
}
callbacks_table.sem_sync(cid_);
act = entry(cid_).act;
a_arg = entry(cid_).a_arg;
fun = entry(cid_).fun;
arg = entry(cid_).arg;
api_data_.phase = 0;
if (act != NULL) act(cid_, &record_, &api_data_, a_arg);
if (fun != NULL) fun(HIP_DOMAIN_ID, cid_, &api_data_, arg);
}
~api_callbacks_spawner_t() {
api_data_.phase = 1;
if (act != NULL) act(cid_, &record_, &api_data_, a_arg);
if (fun != NULL) fun(HIP_DOMAIN_ID, cid_, &api_data_, arg);
callbacks_table.sem_release(cid_);
}
private:
inline api_callbacks_table_t::hip_cb_table_entry_t& entry(const uint32_t& id) {
return callbacks_table.entry(id);
}
hip_api_data_t& api_data_;
hip_api_record_t record_;
hip_act_callback_t act;
void* a_arg;
hip_api_callback_t fun;
void* arg;
};
template <>
class api_callbacks_spawner_t<HIP_API_ID_NUMBER> {
public:
api_callbacks_spawner_t(const hip_api_id_t& cid, hip_api_data_t& api_data) {}
};
#else
#define HIP_CB_SPAWNER_OBJECT(x) do {} while(0)
class api_callbacks_table_t {
public:
typedef void* act_t;
typedef void* fun_t;
bool set_activity(uint32_t id, act_t fun, void* arg) { return true; }
bool set_callback(uint32_t id, fun_t fun, void* arg) { return true; }
};
#endif
#endif // _HIP_PROF_API_H
projects/hip/include/hip/hcc_detail/hip_prof_str.h
+2508
Просмотреть файл
Разница между файлами не показана из-за своего большого размера Загрузить разницу
projects/hip/include/hip/hcc_detail/hip_runtime_api.h
+18
Просмотреть файл
@@ -2610,6 +2610,24 @@ hipError_t hipLaunchByPtr(const void* func);
} /* extern "c" */
#endif
#include <hip/hcc_detail/hip_prof_api.h>
#ifdef __cplusplus
extern "C" {
#endif
/**
* Callback/Activity API
*/
hipError_t hipRegisterApiCallback(uint32_t id, void* fun, void* arg);
hipError_t hipRemoveApiCallback(uint32_t id);
hipError_t hipRegisterActivityCallback(uint32_t id, void* fun, void* arg);
hipError_t hipRemoveActivityCallback(uint32_t id);
static inline const char* hipApiName(const uint32_t& id) { return hip_api_name(id); }
const char* hipKernelNameRef(hipFunction_t f);
#ifdef __cplusplus
} /* extern "C" */
#endif
#ifdef __cplusplus
hipError_t hipBindTexture(size_t* offset, textureReference* tex, const void* devPtr,
projects/hip/src/hip_clang.cpp
+3 -3
Просмотреть файл
@@ -159,7 +159,7 @@ extern "C" void __hipRegisterFunction(
dim3* gridDim,
int* wSize)
{
HIP_INIT_API(modules, hostFunction, deviceFunction, deviceName);
HIP_INIT_API(NONE, modules, hostFunction, deviceFunction, deviceName);
std::vector<hipFunction_t> functions{g_deviceCnt};
assert(modules && modules->size() >= g_deviceCnt);
@@ -214,7 +214,7 @@ hipError_t hipSetupArgument(
size_t size,
size_t offset)
{
HIP_INIT_API(arg, size, offset);
HIP_INIT_API(hipSetupArgument, arg, size, offset);
auto ctx = ihipGetTlsDefaultCtx();
LockedAccessor_CtxCrit_t crit(ctx->criticalData());
auto& arguments = crit->_execStack.top()._arguments;
@@ -229,7 +229,7 @@ hipError_t hipSetupArgument(
hipError_t hipLaunchByPtr(const void *hostFunction)
{
HIP_INIT_API(hostFunction);
HIP_INIT_API(hipLaunchByPtr, hostFunction);
ihipExec_t exec;
{
auto ctx = ihipGetTlsDefaultCtx();
projects/hip/src/hip_context.cpp
+23 -23
Просмотреть файл
@@ -40,7 +40,7 @@ void ihipCtxStackUpdate() {
}
hipError_t hipInit(unsigned int flags) {
HIP_INIT_API(flags);
HIP_INIT_API(hipInit, flags);
hipError_t e = hipSuccess;
@@ -53,7 +53,7 @@ hipError_t hipInit(unsigned int flags) {
}
hipError_t hipCtxCreate(hipCtx_t* ctx, unsigned int flags, hipDevice_t device) {
HIP_INIT_API(ctx, flags, device); // FIXME - review if we want to init
HIP_INIT_API(hipCtxCreate, ctx, flags, device); // FIXME - review if we want to init
hipError_t e = hipSuccess;
auto deviceHandle = ihipGetDevice(device);
{
@@ -71,7 +71,7 @@ hipError_t hipCtxCreate(hipCtx_t* ctx, unsigned int flags, hipDevice_t device) {
}
hipError_t hipDeviceGet(hipDevice_t* device, int deviceId) {
HIP_INIT_API(device, deviceId); // FIXME - review if we want to init
HIP_INIT_API(hipDeviceGet, device, deviceId); // FIXME - review if we want to init
auto deviceHandle = ihipGetDevice(deviceId);
@@ -86,7 +86,7 @@ hipError_t hipDeviceGet(hipDevice_t* device, int deviceId) {
};
hipError_t hipDriverGetVersion(int* driverVersion) {
HIP_INIT_API(driverVersion);
HIP_INIT_API(hipDriverGetVersion, driverVersion);
hipError_t e = hipSuccess;
if (driverVersion) {
*driverVersion = 4;
@@ -98,7 +98,7 @@ hipError_t hipDriverGetVersion(int* driverVersion) {
}
hipError_t hipRuntimeGetVersion(int* runtimeVersion) {
HIP_INIT_API(runtimeVersion);
HIP_INIT_API(hipRuntimeGetVersion, runtimeVersion);
hipError_t e = hipSuccess;
if (runtimeVersion) {
*runtimeVersion = HIP_VERSION_PATCH;
@@ -110,7 +110,7 @@ hipError_t hipRuntimeGetVersion(int* runtimeVersion) {
}
hipError_t hipCtxDestroy(hipCtx_t ctx) {
HIP_INIT_API(ctx);
HIP_INIT_API(hipCtxDestroy, ctx);
hipError_t e = hipSuccess;
ihipCtx_t* currentCtx = ihipGetTlsDefaultCtx();
ihipCtx_t* primaryCtx = ((ihipDevice_t*)ctx->getDevice())->_primaryCtx;
@@ -134,7 +134,7 @@ hipError_t hipCtxDestroy(hipCtx_t ctx) {
}
hipError_t hipCtxPopCurrent(hipCtx_t* ctx) {
HIP_INIT_API(ctx);
HIP_INIT_API(hipCtxPopCurrent, ctx);
hipError_t e = hipSuccess;
ihipCtx_t* currentCtx = ihipGetTlsDefaultCtx();
auto deviceHandle = currentCtx->getDevice();
@@ -155,7 +155,7 @@ hipError_t hipCtxPopCurrent(hipCtx_t* ctx) {
}
hipError_t hipCtxPushCurrent(hipCtx_t ctx) {
HIP_INIT_API(ctx);
HIP_INIT_API(hipCtxPushCurrent, ctx);
hipError_t e = hipSuccess;
if (ctx != NULL) { // TODO- is this check needed?
ihipSetTlsDefaultCtx(ctx);
@@ -168,7 +168,7 @@ hipError_t hipCtxPushCurrent(hipCtx_t ctx) {
}
hipError_t hipCtxGetCurrent(hipCtx_t* ctx) {
HIP_INIT_API(ctx);
HIP_INIT_API(hipCtxGetCurrent, ctx);
hipError_t e = hipSuccess;
if ((tls_getPrimaryCtx) || tls_ctxStack.empty()) {
*ctx = ihipGetTlsDefaultCtx();
@@ -179,7 +179,7 @@ hipError_t hipCtxGetCurrent(hipCtx_t* ctx) {
}
hipError_t hipCtxSetCurrent(hipCtx_t ctx) {
HIP_INIT_API(ctx);
HIP_INIT_API(hipCtxSetCurrent, ctx);
hipError_t e = hipSuccess;
if (ctx == NULL) {
tls_ctxStack.pop();
@@ -192,7 +192,7 @@ hipError_t hipCtxSetCurrent(hipCtx_t ctx) {
}
hipError_t hipCtxGetDevice(hipDevice_t* device) {
HIP_INIT_API(device);
HIP_INIT_API(hipCtxGetDevice, device);
hipError_t e = hipSuccess;
ihipCtx_t* ctx = ihipGetTlsDefaultCtx();
@@ -208,7 +208,7 @@ hipError_t hipCtxGetDevice(hipDevice_t* device) {
}
hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion) {
HIP_INIT_API(apiVersion);
HIP_INIT_API(hipCtxGetApiVersion, apiVersion);
if (apiVersion) {
*apiVersion = 4;
@@ -218,7 +218,7 @@ hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion) {
}
hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig) {
HIP_INIT_API(cacheConfig);
HIP_INIT_API(hipCtxGetCacheConfig, cacheConfig);
*cacheConfig = hipFuncCachePreferNone;
@@ -226,7 +226,7 @@ hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig) {
}
hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig) {
HIP_INIT_API(cacheConfig);
HIP_INIT_API(hipCtxSetCacheConfig, cacheConfig);
// Nop, AMD does not support variable cache configs.
@@ -234,7 +234,7 @@ hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig) {
}
hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config) {
HIP_INIT_API(config);
HIP_INIT_API(hipCtxSetSharedMemConfig, config);
// Nop, AMD does not support variable shared mem configs.
@@ -242,7 +242,7 @@ hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config) {
}
hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig* pConfig) {
HIP_INIT_API(pConfig);
HIP_INIT_API(hipCtxGetSharedMemConfig, pConfig);
*pConfig = hipSharedMemBankSizeFourByte;
@@ -250,12 +250,12 @@ hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig* pConfig) {
}
hipError_t hipCtxSynchronize(void) {
HIP_INIT_API(1);
HIP_INIT_API(hipCtxSynchronize, 1);
return ihipLogStatus(ihipSynchronize()); // TODP Shall check validity of ctx?
}
hipError_t hipCtxGetFlags(unsigned int* flags) {
HIP_INIT_API(flags);
HIP_INIT_API(hipCtxGetFlags, flags);
hipError_t e = hipSuccess;
ihipCtx_t* tempCtx;
tempCtx = ihipGetTlsDefaultCtx();
@@ -264,7 +264,7 @@ hipError_t hipCtxGetFlags(unsigned int* flags) {
}
hipError_t hipDevicePrimaryCtxGetState(hipDevice_t dev, unsigned int* flags, int* active) {
HIP_INIT_API(dev, flags, active);
HIP_INIT_API(hipDevicePrimaryCtxGetState, dev, flags, active);
hipError_t e = hipSuccess;
auto deviceHandle = ihipGetDevice(dev);
@@ -286,7 +286,7 @@ hipError_t hipDevicePrimaryCtxGetState(hipDevice_t dev, unsigned int* flags, int
}
hipError_t hipDevicePrimaryCtxRelease(hipDevice_t dev) {
HIP_INIT_API(dev);
HIP_INIT_API(hipDevicePrimaryCtxRelease, dev);
hipError_t e = hipSuccess;
auto deviceHandle = ihipGetDevice(dev);
@@ -297,7 +297,7 @@ hipError_t hipDevicePrimaryCtxRelease(hipDevice_t dev) {
}
hipError_t hipDevicePrimaryCtxRetain(hipCtx_t* pctx, hipDevice_t dev) {
HIP_INIT_API(pctx, dev);
HIP_INIT_API(hipDevicePrimaryCtxRetain, pctx, dev);
hipError_t e = hipSuccess;
auto deviceHandle = ihipGetDevice(dev);
@@ -309,7 +309,7 @@ hipError_t hipDevicePrimaryCtxRetain(hipCtx_t* pctx, hipDevice_t dev) {
}
hipError_t hipDevicePrimaryCtxReset(hipDevice_t dev) {
HIP_INIT_API(dev);
HIP_INIT_API(hipDevicePrimaryCtxReset, dev);
hipError_t e = hipSuccess;
auto deviceHandle = ihipGetDevice(dev);
@@ -322,7 +322,7 @@ hipError_t hipDevicePrimaryCtxReset(hipDevice_t dev) {
}
hipError_t hipDevicePrimaryCtxSetFlags(hipDevice_t dev, unsigned int flags) {
HIP_INIT_API(dev, flags);
HIP_INIT_API(hipDevicePrimaryCtxSetFlags, dev, flags);
hipError_t e = hipSuccess;
auto deviceHandle = ihipGetDevice(dev);
projects/hip/src/hip_device.cpp
+20 -20
Просмотреть файл
@@ -30,7 +30,7 @@ THE SOFTWARE.
//-------------------------------------------------------------------------------------------------
// TODO - does this initialize HIP runtime?
hipError_t hipGetDevice(int* deviceId) {
HIP_INIT_API(deviceId);
HIP_INIT_API(hipGetDevice, deviceId);
hipError_t e = hipSuccess;
@@ -69,12 +69,12 @@ hipError_t ihipGetDeviceCount(int* count) {
}
hipError_t hipGetDeviceCount(int* count) {
HIP_INIT_API(count);
HIP_INIT_API(hipGetDeviceCount, count);
return ihipLogStatus(ihipGetDeviceCount(count));
}
hipError_t hipDeviceSetCacheConfig(hipFuncCache_t cacheConfig) {
HIP_INIT_API(cacheConfig);
HIP_INIT_API(hipDeviceSetCacheConfig, cacheConfig);
// Nop, AMD does not support variable cache configs.
@@ -82,7 +82,7 @@ hipError_t hipDeviceSetCacheConfig(hipFuncCache_t cacheConfig) {
}
hipError_t hipDeviceGetCacheConfig(hipFuncCache_t* cacheConfig) {
HIP_INIT_API(cacheConfig);
HIP_INIT_API(hipDeviceGetCacheConfig, cacheConfig);
if (cacheConfig == nullptr) {
return ihipLogStatus(hipErrorInvalidValue);
@@ -94,7 +94,7 @@ hipError_t hipDeviceGetCacheConfig(hipFuncCache_t* cacheConfig) {
}
hipError_t hipDeviceGetLimit(size_t* pValue, hipLimit_t limit) {
HIP_INIT_API(pValue, limit);
HIP_INIT_API(hipDeviceGetLimit, pValue, limit);
if (pValue == nullptr) {
return ihipLogStatus(hipErrorInvalidValue);
}
@@ -107,7 +107,7 @@ hipError_t hipDeviceGetLimit(size_t* pValue, hipLimit_t limit) {
}
hipError_t hipFuncSetCacheConfig(const void* func, hipFuncCache_t cacheConfig) {
HIP_INIT_API(cacheConfig);
HIP_INIT_API(hipFuncSetCacheConfig, cacheConfig);
// Nop, AMD does not support variable cache configs.
@@ -115,7 +115,7 @@ hipError_t hipFuncSetCacheConfig(const void* func, hipFuncCache_t cacheConfig) {
}
hipError_t hipDeviceSetSharedMemConfig(hipSharedMemConfig config) {
HIP_INIT_API(config);
HIP_INIT_API(hipDeviceSetSharedMemConfig, config);
// Nop, AMD does not support variable shared mem configs.
@@ -123,7 +123,7 @@ hipError_t hipDeviceSetSharedMemConfig(hipSharedMemConfig config) {
}
hipError_t hipDeviceGetSharedMemConfig(hipSharedMemConfig* pConfig) {
HIP_INIT_API(pConfig);
HIP_INIT_API(hipDeviceGetSharedMemConfig, pConfig);
*pConfig = hipSharedMemBankSizeFourByte;
@@ -131,7 +131,7 @@ hipError_t hipDeviceGetSharedMemConfig(hipSharedMemConfig* pConfig) {
}
hipError_t hipSetDevice(int deviceId) {
HIP_INIT_API(deviceId);
HIP_INIT_API(hipSetDevice, deviceId);
if ((deviceId < 0) || (deviceId >= g_deviceCnt)) {
return ihipLogStatus(hipErrorInvalidDevice);
} else {
@@ -142,12 +142,12 @@ hipError_t hipSetDevice(int deviceId) {
}
hipError_t hipDeviceSynchronize(void) {
HIP_INIT_SPECIAL_API(TRACE_SYNC);
HIP_INIT_SPECIAL_API(hipDeviceSynchronize, TRACE_SYNC);
return ihipLogStatus(ihipSynchronize());
}
hipError_t hipDeviceReset(void) {
HIP_INIT_API();
HIP_INIT_API(hipDeviceReset, );
auto* ctx = ihipGetTlsDefaultCtx();
@@ -287,7 +287,7 @@ hipError_t ihipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device
}
hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) {
HIP_INIT_API(pi, attr, device);
HIP_INIT_API(hipDeviceGetAttribute, pi, attr, device);
if ((device < 0) || (device >= g_deviceCnt)) {
return ihipLogStatus(hipErrorInvalidDevice);
}
@@ -314,7 +314,7 @@ hipError_t ihipGetDeviceProperties(hipDeviceProp_t* props, int device) {
}
hipError_t hipGetDeviceProperties(hipDeviceProp_t* props, int device) {
HIP_INIT_API(props, device);
HIP_INIT_API(hipGetDeviceProperties, props, device);
if ((device < 0) || (device >= g_deviceCnt)) {
return ihipLogStatus(hipErrorInvalidDevice);
}
@@ -322,7 +322,7 @@ hipError_t hipGetDeviceProperties(hipDeviceProp_t* props, int device) {
}
hipError_t hipSetDeviceFlags(unsigned int flags) {
HIP_INIT_API(flags);
HIP_INIT_API(hipSetDeviceFlags, flags);
hipError_t e = hipSuccess;
@@ -367,7 +367,7 @@ hipError_t hipSetDeviceFlags(unsigned int flags) {
};
hipError_t hipDeviceComputeCapability(int* major, int* minor, hipDevice_t device) {
HIP_INIT_API(major, minor, device);
HIP_INIT_API(hipDeviceComputeCapability, major, minor, device);
hipError_t e = hipSuccess;
if ((device < 0) || (device >= g_deviceCnt)) {
e = hipErrorInvalidDevice;
@@ -380,7 +380,7 @@ hipError_t hipDeviceComputeCapability(int* major, int* minor, hipDevice_t device
hipError_t hipDeviceGetName(char* name, int len, hipDevice_t device) {
// Cast to void* here to avoid printing garbage in debug modes.
HIP_INIT_API((void*)name, len, device);
HIP_INIT_API(hipDeviceGetName, (void*)name, len, device);
hipError_t e = hipSuccess;
if ((device < 0) || (device >= g_deviceCnt)) {
e = hipErrorInvalidDevice;
@@ -394,7 +394,7 @@ hipError_t hipDeviceGetName(char* name, int len, hipDevice_t device) {
hipError_t hipDeviceGetPCIBusId(char* pciBusId, int len, int device) {
// Cast to void* here to avoid printing garbage in debug modes.
HIP_INIT_API((void*)pciBusId, len, device);
HIP_INIT_API(hipDeviceGetPCIBusId, (void*)pciBusId, len, device);
hipError_t e = hipErrorInvalidValue;
if ((device < 0) || (device >= g_deviceCnt)) {
e = hipErrorInvalidDevice;
@@ -413,7 +413,7 @@ hipError_t hipDeviceGetPCIBusId(char* pciBusId, int len, int device) {
}
hipError_t hipDeviceTotalMem(size_t* bytes, hipDevice_t device) {
HIP_INIT_API(bytes, device);
HIP_INIT_API(hipDeviceTotalMem, bytes, device);
hipError_t e = hipSuccess;
if ((device < 0) || (device >= g_deviceCnt)) {
e = hipErrorInvalidDevice;
@@ -425,7 +425,7 @@ hipError_t hipDeviceTotalMem(size_t* bytes, hipDevice_t device) {
}
hipError_t hipDeviceGetByPCIBusId(int* device, const char* pciBusId) {
HIP_INIT_API(device, pciBusId);
HIP_INIT_API(hipDeviceGetByPCIBusId, device, pciBusId);
hipDeviceProp_t tempProp;
int deviceCount = 0;
hipError_t e = hipErrorInvalidValue;
@@ -451,7 +451,7 @@ hipError_t hipDeviceGetByPCIBusId(int* device, const char* pciBusId) {
}
hipError_t hipChooseDevice(int* device, const hipDeviceProp_t* prop) {
HIP_INIT_API(device, prop);
HIP_INIT_API(hipChooseDevice, device, prop);
hipDeviceProp_t tempProp;
hipError_t e = hipSuccess;
if ((device == NULL) || (prop == NULL)) {
projects/hip/src/hip_error.cpp
+4 -4
Просмотреть файл
@@ -30,7 +30,7 @@ THE SOFTWARE.
//---
hipError_t hipGetLastError() {
HIP_INIT_API();
HIP_INIT_API(hipGetLastError);
// Return last error, but then reset the state:
hipError_t e = ihipLogStatus(tls_lastHipError);
@@ -39,20 +39,20 @@ hipError_t hipGetLastError() {
}
hipError_t hipPeekAtLastError() {
HIP_INIT_API();
HIP_INIT_API(hipPeekAtLastError);
// peek at last error, but don't reset it.
return ihipLogStatus(tls_lastHipError);
}
const char* hipGetErrorName(hipError_t hip_error) {
HIP_INIT_API(hip_error);
HIP_INIT_API(hipGetErrorName, hip_error);
return ihipErrorString(hip_error);
}
const char* hipGetErrorString(hipError_t hip_error) {
HIP_INIT_API(hip_error);
HIP_INIT_API(hipGetErrorString, hip_error);
// TODO - return a message explaining the error.
// TODO - This should be set up to return the same string reported in the the doxygen comments,
projects/hip/src/hip_event.cpp
+7 -7
Просмотреть файл
@@ -95,20 +95,20 @@ hipError_t ihipEventCreate(hipEvent_t* event, unsigned flags) {
}
hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags) {
HIP_INIT_API(event, flags);
HIP_INIT_API(hipEventCreateWithFlags, event, flags);
return ihipLogStatus(ihipEventCreate(event, flags));
}
hipError_t hipEventCreate(hipEvent_t* event) {
HIP_INIT_API(event);
HIP_INIT_API(hipEventCreate, event);
return ihipLogStatus(ihipEventCreate(event, 0));
}
hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) {
HIP_INIT_SPECIAL_API(TRACE_SYNC, event, stream);
HIP_INIT_SPECIAL_API(hipEventRecord, TRACE_SYNC, event, stream);
auto ecd = event->locked_copyCrit();
@@ -153,7 +153,7 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) {
hipError_t hipEventDestroy(hipEvent_t event) {
HIP_INIT_API(event);
HIP_INIT_API(hipEventDestroy, event);
if (event) {
delete event;
@@ -165,7 +165,7 @@ hipError_t hipEventDestroy(hipEvent_t event) {
}
hipError_t hipEventSynchronize(hipEvent_t event) {
HIP_INIT_SPECIAL_API(TRACE_SYNC, event);
HIP_INIT_SPECIAL_API(hipEventSynchronize, TRACE_SYNC, event);
if (!(event->_flags & hipEventReleaseToSystem)) {
tprintf(DB_WARN,
@@ -198,7 +198,7 @@ hipError_t hipEventSynchronize(hipEvent_t event) {
}
hipError_t hipEventElapsedTime(float* ms, hipEvent_t start, hipEvent_t stop) {
HIP_INIT_API(ms, start, stop);
HIP_INIT_API(hipEventElapsedTime, ms, start, stop);
hipError_t status = hipSuccess;
@@ -255,7 +255,7 @@ hipError_t hipEventElapsedTime(float* ms, hipEvent_t start, hipEvent_t stop) {
}
hipError_t hipEventQuery(hipEvent_t event) {
HIP_INIT_SPECIAL_API(TRACE_QUERY, event);
HIP_INIT_SPECIAL_API(hipEventQuery, TRACE_QUERY, event);
if (!(event->_flags & hipEventReleaseToSystem)) {
tprintf(DB_WARN,
projects/hip/src/hip_hcc.cpp
+7 -6
Просмотреть файл
@@ -179,7 +179,7 @@ uint64_t recordApiTrace(std::string* fullStr, const std::string& apiStr) {
if (COMPILE_HIP_DB && HIP_TRACE_API) {
fprintf(stderr, "%s<<hip-api tid:%s @%lu%s\n", API_COLOR, fullStr->c_str(), apiStartTick,
fprintf(stderr, "%s<<hip-api pid:%d tid:%s @%lu%s\n", API_COLOR, tls_tidInfo.pid(), fullStr->c_str(), apiStartTick,
API_COLOR_END);
}
@@ -237,6 +237,7 @@ hipError_t ihipSynchronize(void) {
//=================================================================================================
TidInfo::TidInfo() : _apiSeqNum(0) {
_shortTid = g_lastShortTid.fetch_add(1);
_pid = getpid();
if (COMPILE_HIP_DB && HIP_TRACE_API) {
std::stringstream tid_ss;
@@ -1535,7 +1536,7 @@ void ihipPrintKernelLaunch(const char* kernelName, const grid_launch_parm* lp,
if ((HIP_TRACE_API & (1 << TRACE_KCMD)) || HIP_PROFILE_API ||
(COMPILE_HIP_DB & HIP_TRACE_API)) {
std::stringstream os;
os << tls_tidInfo.tid() << "." << tls_tidInfo.apiSeqNum() << " hipLaunchKernel '"
os << tls_tidInfo.pid() << " " << tls_tidInfo.tid() << "." << tls_tidInfo.apiSeqNum() << " hipLaunchKernel '"
<< kernelName << "'"
<< " gridDim:" << lp->grid_dim << " groupDim:" << lp->group_dim << " sharedMem:+"
<< lp->dynamic_group_mem_bytes << " " << *stream;
@@ -2288,7 +2289,7 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes
//-------------------------------------------------------------------------------------------------
// Profiler, really these should live elsewhere:
hipError_t hipProfilerStart() {
HIP_INIT_API();
HIP_INIT_API(hipProfilerStart);
#if COMPILE_HIP_ATP_MARKER
amdtResumeProfiling(AMDT_ALL_PROFILING);
#endif
@@ -2298,7 +2299,7 @@ hipError_t hipProfilerStart() {
hipError_t hipProfilerStop() {
HIP_INIT_API();
HIP_INIT_API(hipProfilerStop);
#if COMPILE_HIP_ATP_MARKER
amdtStopProfiling(AMDT_ALL_PROFILING);
#endif
@@ -2313,7 +2314,7 @@ hipError_t hipProfilerStop() {
//---
hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator* acc) {
HIP_INIT_API(deviceId, acc);
HIP_INIT_API(hipHccGetAccelerator, deviceId, acc);
const ihipDevice_t* device = ihipGetDevice(deviceId);
hipError_t err;
@@ -2329,7 +2330,7 @@ hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator* acc) {
//---
hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** av) {
HIP_INIT_API(stream, av);
HIP_INIT_API(hipHccGetAcceleratorView, stream, av);
if (stream == hipStreamNull) {
ihipCtx_t* device = ihipGetTlsDefaultCtx();
projects/hip/src/hip_hcc_internal.h
+32 -28
Просмотреть файл
@@ -94,11 +94,13 @@ class TidInfo {
TidInfo();
int tid() const { return _shortTid; };
pid_t pid() const { return _pid; };
uint64_t incApiSeqNum() { return ++_apiSeqNum; };
uint64_t apiSeqNum() const { return _apiSeqNum; };
private:
int _shortTid;
pid_t _pid;
// monotonically increasing API sequence number for this threa.
uint64_t _apiSeqNum;
@@ -243,14 +245,14 @@ static const DbName dbName[] = {
#if COMPILE_HIP_DB
#define tprintf(trace_level, ...) \
{ \
if (HIP_DB & (1 << (trace_level))) { \
char msgStr[1000]; \
snprintf(msgStr, sizeof(msgStr), __VA_ARGS__); \
fprintf(stderr, " %ship-%s tid:%d:%s%s", dbName[trace_level]._color, \
dbName[trace_level]._shortName, tls_tidInfo.tid(), msgStr, KNRM); \
} \
#define tprintf(trace_level, ...) \
{ \
if (HIP_DB & (1 << (trace_level))) { \
char msgStr[1000]; \
snprintf(msgStr, sizeof(msgStr), __VA_ARGS__); \
fprintf(stderr, " %ship-%s pid:%d tid:%d:%s%s", dbName[trace_level]._color, \
dbName[trace_level]._shortName, tls_tidInfo.pid(), tls_tidInfo.tid(), msgStr, KNRM); \
} \
}
#else
/* Compile to empty code */
@@ -297,38 +299,40 @@ extern uint64_t recordApiTrace(std::string* fullStr, const std::string& apiStr);
// This macro should be called at the beginning of every HIP API.
// It initializes the hip runtime (exactly once), and
// generates a trace string that can be output to stderr or to ATP file.
#define HIP_INIT_API(...) \
#define HIP_INIT_API(cid, ...) \
HIP_INIT() \
API_TRACE(0, __VA_ARGS__);
API_TRACE(0, __VA_ARGS__); \
HIP_CB_SPAWNER_OBJECT(cid);
// Like above, but will trace with a specified "special" bit.
// Replace HIP_INIT_API with this call inside HIP APIs that launch work on the GPU:
// kernel launches, copy commands, memory sets, etc.
#define HIP_INIT_SPECIAL_API(tbit, ...) \
#define HIP_INIT_SPECIAL_API(cid, tbit, ...) \
HIP_INIT() \
API_TRACE((HIP_TRACE_API & (1 << tbit)), __VA_ARGS__);
API_TRACE((HIP_TRACE_API & (1 << tbit)), __VA_ARGS__); \
HIP_CB_SPAWNER_OBJECT(cid);
// This macro should be called at the end of every HIP API, and only at the end of top-level hip
// APIS (not internal hip) It has dual function: logs the last error returned for use by
// hipGetLastError, and also prints the closing message when the debug trace is enabled.
#define ihipLogStatus(hipStatus) \
({ \
hipError_t localHipStatus = hipStatus; /*local copy so hipStatus only evaluated once*/ \
tls_lastHipError = localHipStatus; \
\
if ((COMPILE_HIP_TRACE_API & 0x2) && HIP_TRACE_API & (1 << TRACE_ALL)) { \
auto ticks = getTicks() - hipApiStartTick; \
fprintf(stderr, " %ship-api tid:%d.%lu %-30s ret=%2d (%s)>> +%lu ns%s\n", \
(localHipStatus == 0) ? API_COLOR : KRED, tls_tidInfo.tid(), \
tls_tidInfo.apiSeqNum(), __func__, localHipStatus, \
ihipErrorString(localHipStatus), ticks, API_COLOR_END); \
} \
if (HIP_PROFILE_API) { \
MARKER_END(); \
} \
localHipStatus; \
#define ihipLogStatus(hipStatus) \
({ \
hipError_t localHipStatus = hipStatus; /*local copy so hipStatus only evaluated once*/ \
tls_lastHipError = localHipStatus; \
\
if ((COMPILE_HIP_TRACE_API & 0x2) && HIP_TRACE_API & (1 << TRACE_ALL)) { \
auto ticks = getTicks() - hipApiStartTick; \
fprintf(stderr, " %ship-api pid:%d tid:%d.%lu %-30s ret=%2d (%s)>> +%lu ns%s\n", \
(localHipStatus == 0) ? API_COLOR : KRED, tls_tidInfo.pid(), tls_tidInfo.tid(), \
tls_tidInfo.apiSeqNum(), __func__, localHipStatus, \
ihipErrorString(localHipStatus), ticks, API_COLOR_END); \
} \
if (HIP_PROFILE_API) { \
MARKER_END(); \
} \
localHipStatus; \
})
projects/hip/src/hip_intercept.cpp
+49
Просмотреть файл
@@ -0,0 +1,49 @@
/*
Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "hip/hip_runtime.h"
#include "hip/hcc_detail/hip_prof_api.h"
// HIP API callback/activity
api_callbacks_table_t callbacks_table;
extern std::string& FunctionSymbol(hipFunction_t f);
const char* hipKernelNameRef(const hipFunction_t f) { return FunctionSymbol(f).c_str(); }
hipError_t hipRegisterApiCallback(uint32_t id, void* fun, void* arg) {
return callbacks_table.set_callback(id, reinterpret_cast<api_callbacks_table_t::fun_t>(fun), arg) ?
hipSuccess : hipErrorInvalidValue;
}
hipError_t hipRemoveApiCallback(uint32_t id) {
return callbacks_table.set_callback(id, NULL, NULL) ? hipSuccess : hipErrorInvalidValue;
}
hipError_t hipRegisterActivityCallback(uint32_t id, void* fun, void* arg) {
return callbacks_table.set_activity(id, reinterpret_cast<api_callbacks_table_t::act_t>(fun), arg) ?
hipSuccess : hipErrorInvalidValue;
}
hipError_t hipRemoveActivityCallback(uint32_t id) {
return callbacks_table.set_activity(id, NULL, NULL) ? hipSuccess : hipErrorInvalidValue;
}
projects/hip/src/hip_memory.cpp
+51 -51
Просмотреть файл
@@ -159,7 +159,7 @@ void* allocAndSharePtr(const char* msg, size_t sizeBytes, ihipCtx_t* ctx, bool s
// TODO - add more info here when available.
//
hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void* ptr) {
HIP_INIT_API(attributes, ptr);
HIP_INIT_API(hipPointerGetAttributes, attributes, ptr);
hipError_t e = hipSuccess;
if ((attributes == nullptr) || (ptr == nullptr)) {
@@ -206,7 +206,7 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void
hipError_t hipHostGetDevicePointer(void** devicePointer, void* hostPointer, unsigned flags) {
HIP_INIT_API(devicePointer, hostPointer, flags);
HIP_INIT_API(hipHostGetDevicePointer, devicePointer, hostPointer, flags);
hipError_t e = hipSuccess;
@@ -237,7 +237,7 @@ hipError_t hipHostGetDevicePointer(void** devicePointer, void* hostPointer, unsi
hipError_t hipMalloc(void** ptr, size_t sizeBytes) {
HIP_INIT_SPECIAL_API((TRACE_MEM), ptr, sizeBytes);
HIP_INIT_SPECIAL_API(hipMalloc, (TRACE_MEM), ptr, sizeBytes);
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
@@ -266,7 +266,7 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) {
hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
HIP_INIT_SPECIAL_API((TRACE_MEM), ptr, sizeBytes, flags);
HIP_INIT_SPECIAL_API(hipHostMalloc, (TRACE_MEM), ptr, sizeBytes, flags);
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
@@ -398,7 +398,7 @@ hipError_t ihipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t heigh
// width in bytes
hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height) {
HIP_INIT_SPECIAL_API((TRACE_MEM), ptr, pitch, width, height);
HIP_INIT_SPECIAL_API(hipMallocPitch, (TRACE_MEM), ptr, pitch, width, height);
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
@@ -409,7 +409,7 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height
}
hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent) {
HIP_INIT_API(pitchedDevPtr, &extent);
HIP_INIT_API(hipMalloc3D, pitchedDevPtr, &extent);
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
@@ -444,7 +444,7 @@ extern void getChannelOrderAndType(const hipChannelFormatDesc& desc,
hsa_ext_image_channel_type_t* channelType);
hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray) {
HIP_INIT_SPECIAL_API((TRACE_MEM), array, pAllocateArray);
HIP_INIT_SPECIAL_API(hipArrayCreate, (TRACE_MEM), array, pAllocateArray);
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
if (pAllocateArray->width > 0) {
@@ -554,7 +554,7 @@ hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocat
hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, size_t width,
size_t height, unsigned int flags) {
HIP_INIT_SPECIAL_API((TRACE_MEM), array, desc, width, height, flags);
HIP_INIT_SPECIAL_API(hipMallocArray, (TRACE_MEM), array, desc, width, height, flags);
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
if (width > 0) {
@@ -635,7 +635,7 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, si
}
hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray) {
HIP_INIT_SPECIAL_API((TRACE_MEM), array, pAllocateArray);
HIP_INIT_SPECIAL_API(hipArray3DCreate, (TRACE_MEM), array, pAllocateArray);
hipError_t hip_status = hipSuccess;
auto ctx = ihipGetTlsDefaultCtx();
@@ -762,7 +762,7 @@ hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc*
HIP_INIT_API(array, desc, &extent, flags);
HIP_INIT_API(hipMalloc3DArray, array, desc, &extent, flags);
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
@@ -846,7 +846,7 @@ hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc*
}
hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) {
HIP_INIT_API(flagsPtr, hostPtr);
HIP_INIT_API(hipHostGetFlags, flagsPtr, hostPtr);
hipError_t hip_status = hipSuccess;
@@ -874,7 +874,7 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) {
// TODO - need to fix several issues here related to P2P access, host memory fallback.
hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags) {
HIP_INIT_API(hostPtr, sizeBytes, flags);
HIP_INIT_API(hipHostRegister, hostPtr, sizeBytes, flags);
hipError_t hip_status = hipSuccess;
@@ -939,7 +939,7 @@ hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags)
}
hipError_t hipHostUnregister(void* hostPtr) {
HIP_INIT_API(hostPtr);
HIP_INIT_API(hipHostUnregister, hostPtr);
auto ctx = ihipGetTlsDefaultCtx();
hipError_t hip_status = hipSuccess;
if (hostPtr == NULL) {
@@ -974,7 +974,7 @@ inline hipDeviceptr_t agent_address_for_symbol(const char* symbolName) {
hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, size_t count, size_t offset,
hipMemcpyKind kind) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), symbolName, src, count, offset, kind);
HIP_INIT_SPECIAL_API(hipMemcpyToSymbol, (TRACE_MCMD), symbolName, src, count, offset, kind);
if (symbolName == nullptr) {
return ihipLogStatus(hipErrorInvalidSymbol);
@@ -1006,7 +1006,7 @@ hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, size_t cou
hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName, size_t count, size_t offset,
hipMemcpyKind kind) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), symbolName, dst, count, offset, kind);
HIP_INIT_SPECIAL_API(hipMemcpyFromSymbol, (TRACE_MCMD), symbolName, dst, count, offset, kind);
if (symbolName == nullptr) {
return ihipLogStatus(hipErrorInvalidSymbol);
@@ -1038,7 +1038,7 @@ hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName, size_t count,
hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src, size_t count,
size_t offset, hipMemcpyKind kind, hipStream_t stream) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), symbolName, src, count, offset, kind, stream);
HIP_INIT_SPECIAL_API(hipMemcpyToSymbolAsync, (TRACE_MCMD), symbolName, src, count, offset, kind, stream);
if (symbolName == nullptr) {
return ihipLogStatus(hipErrorInvalidSymbol);
@@ -1073,7 +1073,7 @@ hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src, size_
hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t count, size_t offset,
hipMemcpyKind kind, hipStream_t stream) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), symbolName, dst, count, offset, kind, stream);
HIP_INIT_SPECIAL_API(hipMemcpyFromSymbolAsync, (TRACE_MCMD), symbolName, dst, count, offset, kind, stream);
if (symbolName == nullptr) {
return ihipLogStatus(hipErrorInvalidSymbol);
@@ -1125,7 +1125,7 @@ hipError_t hipGetSymbolSize(size_t* size, const void* symbolName) {
//---
hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, kind);
HIP_INIT_SPECIAL_API(hipMemcpy, (TRACE_MCMD), dst, src, sizeBytes, kind);
hipError_t e = hipSuccess;
@@ -1152,7 +1152,7 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind
hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes);
HIP_INIT_SPECIAL_API(hipMemcpyHtoD, (TRACE_MCMD), dst, src, sizeBytes);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
@@ -1171,7 +1171,7 @@ hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) {
hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes);
HIP_INIT_SPECIAL_API(hipMemcpyDtoH, (TRACE_MCMD), dst, src, sizeBytes);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
@@ -1190,7 +1190,7 @@ hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) {
hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes);
HIP_INIT_SPECIAL_API(hipMemcpyDtoD, (TRACE_MCMD), dst, src, sizeBytes);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
@@ -1208,7 +1208,7 @@ hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeByte
}
hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes);
HIP_INIT_SPECIAL_API(hipMemcpyHtoH, (TRACE_MCMD), dst, src, sizeBytes);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
@@ -1227,13 +1227,13 @@ hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) {
hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind,
hipStream_t stream) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, kind, stream);
HIP_INIT_SPECIAL_API(hipMemcpyAsync, (TRACE_MCMD), dst, src, sizeBytes, kind, stream);
return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, kind, stream));
}
hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, hipStream_t stream) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, stream);
HIP_INIT_SPECIAL_API(hipMemcpyHtoDAsync, (TRACE_MCMD), dst, src, sizeBytes, stream);
return ihipLogStatus(
hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyHostToDevice, stream));
@@ -1241,14 +1241,14 @@ hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, h
hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes,
hipStream_t stream) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, stream);
HIP_INIT_SPECIAL_API(hipMemcpyDtoDAsync, (TRACE_MCMD), dst, src, sizeBytes, stream);
return ihipLogStatus(
hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyDeviceToDevice, stream));
}
hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, stream);
HIP_INIT_SPECIAL_API(hipMemcpyDtoHAsync, (TRACE_MCMD), dst, src, sizeBytes, stream);
return ihipLogStatus(
hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyDeviceToHost, stream));
@@ -1256,7 +1256,7 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, h
hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src,
size_t spitch, size_t width, size_t height, hipMemcpyKind kind) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, wOffset, hOffset, src, spitch, width, height, kind);
HIP_INIT_SPECIAL_API(hipMemcpy2DToArray, (TRACE_MCMD), dst, wOffset, hOffset, src, spitch, width, height, kind);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
@@ -1308,7 +1308,7 @@ hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, con
hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src,
size_t count, hipMemcpyKind kind) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, wOffset, hOffset, src, count, kind);
HIP_INIT_SPECIAL_API(hipMemcpyToArray, (TRACE_MCMD), dst, wOffset, hOffset, src, count, kind);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
@@ -1327,7 +1327,7 @@ hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, const
hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffset, size_t hOffset,
size_t count, hipMemcpyKind kind) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, srcArray, wOffset, hOffset, count, kind);
HIP_INIT_SPECIAL_API(hipMemcpyFromArray, (TRACE_MCMD), dst, srcArray, wOffset, hOffset, count, kind);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
@@ -1345,7 +1345,7 @@ hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffs
}
hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHost, size_t count) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dstArray, dstOffset, srcHost, count);
HIP_INIT_SPECIAL_API(hipMemcpyHtoA, (TRACE_MCMD), dstArray, dstOffset, srcHost, count);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
@@ -1363,7 +1363,7 @@ hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHo
}
hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t count) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, srcArray, srcOffset, count);
HIP_INIT_SPECIAL_API(hipMemcpyAtoH, (TRACE_MCMD), dst, srcArray, srcOffset, count);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
@@ -1382,7 +1382,7 @@ hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t
}
hipError_t hipMemcpy3D(const struct hipMemcpy3DParms* p) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), p);
HIP_INIT_SPECIAL_API(hipMemcpy3D, (TRACE_MCMD), p);
hipError_t e = hipSuccess;
if (p) {
size_t byteSize;
@@ -1650,7 +1650,7 @@ hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch
hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
size_t height, hipMemcpyKind kind) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind);
HIP_INIT_SPECIAL_API(hipMemcpy2D, (TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind);
hipError_t e = hipSuccess;
e = ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind);
return ihipLogStatus(e);
@@ -1658,7 +1658,7 @@ hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch,
hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
size_t height, hipMemcpyKind kind, hipStream_t stream) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind, stream);
HIP_INIT_SPECIAL_API(hipMemcpy2DAsync, (TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind, stream);
if (dst == nullptr || src == nullptr || width > dpitch || width > spitch) return ihipLogStatus(hipErrorInvalidValue);
hipError_t e = hipSuccess;
int isLocked = 0;
@@ -1697,7 +1697,7 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp
}
hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), pCopy);
HIP_INIT_SPECIAL_API(hipMemcpyParam2D, (TRACE_MCMD), pCopy);
hipError_t e = hipSuccess;
if (pCopy == nullptr) {
e = hipErrorInvalidValue;
@@ -1709,7 +1709,7 @@ hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) {
// TODO-sync: function is async unless target is pinned host memory - then these are fully sync.
hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, value, sizeBytes, stream);
HIP_INIT_SPECIAL_API(hipMemsetAsync, (TRACE_MCMD), dst, value, sizeBytes, stream);
hipError_t e = hipSuccess;
@@ -1721,7 +1721,7 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t st
};
hipError_t hipMemset(void* dst, int value, size_t sizeBytes) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, value, sizeBytes);
HIP_INIT_SPECIAL_API(hipMemset, (TRACE_MCMD), dst, value, sizeBytes);
hipError_t e = hipSuccess;
@@ -1737,7 +1737,7 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes) {
}
hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t height) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, pitch, value, width, height);
HIP_INIT_SPECIAL_API(hipMemset2D, (TRACE_MCMD), dst, pitch, value, width, height);
hipError_t e = hipSuccess;
@@ -1756,7 +1756,7 @@ hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t
hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, size_t width, size_t height, hipStream_t stream )
{
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, pitch, value, width, height, stream);
HIP_INIT_SPECIAL_API(hipMemset2DAsync, (TRACE_MCMD), dst, pitch, value, width, height, stream);
hipError_t e = hipSuccess;
@@ -1773,7 +1773,7 @@ hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, size_t width, si
};
hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, value, sizeBytes);
HIP_INIT_SPECIAL_API(hipMemsetD8, (TRACE_MCMD), dst, value, sizeBytes);
hipError_t e = hipSuccess;
@@ -1790,7 +1790,7 @@ hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes
hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent )
{
HIP_INIT_SPECIAL_API((TRACE_MCMD), &pitchedDevPtr, value, &extent);
HIP_INIT_SPECIAL_API(hipMemset3D, (TRACE_MCMD), &pitchedDevPtr, value, &extent);
hipError_t e = hipSuccess;
hipStream_t stream = hipStreamNull;
@@ -1809,7 +1809,7 @@ hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent
hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent ,hipStream_t stream )
{
HIP_INIT_SPECIAL_API((TRACE_MCMD), &pitchedDevPtr, value, &extent);
HIP_INIT_SPECIAL_API(hipMemset3DAsync, (TRACE_MCMD), &pitchedDevPtr, value, &extent);
hipError_t e = hipSuccess;
// TODO - call an ihip memset so HIP_TRACE is correct.
@@ -1825,7 +1825,7 @@ hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int value, hipExtent e
}
hipError_t hipMemGetInfo(size_t* free, size_t* total) {
HIP_INIT_API(free, total);
HIP_INIT_API(hipMemGetInfo, free, total);
hipError_t e = hipSuccess;
@@ -1859,7 +1859,7 @@ hipError_t hipMemGetInfo(size_t* free, size_t* total) {
}
hipError_t hipMemPtrGetInfo(void* ptr, size_t* size) {
HIP_INIT_API(ptr, size);
HIP_INIT_API(hipMemPtrGetInfo, ptr, size);
hipError_t e = hipSuccess;
@@ -1884,7 +1884,7 @@ hipError_t hipMemPtrGetInfo(void* ptr, size_t* size) {
hipError_t hipFree(void* ptr) {
HIP_INIT_SPECIAL_API((TRACE_MEM), ptr);
HIP_INIT_SPECIAL_API(hipFree, (TRACE_MEM), ptr);
hipError_t hipStatus = hipErrorInvalidDevicePointer;
@@ -1916,7 +1916,7 @@ hipError_t hipFree(void* ptr) {
hipError_t hipHostFree(void* ptr) {
HIP_INIT_SPECIAL_API((TRACE_MEM), ptr);
HIP_INIT_SPECIAL_API(hipHostFree, (TRACE_MEM), ptr);
// Synchronize to ensure all work has finished.
ihipGetTlsDefaultCtx()->locked_waitAllStreams(); // ignores non-blocking streams, this waits
@@ -1951,7 +1951,7 @@ hipError_t hipHostFree(void* ptr) {
hipError_t hipFreeHost(void* ptr) { return hipHostFree(ptr); }
hipError_t hipFreeArray(hipArray* array) {
HIP_INIT_SPECIAL_API((TRACE_MEM), array);
HIP_INIT_SPECIAL_API(hipFreeArray, (TRACE_MEM), array);
hipError_t hipStatus = hipErrorInvalidDevicePointer;
@@ -1979,7 +1979,7 @@ hipError_t hipFreeArray(hipArray* array) {
}
hipError_t hipMemGetAddressRange(hipDeviceptr_t* pbase, size_t* psize, hipDeviceptr_t dptr) {
HIP_INIT_API(pbase, psize, dptr);
HIP_INIT_API(hipMemGetAddressRange, pbase, psize, dptr);
hipError_t hipStatus = hipSuccess;
hc::accelerator acc;
#if (__hcc_workweek__ >= 17332)
@@ -2000,7 +2000,7 @@ hipError_t hipMemGetAddressRange(hipDeviceptr_t* pbase, size_t* psize, hipDevice
// TODO: IPC implementaiton:
hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr) {
HIP_INIT_API(handle, devPtr);
HIP_INIT_API(hipIpcGetMemHandle, handle, devPtr);
hipError_t hipStatus = hipSuccess;
// Get the size of allocated pointer
size_t psize = 0u;
@@ -2036,7 +2036,7 @@ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr) {
}
hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags) {
HIP_INIT_API(devPtr, &handle, flags);
HIP_INIT_API(hipIpcOpenMemHandle, devPtr, &handle, flags);
hipError_t hipStatus = hipSuccess;
if (devPtr == NULL) {
hipStatus = hipErrorInvalidValue;
@@ -2066,7 +2066,7 @@ hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned
}
hipError_t hipIpcCloseMemHandle(void* devPtr) {
HIP_INIT_API(devPtr);
HIP_INIT_API(hipIpcCloseMemHandle, devPtr);
hipError_t hipStatus = hipSuccess;
if (devPtr == NULL) {
hipStatus = hipErrorInvalidValue;
projects/hip/src/hip_module.cpp
+10 -9
Просмотреть файл
@@ -84,6 +84,7 @@ string ToString(hipFunction_t v) {
return ss.str();
};
std::string& FunctionSymbol(hipFunction_t f) { return f->_name; };
#define CHECK_HSA(hsaStatus, hipStatus) \
if (hsaStatus != HSA_STATUS_SUCCESS) { \
@@ -96,7 +97,7 @@ string ToString(hipFunction_t v) {
}
hipError_t hipModuleUnload(hipModule_t hmod) {
HIP_INIT_API(hmod);
HIP_INIT_API(hipModuleUnload, hmod);
// TODO - improve this synchronization so it is thread-safe.
// Currently we want for all inflight activity to complete, but don't prevent another
@@ -230,7 +231,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, uint32_t gridDimX, uint32_t gr
uint32_t gridDimZ, uint32_t blockDimX, uint32_t blockDimY,
uint32_t blockDimZ, uint32_t sharedMemBytes, hipStream_t hStream,
void** kernelParams, void** extra) {
HIP_INIT_API(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes,
HIP_INIT_API(hipModuleLaunchKernel, f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes,
hStream, kernelParams, extra);
return ihipLogStatus(ihipModuleLaunchKernel(
f, blockDimX * gridDimX, blockDimY * gridDimY, gridDimZ * blockDimZ, blockDimX, blockDimY,
@@ -244,7 +245,7 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
uint32_t localWorkSizeZ, size_t sharedMemBytes,
hipStream_t hStream, void** kernelParams, void** extra,
hipEvent_t startEvent, hipEvent_t stopEvent) {
HIP_INIT_API(f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX,
HIP_INIT_API(hipHccModuleLaunchKernel, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX,
localWorkSizeY, localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra);
return ihipLogStatus(ihipModuleLaunchKernel(
f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY,
@@ -464,13 +465,13 @@ hipError_t ihipModuleGetFunction(hipFunction_t* func, hipModule_t hmod, const ch
}
hipError_t hipModuleGetFunction(hipFunction_t* hfunc, hipModule_t hmod, const char* name) {
HIP_INIT_API(hfunc, hmod, name);
HIP_INIT_API(hipModuleGetFunction, hfunc, hmod, name);
return ihipLogStatus(ihipModuleGetFunction(hfunc, hmod, name));
}
hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t hmod,
const char* name) {
HIP_INIT_API(dptr, bytes, hmod, name);
HIP_INIT_API(hipModuleGetGlobal, dptr, bytes, hmod, name);
return ihipLogStatus(ihipModuleGetGlobal(dptr, bytes, hmod, name));
}
@@ -567,12 +568,12 @@ hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) {
}
hipError_t hipModuleLoadData(hipModule_t* module, const void* image) {
HIP_INIT_API(module, image);
HIP_INIT_API(hipModuleLoadData, module, image);
return ihipLogStatus(ihipModuleLoadData(module,image));
}
hipError_t hipModuleLoad(hipModule_t* module, const char* fname) {
HIP_INIT_API(module, fname);
HIP_INIT_API(hipModuleLoad, module, fname);
if (!fname) return ihipLogStatus(hipErrorInvalidValue);
@@ -587,12 +588,12 @@ hipError_t hipModuleLoad(hipModule_t* module, const char* fname) {
hipError_t hipModuleLoadDataEx(hipModule_t* module, const void* image, unsigned int numOptions,
hipJitOption* options, void** optionValues) {
HIP_INIT_API(module, image, numOptions, options, optionValues);
HIP_INIT_API(hipModuleLoadDataEx, module, image, numOptions, options, optionValues);
return ihipLogStatus(ihipModuleLoadData(module, image));
}
hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const char* name) {
HIP_INIT_API(texRef, hmod, name);
HIP_INIT_API(hipModuleGetTexRef, texRef, hmod, name);
hipError_t ret = hipErrorNotFound;
if (!texRef) return ihipLogStatus(hipErrorInvalidValue);
projects/hip/src/hip_peer.cpp
+10 -10
Просмотреть файл
@@ -73,7 +73,7 @@ hipError_t ihipDeviceCanAccessPeer(int* canAccessPeer, hipCtx_t thisCtx, hipCtx_
*/
//---
hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, hipCtx_t thisCtx, hipCtx_t peerCtx) {
HIP_INIT_API(canAccessPeer, thisCtx, peerCtx);
HIP_INIT_API(hipDeviceCanAccessPeer2, canAccessPeer, thisCtx, peerCtx);
return ihipLogStatus(ihipDeviceCanAccessPeer(canAccessPeer, thisCtx, peerCtx));
}
@@ -150,7 +150,7 @@ hipError_t ihipEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags) {
//---
hipError_t hipMemcpyPeer(void* dst, hipCtx_t dstCtx, const void* src, hipCtx_t srcCtx,
size_t sizeBytes) {
HIP_INIT_API(dst, dstCtx, src, srcCtx, sizeBytes);
HIP_INIT_API(hipMemcpyPeer2, dst, dstCtx, src, srcCtx, sizeBytes);
// TODO - move to ihip memory copy implementaion.
// HCC has a unified memory architecture so device specifiers are not required.
@@ -161,7 +161,7 @@ hipError_t hipMemcpyPeer(void* dst, hipCtx_t dstCtx, const void* src, hipCtx_t s
//---
hipError_t hipMemcpyPeerAsync(void* dst, hipCtx_t dstDevice, const void* src, hipCtx_t srcDevice,
size_t sizeBytes, hipStream_t stream) {
HIP_INIT_API(dst, dstDevice, src, srcDevice, sizeBytes, stream);
HIP_INIT_API(hipMemcpyPeerAsync2, dst, dstDevice, src, srcDevice, sizeBytes, stream);
// TODO - move to ihip memory copy implementaion.
// HCC has a unified memory architecture so device specifiers are not required.
@@ -175,21 +175,21 @@ hipError_t hipMemcpyPeerAsync(void* dst, hipCtx_t dstDevice, const void* src, hi
//=============================================================================
hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, int deviceId, int peerDeviceId) {
HIP_INIT_API(canAccessPeer, deviceId, peerDeviceId);
HIP_INIT_API(hipDeviceCanAccessPeer, canAccessPeer, deviceId, peerDeviceId);
return ihipLogStatus(ihipDeviceCanAccessPeer(canAccessPeer, ihipGetPrimaryCtx(deviceId),
ihipGetPrimaryCtx(peerDeviceId)));
}
hipError_t hipDeviceDisablePeerAccess(int peerDeviceId) {
HIP_INIT_API(peerDeviceId);
HIP_INIT_API(hipDeviceDisablePeerAccess, peerDeviceId);
return ihipLogStatus(ihipDisablePeerAccess(ihipGetPrimaryCtx(peerDeviceId)));
}
hipError_t hipDeviceEnablePeerAccess(int peerDeviceId, unsigned int flags) {
HIP_INIT_API(peerDeviceId, flags);
HIP_INIT_API(hipDeviceEnablePeerAccess, peerDeviceId, flags);
return ihipLogStatus(ihipEnablePeerAccess(ihipGetPrimaryCtx(peerDeviceId), flags));
}
@@ -197,7 +197,7 @@ hipError_t hipDeviceEnablePeerAccess(int peerDeviceId, unsigned int flags) {
hipError_t hipMemcpyPeer(void* dst, int dstDevice, const void* src, int srcDevice,
size_t sizeBytes) {
HIP_INIT_API(dst, dstDevice, src, srcDevice, sizeBytes);
HIP_INIT_API(hipMemcpyPeer, dst, dstDevice, src, srcDevice, sizeBytes);
return ihipLogStatus(hipMemcpyPeer(dst, ihipGetPrimaryCtx(dstDevice), src,
ihipGetPrimaryCtx(srcDevice), sizeBytes));
}
@@ -205,18 +205,18 @@ hipError_t hipMemcpyPeer(void* dst, int dstDevice, const void* src, int srcDevic
hipError_t hipMemcpyPeerAsync(void* dst, int dstDevice, const void* src, int srcDevice,
size_t sizeBytes, hipStream_t stream) {
HIP_INIT_API(dst, dstDevice, src, srcDevice, sizeBytes, stream);
HIP_INIT_API(hipMemcpyPeerAsync, dst, dstDevice, src, srcDevice, sizeBytes, stream);
return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyDefault, stream));
}
hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags) {
HIP_INIT_API(peerCtx, flags);
HIP_INIT_API(hipCtxEnablePeerAccess, peerCtx, flags);
return ihipLogStatus(ihipEnablePeerAccess(peerCtx, flags));
}
hipError_t hipCtxDisablePeerAccess(hipCtx_t peerCtx) {
HIP_INIT_API(peerCtx);
HIP_INIT_API(hipCtxDisablePeerAccess, peerCtx);
return ihipLogStatus(ihipDisablePeerAccess(peerCtx));
}
projects/hip/src/hip_stream.cpp
+11 -11
Просмотреть файл
@@ -90,21 +90,21 @@ hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags, int priorit
//---
hipError_t hipStreamCreateWithFlags(hipStream_t* stream, unsigned int flags) {
HIP_INIT_API(stream, flags);
HIP_INIT_API(hipStreamCreateWithFlags, stream, flags);
return ihipLogStatus(ihipStreamCreate(stream, flags, priority_normal));
}
//---
hipError_t hipStreamCreate(hipStream_t* stream) {
HIP_INIT_API(stream);
HIP_INIT_API(hipStreamCreate, stream);
return ihipLogStatus(ihipStreamCreate(stream, hipStreamDefault, priority_normal));
}
//---
hipError_t hipStreamCreateWithPriority(hipStream_t* stream, unsigned int flags, int priority) {
HIP_INIT_API(stream, flags, priority);
HIP_INIT_API(hipStreamCreateWithPriority, stream, flags, priority);
// clamp priority to range [priority_high:priority_low]
priority = (priority < priority_high ? priority_high : (priority > priority_low ? priority_low : priority));
@@ -113,7 +113,7 @@ hipError_t hipStreamCreateWithPriority(hipStream_t* stream, unsigned int flags,
//---
hipError_t hipDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPriority) {
HIP_INIT_API(leastPriority, greatestPriority);
HIP_INIT_API(hipDeviceGetStreamPriorityRange, leastPriority, greatestPriority);
if (leastPriority != NULL) *leastPriority = priority_low;
if (greatestPriority != NULL) *greatestPriority = priority_high;
@@ -121,7 +121,7 @@ hipError_t hipDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPrio
}
hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags) {
HIP_INIT_SPECIAL_API(TRACE_SYNC, stream, event, flags);
HIP_INIT_SPECIAL_API(hipStreamWaitEvent, TRACE_SYNC, stream, event, flags);
hipError_t e = hipSuccess;
@@ -152,7 +152,7 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
//---
hipError_t hipStreamQuery(hipStream_t stream) {
HIP_INIT_SPECIAL_API(TRACE_QUERY, stream);
HIP_INIT_SPECIAL_API(hipStreamQuery, TRACE_QUERY, stream);
// Use default stream if 0 specified:
if (stream == hipStreamNull) {
@@ -175,7 +175,7 @@ hipError_t hipStreamQuery(hipStream_t stream) {
//---
hipError_t hipStreamSynchronize(hipStream_t stream) {
HIP_INIT_SPECIAL_API(TRACE_SYNC, stream);
HIP_INIT_SPECIAL_API(hipStreamSynchronize, TRACE_SYNC, stream);
return ihipLogStatus(ihipStreamSynchronize(stream));
}
@@ -186,7 +186,7 @@ hipError_t hipStreamSynchronize(hipStream_t stream) {
* @return #hipSuccess, #hipErrorInvalidResourceHandle
*/
hipError_t hipStreamDestroy(hipStream_t stream) {
HIP_INIT_API(stream);
HIP_INIT_API(hipStreamDestroy, stream);
hipError_t e = hipSuccess;
@@ -214,7 +214,7 @@ hipError_t hipStreamDestroy(hipStream_t stream) {
//---
hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int* flags) {
HIP_INIT_API(stream, flags);
HIP_INIT_API(hipStreamGetFlags, stream, flags);
if (flags == NULL) {
return ihipLogStatus(hipErrorInvalidValue);
@@ -229,7 +229,7 @@ hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int* flags) {
//--
hipError_t hipStreamGetPriority(hipStream_t stream, int* priority) {
HIP_INIT_API(stream, priority);
HIP_INIT_API(hipStreamGetPriority, stream, priority);
if (priority == NULL) {
return ihipLogStatus(hipErrorInvalidValue);
@@ -250,7 +250,7 @@ hipError_t hipStreamGetPriority(hipStream_t stream, int* priority) {
//---
hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData,
unsigned int flags) {
HIP_INIT_API(stream, callback, userData, flags);
HIP_INIT_API(hipStreamAddCallback, stream, callback, userData, flags);
hipError_t e = hipSuccess;
// Create a thread in detached mode to handle callback
projects/hip/src/hip_surface.cpp
+2 -2
Просмотреть файл
@@ -41,7 +41,7 @@ void saveSurfaceInfo(const hipSurface* pSurface, const hipResourceDesc* pResDesc
// Surface Object APIs
hipError_t hipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject,
const hipResourceDesc* pResDesc) {
HIP_INIT_API(pSurfObject, pResDesc);
HIP_INIT_API(hipCreateSurfaceObject, pSurfObject, pResDesc);
hipError_t hip_status = hipSuccess;
auto ctx = ihipGetTlsDefaultCtx();
@@ -71,7 +71,7 @@ hipError_t hipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject,
}
hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject) {
HIP_INIT_API(surfaceObject);
HIP_INIT_API(hipDestroySurfaceObject, surfaceObject);
hipError_t hip_status = hipSuccess;
projects/hip/src/hip_texture.cpp
+20 -20
Просмотреть файл
@@ -202,7 +202,7 @@ bool getHipTextureObject(hipTextureObject_t* pTexObject, hsa_ext_image_t& image,
hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResourceDesc* pResDesc,
const hipTextureDesc* pTexDesc,
const hipResourceViewDesc* pResViewDesc) {
HIP_INIT_API(pTexObject, pResDesc, pTexDesc, pResViewDesc);
HIP_INIT_API(hipCreateTextureObject, pTexObject, pResDesc, pTexDesc, pResViewDesc);
hipError_t hip_status = hipSuccess;
auto ctx = ihipGetTlsDefaultCtx();
@@ -311,7 +311,7 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou
}
hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject) {
HIP_INIT_API(textureObject);
HIP_INIT_API(hipDestroyTextureObject, textureObject);
hipError_t hip_status = hipSuccess;
@@ -335,7 +335,7 @@ hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject) {
hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDesc,
hipTextureObject_t textureObject) {
HIP_INIT_API(pResDesc, textureObject);
HIP_INIT_API(hipGetTextureObjectResourceDesc, pResDesc, textureObject);
hipError_t hip_status = hipSuccess;
auto ctx = ihipGetTlsDefaultCtx();
@@ -350,7 +350,7 @@ hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDesc,
hipError_t hipGetTextureObjectResourceViewDesc(hipResourceViewDesc* pResViewDesc,
hipTextureObject_t textureObject) {
HIP_INIT_API(pResViewDesc, textureObject);
HIP_INIT_API(hipGetTextureObjectResourceViewDesc, pResViewDesc, textureObject);
hipError_t hip_status = hipSuccess;
auto ctx = ihipGetTlsDefaultCtx();
@@ -366,7 +366,7 @@ hipError_t hipGetTextureObjectResourceViewDesc(hipResourceViewDesc* pResViewDesc
hipError_t hipGetTextureObjectTextureDesc(hipTextureDesc* pTexDesc,
hipTextureObject_t textureObject) {
HIP_INIT_API(pTexDesc, textureObject);
HIP_INIT_API(hipGetTextureObjectTextureDesc, pTexDesc, textureObject);
hipError_t hip_status = hipSuccess;
@@ -444,7 +444,7 @@ hipError_t ihipBindTextureImpl(int dim, enum hipTextureReadMode readMode, size_t
hipError_t hipBindTexture(size_t* offset, textureReference* tex, const void* devPtr,
const hipChannelFormatDesc* desc, size_t size) {
HIP_INIT_API(offset, tex, devPtr, desc, size);
HIP_INIT_API(hipBindTexture, offset, tex, devPtr, desc, size);
hipError_t hip_status = hipSuccess;
// TODO: hipReadModeElementType is default.
hip_status = ihipBindTextureImpl(hipTextureType1D, hipReadModeElementType, offset, devPtr, desc,
@@ -517,7 +517,7 @@ hipError_t ihipBindTexture2DImpl(int dim, enum hipTextureReadMode readMode, size
hipError_t hipBindTexture2D(size_t* offset, textureReference* tex, const void* devPtr,
const hipChannelFormatDesc* desc, size_t width, size_t height,
size_t pitch) {
HIP_INIT_API(offset, tex, devPtr, desc, width, height, pitch);
HIP_INIT_API(hipBindTexture2D, offset, tex, devPtr, desc, width, height, pitch);
hipError_t hip_status = hipSuccess;
hip_status = ihipBindTexture2DImpl(hipTextureType2D, hipReadModeElementType, offset, devPtr,
desc, width, height, tex);
@@ -613,7 +613,7 @@ hipError_t ihipBindTextureToArrayImpl(int dim, enum hipTextureReadMode readMode,
hipError_t hipBindTextureToArray(textureReference* tex, hipArray_const_t array,
const hipChannelFormatDesc* desc) {
HIP_INIT_API(tex, array, desc);
HIP_INIT_API(hipBindTextureToArray, tex, array, desc);
hipError_t hip_status = hipSuccess;
// TODO: hipReadModeElementType is default.
hip_status =
@@ -624,7 +624,7 @@ hipError_t hipBindTextureToArray(textureReference* tex, hipArray_const_t array,
hipError_t hipBindTextureToMipmappedArray(textureReference* tex,
hipMipmappedArray_const_t mipmappedArray,
const hipChannelFormatDesc* desc) {
HIP_INIT_API(tex, mipmappedArray, desc);
HIP_INIT_API(hipBindTextureToMipmappedArray, tex, mipmappedArray, desc);
hipError_t hip_status = hipSuccess;
return ihipLogStatus(hip_status);
}
@@ -652,14 +652,14 @@ hipError_t ihipUnbindTextureImpl(const hipTextureObject_t& textureObject) {
}
hipError_t hipUnbindTexture(const textureReference* tex) {
HIP_INIT_API(tex);
HIP_INIT_API(hipUnbindTexture, tex);
hipError_t hip_status = hipSuccess;
hip_status = ihipUnbindTextureImpl(tex->textureObject);
return ihipLogStatus(hip_status);
}
hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_const_t array) {
HIP_INIT_API(desc, array);
HIP_INIT_API(hipGetChannelDesc, desc, array);
hipError_t hip_status = hipSuccess;
auto ctx = ihipGetTlsDefaultCtx();
@@ -670,7 +670,7 @@ hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_const_t array)
}
hipError_t hipGetTextureAlignmentOffset(size_t* offset, const textureReference* tex) {
HIP_INIT_API(offset, tex);
HIP_INIT_API(hipGetTextureAlignmentOffset, offset, tex);
hipError_t hip_status = hipSuccess;
@@ -683,7 +683,7 @@ hipError_t hipGetTextureAlignmentOffset(size_t* offset, const textureReference*
}
hipError_t hipGetTextureReference(const textureReference** tex, const void* symbol) {
HIP_INIT_API(tex, symbol);
HIP_INIT_API(hipGetTextureReference, tex, symbol);
hipError_t hip_status = hipSuccess;
@@ -694,7 +694,7 @@ hipError_t hipGetTextureReference(const textureReference** tex, const void* symb
}
hipError_t hipTexRefSetFormat(textureReference* tex, hipArray_Format fmt, int NumPackedComponents) {
HIP_INIT_API(tex, fmt, NumPackedComponents);
HIP_INIT_API(hipTexRefSetFormat, tex, fmt, NumPackedComponents);
hipError_t hip_status = hipSuccess;
tex->format = fmt;
tex->numChannels = NumPackedComponents;
@@ -702,28 +702,28 @@ hipError_t hipTexRefSetFormat(textureReference* tex, hipArray_Format fmt, int Nu
}
hipError_t hipTexRefSetFlags(textureReference* tex, unsigned int flags) {
HIP_INIT_API(tex, flags);
HIP_INIT_API(hipTexRefSetFlags, tex, flags);
hipError_t hip_status = hipSuccess;
tex->normalized = flags;
return ihipLogStatus(hip_status);
}
hipError_t hipTexRefSetFilterMode(textureReference* tex, hipTextureFilterMode fm) {
HIP_INIT_API(tex, fm);
HIP_INIT_API(hipTexRefSetFilterMode, tex, fm);
hipError_t hip_status = hipSuccess;
tex->filterMode = fm;
return ihipLogStatus(hip_status);
}
hipError_t hipTexRefSetAddressMode(textureReference* tex, int dim, hipTextureAddressMode am) {
HIP_INIT_API(tex, dim, am);
HIP_INIT_API(hipTexRefSetAddressMode, tex, dim, am);
hipError_t hip_status = hipSuccess;
tex->addressMode[dim] = am;
return ihipLogStatus(hip_status);
}
hipError_t hipTexRefSetArray(textureReference* tex, hipArray_const_t array, unsigned int flags) {
HIP_INIT_API(tex, array, flags);
HIP_INIT_API(hipTexRefSetArray, tex, array, flags);
hipError_t hip_status = hipSuccess;
hip_status = ihipBindTextureToArrayImpl(array->textureType, hipReadModeElementType, array,
@@ -734,7 +734,7 @@ hipError_t hipTexRefSetArray(textureReference* tex, hipArray_const_t array, unsi
hipError_t hipTexRefSetAddress(size_t* offset, textureReference* tex, hipDeviceptr_t devPtr,
size_t size) {
HIP_INIT_API(offset, tex, devPtr, size);
HIP_INIT_API(hipTexRefSetAddress, offset, tex, devPtr, size);
hipError_t hip_status = hipSuccess;
// TODO: hipReadModeElementType is default.
hip_status = ihipBindTextureImpl(hipTextureType1D, hipReadModeElementType, offset, devPtr, NULL,
@@ -744,7 +744,7 @@ hipError_t hipTexRefSetAddress(size_t* offset, textureReference* tex, hipDevicep
hipError_t hipTexRefSetAddress2D(textureReference* tex, const HIP_ARRAY_DESCRIPTOR* desc,
hipDeviceptr_t devPtr, size_t pitch) {
HIP_INIT_API(tex, desc, devPtr, pitch);
HIP_INIT_API(hipTexRefSetAddress2D, tex, desc, devPtr, pitch);
size_t offset;
hipError_t hip_status = hipSuccess;
// TODO: hipReadModeElementType is default.