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

P4 to Git Change 2009236 by eshcherb@evgeny-hip on 2019/10/07 11:49:55

Browse Source 
SWDEV-197287 - HIP tracing layer instrumentation

	hip_prof_str.h - generated header; as a next step a genrator will be integrated in HIP makefile (similar to HIP-HCC)
	prof_protocol.h - temporarily included in HIP sources. Provided by roctracer-proto package. HIP-HCC CMake has parameter to set a path to the header.

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/build/Makefile.hip#23 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_context.cpp#23 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_device.cpp#24 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_device_runtime.cpp#19 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_error.cpp#3 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_event.cpp#17 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_intercept.cpp#1 add
... //depot/stg/opencl/drivers/opencl/api/hip/hip_internal.hpp#44 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#74 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_module.cpp#43 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_peer.cpp#7 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_platform.cpp#44 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_prof_api.h#1 add
... //depot/stg/opencl/drivers/opencl/api/hip/hip_profile.cpp#3 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_stream.cpp#25 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_surface.cpp#5 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_texture.cpp#21 edit
... //depot/stg/opencl/drivers/opencl/api/hip/prof_protocol.h#1 add
... //depot/stg/opencl/drivers/opencl/hip_prof_gen.py#1 add
... //depot/stg/opencl/drivers/opencl/make/hip.git/include/hip/hcc_detail/hip_prof_str.h#2 delete


[ROCm/hip commit: 6e7f1dea66]
This commit is contained in:
foreman
2019-10-07 11:55:30 -04:00
parent 5f8c952ecf
commit 1ae22b2688
17 changed files with 596 additions and 186 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/api/hip/hip_context.cpp
+22 -22
View File
@@ -112,13 +112,13 @@ amd::HostQueue* getNullStream() {
using namespace hip;
hipError_t hipInit(unsigned int flags) {
HIP_INIT_API(flags);
HIP_INIT_API(hipInit, flags);
HIP_RETURN(hipSuccess);
}
hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags, hipDevice_t device) {
HIP_INIT_API(ctx, flags, device);
HIP_INIT_API(hipCtxCreate, ctx, flags, device);
if (static_cast<size_t>(device) >= g_devices.size()) {
HIP_RETURN(hipErrorInvalidValue);
@@ -134,7 +134,7 @@ hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags, hipDevice_t device)
}
hipError_t hipCtxSetCurrent(hipCtx_t ctx) {
HIP_INIT_API(ctx);
HIP_INIT_API(hipCtxSetCurrent, ctx);
if (ctx == nullptr) {
if(!g_ctxtStack.empty()) {
@@ -152,7 +152,7 @@ hipError_t hipCtxSetCurrent(hipCtx_t ctx) {
}
hipError_t hipCtxGetCurrent(hipCtx_t* ctx) {
HIP_INIT_API(ctx);
HIP_INIT_API(hipCtxGetCurrent, ctx);
*ctx = reinterpret_cast<hipCtx_t>(hip::getCurrentContext());
@@ -160,7 +160,7 @@ hipError_t hipCtxGetCurrent(hipCtx_t* ctx) {
}
hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig* pConfig) {
HIP_INIT_API(pConfig);
HIP_INIT_API(hipCtxGetSharedMemConfig, pConfig);
*pConfig = hipSharedMemBankSizeFourByte;
@@ -168,7 +168,7 @@ hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig* pConfig) {
}
hipError_t hipRuntimeGetVersion(int *runtimeVersion) {
HIP_INIT_API(runtimeVersion);
HIP_INIT_API(hipRuntimeGetVersion, runtimeVersion);
if (!runtimeVersion) {
HIP_RETURN(hipErrorInvalidValue);
@@ -180,7 +180,7 @@ hipError_t hipRuntimeGetVersion(int *runtimeVersion) {
}
hipError_t hipCtxDestroy(hipCtx_t ctx) {
HIP_INIT_API(ctx);
HIP_INIT_API(hipCtxDestroy, ctx);
amd::Context* amdContext = reinterpret_cast<amd::Context*>(as_amd(ctx));
if (amdContext == nullptr) {
@@ -207,7 +207,7 @@ hipError_t hipCtxDestroy(hipCtx_t ctx) {
}
hipError_t hipCtxPopCurrent(hipCtx_t* ctx) {
HIP_INIT_API(ctx);
HIP_INIT_API(hipCtxPopCurrent, ctx);
amd::Context* amdContext = reinterpret_cast<amd::Context*>(as_amd(ctx));
if (amdContext == nullptr) {
@@ -225,7 +225,7 @@ hipError_t hipCtxPopCurrent(hipCtx_t* ctx) {
}
hipError_t hipCtxPushCurrent(hipCtx_t ctx) {
HIP_INIT_API(ctx);
HIP_INIT_API(hipCtxPushCurrent, ctx);
amd::Context* amdContext = reinterpret_cast<amd::Context*>(as_amd(ctx));
if (amdContext == nullptr) {
@@ -239,7 +239,7 @@ hipError_t hipCtxPushCurrent(hipCtx_t ctx) {
}
hipError_t hipDriverGetVersion(int* driverVersion) {
HIP_INIT_API(driverVersion);
HIP_INIT_API(hipDriverGetVersion, driverVersion);
auto* deviceHandle = g_devices[0]->devices()[0];
const auto& info = deviceHandle->info();
@@ -255,7 +255,7 @@ hipError_t hipDriverGetVersion(int* driverVersion) {
}
hipError_t hipCtxGetDevice(hipDevice_t* device) {
HIP_INIT_API(device);
HIP_INIT_API(hipCtxGetDevice, device);
if (device != nullptr) {
for (unsigned int i = 0; i < g_devices.size(); i++) {
@@ -272,7 +272,7 @@ hipError_t hipCtxGetDevice(hipDevice_t* device) {
}
hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion) {
HIP_INIT_API(apiVersion);
HIP_INIT_API(hipCtxGetApiVersion, apiVersion);
assert(0 && "Unimplemented");
@@ -280,7 +280,7 @@ hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion) {
}
hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig) {
HIP_INIT_API(cacheConfig);
HIP_INIT_API(hipCtxGetCacheConfig, cacheConfig);
assert(0 && "Unimplemented");
@@ -288,7 +288,7 @@ hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig) {
}
hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig) {
HIP_INIT_API(cacheConfig);
HIP_INIT_API(hipCtxSetCacheConfig, cacheConfig);
assert(0 && "Unimplemented");
@@ -296,7 +296,7 @@ hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig) {
}
hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config) {
HIP_INIT_API(config);
HIP_INIT_API(hipCtxSetSharedMemConfig, config);
assert(0 && "Unimplemented");
@@ -304,7 +304,7 @@ hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config) {
}
hipError_t hipCtxSynchronize(void) {
HIP_INIT_API(1);
HIP_INIT_API(hipCtxSynchronize, 1);
assert(0 && "Unimplemented");
@@ -312,7 +312,7 @@ hipError_t hipCtxSynchronize(void) {
}
hipError_t hipCtxGetFlags(unsigned int* flags) {
HIP_INIT_API(flags);
HIP_INIT_API(hipCtxGetFlags, flags);
assert(0 && "Unimplemented");
@@ -320,7 +320,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);
if (static_cast<unsigned int>(dev) >= g_devices.size()) {
HIP_RETURN(hipErrorInvalidDevice);
@@ -338,7 +338,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);
if (static_cast<unsigned int>(dev) >= g_devices.size()) {
HIP_RETURN(hipErrorInvalidDevice);
@@ -348,7 +348,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);
if (static_cast<unsigned int>(dev) >= g_devices.size()) {
HIP_RETURN(hipErrorInvalidDevice);
@@ -363,13 +363,13 @@ hipError_t hipDevicePrimaryCtxRetain(hipCtx_t* pctx, hipDevice_t dev) {
}
hipError_t hipDevicePrimaryCtxReset(hipDevice_t dev) {
HIP_INIT_API(dev);
HIP_INIT_API(hipDevicePrimaryCtxReset, dev);
HIP_RETURN(hipSuccess);
}
hipError_t hipDevicePrimaryCtxSetFlags(hipDevice_t dev, unsigned int flags) {
HIP_INIT_API(dev, flags);
HIP_INIT_API(hipDevicePrimaryCtxSetFlags, dev, flags);
if (static_cast<unsigned int>(dev) >= g_devices.size()) {
HIP_RETURN(hipErrorInvalidDevice);
projects/hip/api/hip/hip_device.cpp
+9 -9
View File
@@ -25,7 +25,7 @@ THE SOFTWARE.
#include "hip_internal.hpp"
hipError_t hipDeviceGet(hipDevice_t *device, int deviceId) {
HIP_INIT_API(device, deviceId);
HIP_INIT_API(hipDeviceGet, device, deviceId);
if (device != nullptr) {
*device = deviceId;
@@ -38,7 +38,7 @@ hipError_t hipDeviceGet(hipDevice_t *device, int deviceId) {
hipError_t hipFuncSetCacheConfig (const void* func, hipFuncCache_t cacheConfig) {
HIP_INIT_API(cacheConfig);
HIP_INIT_API(hipFuncSetCacheConfig, cacheConfig);
// No way to set cache config yet.
@@ -47,7 +47,7 @@ hipError_t hipFuncSetCacheConfig (const void* func, hipFuncCache_t cacheConfig)
hipError_t hipDeviceTotalMem (size_t *bytes, hipDevice_t device) {
HIP_INIT_API(bytes, device);
HIP_INIT_API(hipDeviceTotalMem, bytes, device);
if (device < 0 || static_cast<size_t>(device) >= g_devices.size()) {
HIP_RETURN(hipErrorInvalidDevice);
@@ -67,7 +67,7 @@ hipError_t hipDeviceTotalMem (size_t *bytes, hipDevice_t device) {
hipError_t hipDeviceComputeCapability(int *major, int *minor, hipDevice_t device) {
HIP_INIT_API(major, minor, device);
HIP_INIT_API(hipDeviceComputeCapability, major, minor, device);
if (device < 0 || static_cast<size_t>(device) >= g_devices.size()) {
HIP_RETURN(hipErrorInvalidDevice);
@@ -86,7 +86,7 @@ hipError_t hipDeviceComputeCapability(int *major, int *minor, hipDevice_t device
}
hipError_t hipDeviceGetCount(int* count) {
HIP_INIT_API(count);
HIP_INIT_API(NONE, count);
HIP_RETURN(ihipDeviceGetCount(count));
}
@@ -108,7 +108,7 @@ hipError_t ihipDeviceGetCount(int* count) {
hipError_t hipDeviceGetName(char *name, int len, hipDevice_t device) {
HIP_INIT_API((void*)name, len, device);
HIP_INIT_API(hipDeviceGetName, (void*)name, len, device);
if (device < 0 || static_cast<size_t>(device) >= g_devices.size()) {
HIP_RETURN(hipErrorInvalidDevice);
@@ -134,7 +134,7 @@ hipError_t hipDeviceGetName(char *name, int len, hipDevice_t device) {
}
hipError_t hipGetDeviceProperties ( hipDeviceProp_t* props, hipDevice_t device ) {
HIP_INIT_API(props, device);
HIP_INIT_API(hipGetDeviceProperties, props, device);
if (props == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
@@ -218,7 +218,7 @@ hipError_t hipGetDeviceProperties ( hipDeviceProp_t* props, hipDevice_t device )
}
hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator* acc) {
HIP_INIT_API(deviceId, acc);
HIP_INIT_API(NONE, deviceId, acc);
assert(0 && "Unimplemented");
@@ -226,7 +226,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(NONE, stream, av);
assert(0 && "Unimplemented");
projects/hip/api/hip/hip_device_runtime.cpp
+19 -19
View File
@@ -26,7 +26,7 @@ THE SOFTWARE.
hipError_t hipChooseDevice(int* device, const hipDeviceProp_t* properties) {
HIP_INIT_API(device, properties);
HIP_INIT_API(hipChooseDevice, device, properties);
if (device == nullptr || properties == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
@@ -140,7 +140,7 @@ hipError_t hipChooseDevice(int* device, const hipDeviceProp_t* properties) {
hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) {
HIP_INIT_API(pi, attr, device);
HIP_INIT_API(hipDeviceGetAttribute, pi, attr, device);
if (pi == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
@@ -289,7 +289,7 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device)
hipError_t hipDeviceGetByPCIBusId(int* device, const char*pciBusIdstr) {
HIP_INIT_API(device, pciBusIdstr);
HIP_INIT_API(hipDeviceGetByPCIBusId, device, pciBusIdstr);
if (device == nullptr || pciBusIdstr == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
@@ -319,7 +319,7 @@ hipError_t hipDeviceGetByPCIBusId(int* device, const char*pciBusIdstr) {
}
hipError_t hipDeviceGetCacheConfig ( hipFuncCache_t * cacheConfig ) {
HIP_INIT_API(cacheConfig);
HIP_INIT_API(hipDeviceGetCacheConfig, cacheConfig);
if(cacheConfig == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
@@ -332,7 +332,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) {
HIP_RETURN(hipErrorInvalidValue);
@@ -357,7 +357,7 @@ hipError_t hipDeviceGetP2PAttribute ( int* value, hipDeviceP2PAttr attr, int sr
hipError_t hipDeviceGetPCIBusId ( char* pciBusId, int len, int device ) {
HIP_INIT_API((void*)pciBusId, len, device);
HIP_INIT_API(hipDeviceGetPCIBusId, (void*)pciBusId, len, device);
int count;
ihipDeviceGetCount(&count);
@@ -381,7 +381,7 @@ hipError_t hipDeviceGetPCIBusId ( char* pciBusId, int len, int device ) {
}
hipError_t hipDeviceGetSharedMemConfig ( hipSharedMemConfig * pConfig ) {
HIP_INIT_API(pConfig);
HIP_INIT_API(hipDeviceGetSharedMemConfig, pConfig);
*pConfig = hipSharedMemBankSizeFourByte;
@@ -389,7 +389,7 @@ hipError_t hipDeviceGetSharedMemConfig ( hipSharedMemConfig * pConfig ) {
}
hipError_t hipDeviceReset ( void ) {
HIP_INIT_API();
HIP_INIT_API(hipDeviceReset);
/* FIXME */
@@ -397,7 +397,7 @@ hipError_t hipDeviceReset ( void ) {
}
hipError_t hipDeviceSetCacheConfig ( hipFuncCache_t cacheConfig ) {
HIP_INIT_API(cacheConfig);
HIP_INIT_API(hipDeviceSetCacheConfig, cacheConfig);
// No way to set cache config yet.
@@ -409,7 +409,7 @@ hipError_t hipDeviceSetLimit ( hipLimit_t limit, size_t value ) {
}
hipError_t hipDeviceSetSharedMemConfig ( hipSharedMemConfig config ) {
HIP_INIT_API(config);
HIP_INIT_API(hipDeviceSetSharedMemConfig, config);
// No way to set cache config yet.
@@ -417,7 +417,7 @@ hipError_t hipDeviceSetSharedMemConfig ( hipSharedMemConfig config ) {
}
hipError_t hipDeviceSynchronize ( void ) {
HIP_INIT_API();
HIP_INIT_API(hipDeviceSynchronize);
hip::syncStreams();
@@ -442,7 +442,7 @@ int ihipGetDevice() {
}
hipError_t hipGetDevice ( int* deviceId ) {
HIP_INIT_API(deviceId);
HIP_INIT_API(hipGetDevice, deviceId);
if (deviceId != nullptr) {
int dev = ihipGetDevice();
@@ -455,7 +455,7 @@ hipError_t hipGetDevice ( int* deviceId ) {
}
hipError_t hipGetDeviceCount ( int* count ) {
HIP_INIT_API(count);
HIP_INIT_API(hipGetDeviceCount, count);
HIP_RETURN(ihipDeviceGetCount(count));
}
@@ -465,7 +465,7 @@ hipError_t hipGetDeviceFlags ( unsigned int* flags ) {
}
hipError_t hipIpcGetEventHandle ( hipIpcEventHandle_t* handle, hipEvent_t event ) {
HIP_INIT_API(handle, event);
HIP_INIT_API(NONE, handle, event);
assert(0 && "Unimplemented");
@@ -473,7 +473,7 @@ hipError_t hipIpcGetEventHandle ( hipIpcEventHandle_t* handle, hipEvent_t event
}
hipError_t hipIpcOpenEventHandle ( hipEvent_t* event, hipIpcEventHandle_t handle ) {
HIP_INIT_API(event, handle);
HIP_INIT_API(NONE, event, handle);
assert(0 && "Unimplemented");
@@ -481,7 +481,7 @@ hipError_t hipIpcOpenEventHandle ( hipEvent_t* event, hipIpcEventHandle_t handle
}
hipError_t hipSetDevice ( int device ) {
HIP_INIT_API(device);
HIP_INIT_API(hipSetDevice, device);
if (static_cast<unsigned int>(device) < g_devices.size()) {
hip::setCurrentContext(device);
@@ -492,7 +492,7 @@ hipError_t hipSetDevice ( int device ) {
}
hipError_t hipSetDeviceFlags ( unsigned int flags ) {
HIP_INIT_API(flags);
HIP_INIT_API(hipSetDeviceFlags, flags);
/* FIXME */
/* Not all of Ctx may be implemented */
@@ -508,7 +508,7 @@ hipError_t hipSetDeviceFlags ( unsigned int flags ) {
}
hipError_t hipSetValidDevices ( int* device_arr, int len ) {
HIP_INIT_API(device_arr, len);
HIP_INIT_API(NONE, device_arr, len);
assert(0 && "Unimplemented");
@@ -516,7 +516,7 @@ hipError_t hipSetValidDevices ( int* device_arr, int len ) {
}
hipError_t hipExtGetLinkTypeAndHopCount(int device1, int device2, uint32_t* linktype, uint32_t* hopcount) {
HIP_INIT_API(device1, device2, linktype, hopcount);
HIP_INIT_API(hipExtGetLinkTypeAndHopCount, device1, device2, linktype, hopcount);
const int numDevices = static_cast<int>(g_devices.size());
projects/hip/api/hip/hip_error.cpp
+2 -2
View File
@@ -26,7 +26,7 @@ THE SOFTWARE.
hipError_t hipGetLastError()
{
HIP_INIT_API();
HIP_INIT_API(hipGetLastError);
hipError_t err = hip::g_lastError;
hip::g_lastError = hipSuccess;
return err;
@@ -34,7 +34,7 @@ hipError_t hipGetLastError()
hipError_t hipPeekAtLastError()
{
HIP_INIT_API();
HIP_INIT_API(hipPeekAtLastError);
hipError_t err = hip::g_lastError;
HIP_RETURN(err);
}
projects/hip/api/hip/hip_event.cpp
+8 -8
View File
@@ -182,19 +182,19 @@ hipError_t ihipEventQuery(hipEvent_t event) {
}
hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags) {
HIP_INIT_API(event, flags);
HIP_INIT_API(hipEventCreateWithFlags, event, flags);
HIP_RETURN(ihipEventCreateWithFlags(event, flags));
}
}
hipError_t hipEventCreate(hipEvent_t* event) {
HIP_INIT_API(event);
HIP_INIT_API(hipEventCreate, event);
HIP_RETURN(ihipEventCreateWithFlags(event, 0));
}
hipError_t hipEventDestroy(hipEvent_t event) {
HIP_INIT_API(event);
HIP_INIT_API(hipEventDestroy, event);
if (event == nullptr) {
HIP_RETURN(hipErrorInvalidResourceHandle);
@@ -206,7 +206,7 @@ hipError_t hipEventDestroy(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);
if (start == nullptr || stop == nullptr) {
HIP_RETURN(hipErrorInvalidResourceHandle);
@@ -223,7 +223,7 @@ hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop) {
}
hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) {
HIP_INIT_API(event, stream);
HIP_INIT_API(hipEventRecord, event, stream);
if (event == nullptr) {
HIP_RETURN(hipErrorInvalidResourceHandle);
@@ -246,7 +246,7 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) {
}
hipError_t hipEventSynchronize(hipEvent_t event) {
HIP_INIT_API(event);
HIP_INIT_API(hipEventSynchronize, event);
if (event == nullptr) {
HIP_RETURN(hipErrorInvalidResourceHandle);
@@ -258,7 +258,7 @@ hipError_t hipEventSynchronize(hipEvent_t event) {
}
hipError_t hipEventQuery(hipEvent_t event) {
HIP_INIT_API(event);
HIP_INIT_API(hipEventQuery, event);
HIP_RETURN(ihipEventQuery(event));
}
projects/hip/api/hip/hip_intercept.cpp
+58
View File
@@ -0,0 +1,58 @@
/*
Copyright (c) 2019 - 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_prof_api.h"
// HIP API callback/activity
api_callbacks_table_t callbacks_table;
extern std::string& FunctionName(const hipFunction_t f);
const char* hipKernelNameRef(const hipFunction_t f) { return FunctionName(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;
}
hipError_t hipEnableTracing(bool enabled) {
callbacks_table.set_enabled(enabled);
return hipSuccess;
}
const char* hipApiName(uint32_t id) {
return hip_api_name(id);
}
projects/hip/api/hip/hip_internal.hpp
+4 -2
View File
@@ -24,6 +24,7 @@ THE SOFTWARE.
#define HIP_SRC_HIP_INTERNAL_H
#include "cl_common.hpp"
#include "hip_prof_api.h"
#include "trace_helper.h"
#include "utils/debug.hpp"
#include <unordered_set>
@@ -51,13 +52,14 @@ typedef struct ihipIpcMemHandle_st {
}
// This macro should be called at the beginning of every HIP API.
#define HIP_INIT_API(...) \
#define HIP_INIT_API(cid, ...) \
LogPrintfInfo("[%zx] %s ( %s )", std::this_thread::get_id(), __func__, ToString( __VA_ARGS__ ).c_str()); \
amd::Thread* thread = amd::Thread::current(); \
if (!CL_CHECK_THREAD(thread)) { \
HIP_RETURN(hipErrorOutOfMemory); \
} \
HIP_INIT();
HIP_INIT(); \
HIP_CB_SPAWNER_OBJECT(cid);
#define HIP_RETURN(ret) \
hip::g_lastError = ret; \
projects/hip/api/hip/hip_memory.cpp
+53 -53
View File
@@ -200,7 +200,7 @@ hipError_t ihipMemset(void* dst, int value, size_t valueSize, size_t sizeBytes,
}
hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flags) {
HIP_INIT_API(ptr, sizeBytes, flags);
HIP_INIT_API(hipExtMallocWithFlags, ptr, sizeBytes, flags);
if (flags != hipDeviceMallocDefault &&
flags != hipDeviceMallocFinegrained) {
@@ -211,13 +211,13 @@ hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flag
}
hipError_t hipMalloc(void** ptr, size_t sizeBytes) {
HIP_INIT_API(ptr, sizeBytes);
HIP_INIT_API(hipMalloc, ptr, sizeBytes);
HIP_RETURN(ihipMalloc(ptr, sizeBytes, 0));
}
hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
HIP_INIT_API(ptr, sizeBytes, flags);
HIP_INIT_API(hipHostMalloc, ptr, sizeBytes, flags);
if (ptr == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
@@ -241,7 +241,7 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
}
hipError_t hipFree(void* ptr) {
HIP_INIT_API(ptr);
HIP_INIT_API(hipFree, ptr);
if (ptr == nullptr) {
HIP_RETURN(hipSuccess);
@@ -256,7 +256,7 @@ hipError_t hipFree(void* ptr) {
}
hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) {
HIP_INIT_API(dst, src, sizeBytes, kind);
HIP_INIT_API(hipMemcpy, dst, src, sizeBytes, kind);
hip::syncStreams();
amd::HostQueue* queue = hip::getNullStream();
@@ -264,31 +264,31 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind
}
hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream) {
HIP_INIT_API(dst, value, sizeBytes, stream);
HIP_INIT_API(hipMemsetAsync, dst, value, sizeBytes, stream);
HIP_RETURN(ihipMemset(dst, value, sizeof(char), sizeBytes, stream, true));
}
hipError_t hipMemsetD32Async(hipDeviceptr_t dst, int value, size_t count, hipStream_t stream) {
HIP_INIT_API(dst, value, count, stream);
HIP_INIT_API(hipMemsetD32Async, dst, value, count, stream);
HIP_RETURN(ihipMemset(dst, value, sizeof(int), count * sizeof(int), stream, true));
}
hipError_t hipMemset(void* dst, int value, size_t sizeBytes) {
HIP_INIT_API(dst, value, sizeBytes);
HIP_INIT_API(hipMemset, dst, value, sizeBytes);
HIP_RETURN(ihipMemset(dst, value, sizeof(char), sizeBytes, nullptr));
}
hipError_t hipMemsetD32(hipDeviceptr_t dst, int value, size_t count) {
HIP_INIT_API(dst, value, count);
HIP_INIT_API(hipMemsetD32, dst, value, count);
HIP_RETURN(ihipMemset(dst, value, sizeof(int), count * sizeof(int), nullptr));
}
hipError_t hipMemPtrGetInfo(void *ptr, size_t *size) {
HIP_INIT_API(ptr, size);
HIP_INIT_API(hipMemPtrGetInfo, ptr, size);
size_t offset = 0;
amd::Memory* svmMem = getMemoryObject(ptr, offset);
@@ -303,7 +303,7 @@ hipError_t hipMemPtrGetInfo(void *ptr, size_t *size) {
}
hipError_t hipHostFree(void* ptr) {
HIP_INIT_API(ptr);
HIP_INIT_API(hipHostFree, ptr);
if (amd::SvmBuffer::malloced(ptr)) {
amd::SvmBuffer::free(*hip::getCurrentContext(), ptr);
@@ -313,7 +313,7 @@ hipError_t hipHostFree(void* ptr) {
}
hipError_t hipFreeArray(hipArray* array) {
HIP_INIT_API(array);
HIP_INIT_API(hipFreeArray, array);
if (amd::SvmBuffer::malloced(array->data)) {
amd::SvmBuffer::free(*hip::getCurrentContext(), array->data);
@@ -323,7 +323,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);
// Since we are using SVM buffer DevicePtr and HostPtr is the same
void* ptr = dptr;
@@ -341,7 +341,7 @@ hipError_t hipMemGetAddressRange(hipDeviceptr_t* pbase, size_t* psize, hipDevice
}
hipError_t hipMemGetInfo(size_t* free, size_t* total) {
HIP_INIT_API(free, total);
HIP_INIT_API(hipMemGetInfo, free, total);
size_t freeMemory[2];
amd::Device* device = hip::getCurrentContext()->devices()[0];
@@ -394,14 +394,14 @@ hipError_t ihipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t heigh
hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height) {
HIP_INIT_API(ptr, pitch, width, height);
HIP_INIT_API(hipMallocPitch, ptr, pitch, width, height);
const cl_image_format image_format = { CL_R, CL_UNSIGNED_INT8 };
HIP_RETURN(ihipMallocPitch(ptr, pitch, width, height, 1, CL_MEM_OBJECT_IMAGE2D, &image_format));
}
hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent) {
HIP_INIT_API(pitchedDevPtr, &extent);
HIP_INIT_API(hipMalloc3D, pitchedDevPtr, &extent);
size_t pitch = 0;
@@ -424,7 +424,7 @@ hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent) {
}
hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent) {
HIP_INIT_API(pitchedDevPtr, value, &extent);
HIP_INIT_API(hipMemset3D, pitchedDevPtr, value, &extent);
void *dst = pitchedDevPtr.ptr;
size_t sizeBytes = pitchedDevPtr.pitch * extent.height * extent.depth;
@@ -433,7 +433,7 @@ hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent)
}
hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray) {
HIP_INIT_API(array, pAllocateArray);
HIP_INIT_API(hipArrayCreate, array, pAllocateArray);
if (pAllocateArray->Width == 0) {
HIP_RETURN(hipErrorInvalidValue);
@@ -461,7 +461,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_API(array, desc, width, height, flags);
HIP_INIT_API(hipMallocArray, array, desc, width, height, flags);
if (width == 0) {
HIP_RETURN(hipErrorInvalidValue);
@@ -503,7 +503,7 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc,
}
hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY3D_DESCRIPTOR* pAllocateArray) {
HIP_INIT_API(array, pAllocateArray);
HIP_INIT_API(hipArray3DCreate, array, pAllocateArray);
*array = (hipArray*)malloc(sizeof(hipArray));
array[0]->type = pAllocateArray->Flags;
@@ -531,7 +531,7 @@ hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY3D_DESCRIPTOR* pAll
hipError_t hipMalloc3DArray(hipArray_t* array, const struct hipChannelFormatDesc* desc,
struct hipExtent extent, unsigned int flags) {
HIP_INIT_API(array, desc, &extent, flags);
HIP_INIT_API(hipMalloc3DArray, array, desc, &extent, flags);
*array = (hipArray*)malloc(sizeof(hipArray));
array[0]->type = flags;
@@ -569,7 +569,7 @@ hipError_t hipMalloc3DArray(hipArray_t* array, const struct hipChannelFormatDesc
}
hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) {
HIP_INIT_API(flagsPtr, hostPtr);
HIP_INIT_API(hipHostGetFlags, flagsPtr, hostPtr);
if (flagsPtr == nullptr ||
hostPtr == nullptr) {
@@ -589,7 +589,7 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) {
}
hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags) {
HIP_INIT_API(hostPtr, sizeBytes, flags);
HIP_INIT_API(hipHostRegister, hostPtr, sizeBytes, flags);
if(hostPtr != nullptr) {
amd::Memory* mem = new (*hip::host_context) amd::Buffer(*hip::host_context, CL_MEM_USE_HOST_PTR | CL_MEM_SVM_ATOMICS, sizeBytes);
@@ -616,7 +616,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);
if (amd::SvmBuffer::malloced(hostPtr)) {
hip::syncStreams();
@@ -651,7 +651,7 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) {
hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, size_t count,
size_t offset, hipMemcpyKind kind) {
HIP_INIT_API(symbolName, src, count, offset, kind);
HIP_INIT_API(hipMemcpyToSymbol, symbolName, src, count, offset, kind);
size_t sym_size = 0;
hipDeviceptr_t device_ptr = nullptr;
@@ -675,7 +675,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_API(symbolName, dst, count, offset, kind);
HIP_INIT_API(hipMemcpyFromSymbol, symbolName, dst, count, offset, kind);
size_t sym_size = 0;
hipDeviceptr_t device_ptr = nullptr;
@@ -699,7 +699,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_API(symbolName, src, count, offset, kind, stream);
HIP_INIT_API(hipMemcpyToSymbolAsync, symbolName, src, count, offset, kind, stream);
size_t sym_size = 0;
hipDeviceptr_t device_ptr = nullptr;
@@ -723,7 +723,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_API(symbolName, dst, count, offset, kind, stream);
HIP_INIT_API(hipMemcpyFromSymbolAsync, symbolName, dst, count, offset, kind, stream);
size_t sym_size = 0;
hipDeviceptr_t device_ptr = nullptr;
@@ -746,7 +746,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t co
}
hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) {
HIP_INIT_API(dst, src, sizeBytes);
HIP_INIT_API(hipMemcpyHtoD, dst, src, sizeBytes);
hip::syncStreams();
amd::HostQueue* queue = hip::getNullStream();
@@ -755,7 +755,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_API(dst, src, sizeBytes);
HIP_INIT_API(hipMemcpyDtoH, dst, src, sizeBytes);
hip::syncStreams();
amd::HostQueue* queue = hip::getNullStream();
@@ -764,7 +764,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_API(dst, src, sizeBytes);
HIP_INIT_API(hipMemcpyDtoD, dst, src, sizeBytes);
hip::syncStreams();
amd::HostQueue* queue = hip::getNullStream();
@@ -773,7 +773,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_API(dst, src, sizeBytes);
HIP_INIT_API(NONE, dst, src, sizeBytes);
hip::syncStreams();
amd::HostQueue* queue = hip::getNullStream();
@@ -783,7 +783,7 @@ 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_API(dst, src, sizeBytes, kind, stream);
HIP_INIT_API(hipMemcpyAsync, dst, src, sizeBytes, kind, stream);
amd::HostQueue* queue = hip::getQueue(stream);
@@ -793,7 +793,7 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes,
hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes,
hipStream_t stream) {
HIP_INIT_API(dst, src, sizeBytes, stream);
HIP_INIT_API(hipMemcpyHtoDAsync, dst, src, sizeBytes, stream);
amd::HostQueue* queue = hip::getQueue(stream);
@@ -803,7 +803,7 @@ hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes,
hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes,
hipStream_t stream) {
HIP_INIT_API(dst, src, sizeBytes, stream);
HIP_INIT_API(hipMemcpyDtoDAsync, dst, src, sizeBytes, stream);
amd::HostQueue* queue = hip::getQueue(stream);
@@ -813,7 +813,7 @@ hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t siz
hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes,
hipStream_t stream) {
HIP_INIT_API(dst, src, sizeBytes, stream);
HIP_INIT_API(hipMemcpyDtoHAsync, dst, src, sizeBytes, stream);
amd::HostQueue* queue = hip::getQueue(stream);
@@ -888,7 +888,7 @@ hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch
}
hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) {
HIP_INIT_API(pCopy);
HIP_INIT_API(hipMemcpyParam2D, pCopy);
hipError_t e = hipSuccess;
if (pCopy == nullptr) {
e = hipErrorInvalidValue;
@@ -903,7 +903,7 @@ hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) {
hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
size_t height, hipMemcpyKind kind) {
HIP_INIT_API(dst, dpitch, src, spitch, width, height, kind);
HIP_INIT_API(hipMemcpy2D, dst, dpitch, src, spitch, width, height, kind);
hip::syncStreams();
amd::HostQueue* queue = hip::getNullStream();
@@ -914,7 +914,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_API(dst, dpitch, src, spitch, width, height, kind, stream);
HIP_INIT_API(hipMemcpy2DAsync, dst, dpitch, src, spitch, width, height, kind, stream);
amd::HostQueue* queue = hip::getQueue(stream);
@@ -923,7 +923,7 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp
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_API(dst, wOffset, hOffset, src, spitch, width, height, kind);
HIP_INIT_API(hipMemcpy2DToArray, dst, wOffset, hOffset, src, spitch, width, height, kind);
if (dst->data == nullptr) {
HIP_RETURN(hipErrorUnknown);
@@ -1018,7 +1018,7 @@ hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, con
hipError_t hipMemcpyToArray(hipArray* dstArray, size_t wOffset, size_t hOffset, const void* src,
size_t count, hipMemcpyKind kind) {
HIP_INIT_API(dstArray, wOffset, hOffset, src, count, kind);
HIP_INIT_API(hipMemcpyToArray, dstArray, wOffset, hOffset, src, count, kind);
hip::syncStreams();
amd::HostQueue* queue = hip::getNullStream();
@@ -1058,7 +1058,7 @@ hipError_t hipMemcpyToArray(hipArray* dstArray, size_t wOffset, size_t hOffset,
hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffset, size_t hOffset,
size_t count, hipMemcpyKind kind) {
HIP_INIT_API(dst, srcArray, wOffset, hOffset, count, kind);
HIP_INIT_API(hipMemcpyFromArray, dst, srcArray, wOffset, hOffset, count, kind);
hip::syncStreams();
amd::HostQueue* queue = hip::getNullStream();
@@ -1097,7 +1097,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_API(dstArray, dstOffset, srcHost, count);
HIP_INIT_API(hipMemcpyHtoA, dstArray, dstOffset, srcHost, count);
hip::syncStreams();
amd::HostQueue* queue = hip::getNullStream();
@@ -1134,7 +1134,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_API(dst, srcArray, srcOffset, count);
HIP_INIT_API(hipMemcpyAtoH, dst, srcArray, srcOffset, count);
hip::syncStreams();
amd::HostQueue* queue = hip::getNullStream();
@@ -1171,7 +1171,7 @@ hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t
}
hipError_t hipMemcpy3D(const struct hipMemcpy3DParms* p) {
HIP_INIT_API(p);
HIP_INIT_API(hipMemcpy3D, p);
hip::syncStreams();
amd::HostQueue* queue = hip::getNullStream();
@@ -1325,7 +1325,7 @@ hipError_t ihipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t
}
hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t height) {
HIP_INIT_API(dst, pitch, value, width, height);
HIP_INIT_API(hipMemset2D, dst, pitch, value, width, height);
hip::syncStreams();
amd::HostQueue* queue = hip::getNullStream();
@@ -1334,7 +1334,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_API(dst, pitch, value, width, height, stream);
HIP_INIT_API(hipMemset2DAsync, dst, pitch, value, width, height, stream);
amd::HostQueue* queue = hip::getQueue(stream);
@@ -1342,7 +1342,7 @@ hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value,
}
hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes) {
HIP_INIT_API(dst, value, sizeBytes);
HIP_INIT_API(hipMemsetD8, dst, value, sizeBytes);
if (dst == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
@@ -1377,7 +1377,7 @@ hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes
}
hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* dev_ptr) {
HIP_INIT_API(handle, dev_ptr);
HIP_INIT_API(hipIpcGetMemHandle, handle, dev_ptr);
size_t offset = 0;
amd::Memory* amd_mem_obj = nullptr;
@@ -1408,7 +1408,7 @@ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* dev_ptr) {
}
hipError_t hipIpcOpenMemHandle(void** dev_ptr, hipIpcMemHandle_t handle, unsigned int flags) {
HIP_INIT_API(dev_ptr, &handle, flags);
HIP_INIT_API(hipIpcOpenMemHandle, dev_ptr, &handle, flags);
amd::Memory* amd_mem_obj = nullptr;
amd::Device* device = nullptr;
@@ -1434,7 +1434,7 @@ hipError_t hipIpcOpenMemHandle(void** dev_ptr, hipIpcMemHandle_t handle, unsigne
}
hipError_t hipIpcCloseMemHandle(void* dev_ptr) {
HIP_INIT_API(dev_ptr);
HIP_INIT_API(hipIpcCloseMemHandle, dev_ptr);
size_t offset = 0;
amd::Device* device = nullptr;
@@ -1479,7 +1479,7 @@ hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannel
}
hipError_t hipHostGetDevicePointer(void** devicePointer, void* hostPointer, unsigned flags) {
HIP_INIT_API(devicePointer, hostPointer, flags);
HIP_INIT_API(hipHostGetDevicePointer, devicePointer, hostPointer, flags);
size_t offset = 0;
@@ -1493,7 +1493,7 @@ hipError_t hipHostGetDevicePointer(void** devicePointer, void* hostPointer, unsi
}
hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void* ptr) {
HIP_INIT_API(attributes, ptr);
HIP_INIT_API(hipPointerGetAttributes, attributes, ptr);
size_t offset = 0;
amd::Memory* memObj = getMemoryObject(ptr, offset);
projects/hip/api/hip/hip_module.cpp
+20 -15
View File
@@ -30,6 +30,11 @@ THE SOFTWARE.
hipError_t ihipModuleLoadData(hipModule_t *module, const void *image);
const std::string& FunctionName(const hipFunction_t f)
{
return hip::Function::asFunction(f)->function_->name();
}
static uint64_t ElfSize(const void *emi)
{
const Elf64_Ehdr *ehdr = (const Elf64_Ehdr*)emi;
@@ -53,7 +58,7 @@ static uint64_t ElfSize(const void *emi)
hipError_t hipModuleLoad(hipModule_t* module, const char* fname)
{
HIP_INIT_API(module, fname);
HIP_INIT_API(hipModuleLoad, module, fname);
if (!fname) {
HIP_RETURN(hipErrorInvalidValue);
@@ -81,7 +86,7 @@ bool ihipModuleUnregisterGlobal(hipModule_t hmod) {
hipError_t hipModuleUnload(hipModule_t hmod)
{
HIP_INIT_API(hmod);
HIP_INIT_API(hipModuleUnload, hmod);
if (hmod == nullptr) {
HIP_RETURN(hipErrorUnknown);
@@ -100,7 +105,7 @@ hipError_t hipModuleUnload(hipModule_t hmod)
hipError_t hipModuleLoadData(hipModule_t *module, const void *image)
{
HIP_INIT_API(module, image);
HIP_INIT_API(hipModuleLoadData, module, image);
HIP_RETURN(ihipModuleLoadData(module, image));
}
@@ -200,7 +205,7 @@ hipError_t ihipModuleLoadData(hipModule_t *module, const void *image)
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);
amd::Program* program = as_amd(reinterpret_cast<cl_program>(hmod));
@@ -222,7 +227,7 @@ hipError_t hipModuleGetFunction(hipFunction_t *hfunc, hipModule_t hmod, const ch
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);
/* Get address and size for the global symbol */
if (!PlatformState::instance().getGlobalVar(name, ihipGetDevice(), dptr,
@@ -235,7 +240,7 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t h
hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func)
{
HIP_INIT_API(attr, func);
HIP_INIT_API(hipFuncGetAttributes, attr, func);
if (!PlatformState::instance().getFuncAttr(func, attr)) {
HIP_RETURN(hipErrorUnknown);
@@ -253,7 +258,7 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f,
hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags = 0,
uint32_t params = 0)
{
HIP_INIT_API(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ,
HIP_INIT_API(NONE, f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ,
sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, flags, params);
hip::Function* function = hip::Function::asFunction(f);
@@ -347,7 +352,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
uint32_t sharedMemBytes, hipStream_t hStream,
void **kernelParams, void **extra)
{
HIP_INIT_API(f, gridDimX, gridDimY, gridDimZ,
HIP_INIT_API(hipModuleLaunchKernel, f, gridDimX, gridDimY, gridDimZ,
blockDimX, blockDimY, blockDimZ,
sharedMemBytes, hStream,
kernelParams, extra);
@@ -364,7 +369,7 @@ hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
hipStream_t hStream, void** kernelParams, void** extra,
hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags)
{
HIP_INIT_API(f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ,
HIP_INIT_API(NONE, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ,
localWorkSizeX, localWorkSizeY, localWorkSizeZ,
sharedMemBytes, hStream,
kernelParams, extra, startEvent, stopEvent, flags);
@@ -383,7 +388,7 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t gridDimX,
hipEvent_t startEvent,
hipEvent_t stopEvent)
{
HIP_INIT_API(f, gridDimX, gridDimY, gridDimZ,
HIP_INIT_API(NONE, f, gridDimX, gridDimY, gridDimZ,
blockDimX, blockDimY, blockDimZ,
sharedMemBytes, hStream,
kernelParams, extra, startEvent, stopEvent);
@@ -400,7 +405,7 @@ hipError_t hipModuleLaunchKernelExt(hipFunction_t f, uint32_t gridDimX,
hipEvent_t startEvent,
hipEvent_t stopEvent)
{
HIP_INIT_API(f, gridDimX, gridDimY, gridDimZ,
HIP_INIT_API(NONE, f, gridDimX, gridDimY, gridDimZ,
blockDimX, blockDimY, blockDimZ,
sharedMemBytes, hStream,
kernelParams, extra, startEvent, stopEvent);
@@ -413,7 +418,7 @@ hipError_t hipLaunchCooperativeKernel(const void* f,
dim3 gridDim, dim3 blockDim,
void **kernelParams, uint32_t sharedMemBytes, hipStream_t hStream)
{
HIP_INIT_API(f, gridDim, blockDim,
HIP_INIT_API(hipLaunchCooperativeKernel, f, gridDim, blockDim,
sharedMemBytes, hStream);
int deviceId = ihipGetDevice();
@@ -464,7 +469,7 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL
hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList,
int numDevices, unsigned int flags)
{
HIP_INIT_API(launchParamsList, numDevices, flags);
HIP_INIT_API(hipLaunchCooperativeKernelMultiDevice, launchParamsList, numDevices, flags);
return ihipLaunchCooperativeKernelMultiDevice(launchParamsList, numDevices, flags,
(amd::NDRangeKernelCommand::CooperativeGroups |
@@ -473,13 +478,13 @@ hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsLi
hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList,
int numDevices, unsigned int flags) {
HIP_INIT_API(launchParamsList, numDevices, flags);
HIP_INIT_API(hipExtLaunchMultiKernelMultiDevice, launchParamsList, numDevices, flags);
return ihipLaunchCooperativeKernelMultiDevice(launchParamsList, numDevices, flags, 0);
}
hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const char* name) {
HIP_INIT_API(texRef, hmod, name);
HIP_INIT_API(hipModuleGetTexRef, texRef, hmod, name);
/* input args check */
if ((texRef == nullptr) || (name == nullptr)) {
projects/hip/api/hip/hip_peer.cpp
+10 -10
View File
@@ -25,7 +25,7 @@ THE SOFTWARE.
#include "hip_internal.hpp"
hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, hipCtx_t thisCtx, hipCtx_t peerCtx) {
HIP_INIT_API(canAccessPeer, thisCtx, peerCtx);
HIP_INIT_API(NONE, canAccessPeer, thisCtx, peerCtx);
assert(0 && "Unimplemented");
@@ -34,7 +34,7 @@ hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, hipCtx_t thisCtx, hipCtx_t
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(NONE, dst, dstCtx, src, srcCtx, sizeBytes);
assert(0 && "Unimplemented");
@@ -43,7 +43,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(NONE, dst, dstDevice, src, srcDevice, sizeBytes, stream);
assert(0 && "Unimplemented");
@@ -51,7 +51,7 @@ 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);
amd::Device* device = nullptr;
amd::Device* peer_device = nullptr;
@@ -83,39 +83,39 @@ hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, int deviceId, int peerDevi
}
hipError_t hipDeviceDisablePeerAccess(int peerDeviceId) {
HIP_INIT_API(peerDeviceId);
HIP_INIT_API(hipDeviceDisablePeerAccess, peerDeviceId);
HIP_RETURN(hipSuccess);
}
hipError_t hipDeviceEnablePeerAccess(int peerDeviceId, unsigned int flags) {
HIP_INIT_API(peerDeviceId, flags);
HIP_INIT_API(hipDeviceEnablePeerAccess, peerDeviceId, flags);
HIP_RETURN(hipSuccess);
}
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);
HIP_RETURN(hipMemcpy(dst, src, sizeBytes, hipMemcpyDeviceToDevice));
}
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);
HIP_RETURN(hipMemcpyAsync(dst, src, sizeBytes, hipMemcpyDeviceToDevice, stream));
}
hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags) {
HIP_INIT_API(peerCtx, flags);
HIP_INIT_API(hipCtxEnablePeerAccess, peerCtx, flags);
HIP_RETURN(hipSuccess);
}
hipError_t hipCtxDisablePeerAccess(hipCtx_t peerCtx) {
HIP_INIT_API(peerCtx);
HIP_INIT_API(hipCtxDisablePeerAccess, peerCtx);
HIP_RETURN(hipSuccess);
}
projects/hip/api/hip/hip_platform.cpp
+7 -7
View File
@@ -425,7 +425,7 @@ extern "C" hipError_t hipConfigureCall(
size_t sharedMem,
hipStream_t stream)
{
HIP_INIT_API(gridDim, blockDim, sharedMem, stream);
HIP_INIT_API(NONE, gridDim, blockDim, sharedMem, stream);
PlatformState::instance().configureCall(gridDim, blockDim, sharedMem, stream);
@@ -438,7 +438,7 @@ extern "C" hipError_t __hipPushCallConfiguration(
size_t sharedMem,
hipStream_t stream)
{
HIP_INIT_API(gridDim, blockDim, sharedMem, stream);
HIP_INIT_API(NONE, gridDim, blockDim, sharedMem, stream);
PlatformState::instance().configureCall(gridDim, blockDim, sharedMem, stream);
@@ -449,7 +449,7 @@ extern "C" hipError_t __hipPopCallConfiguration(dim3 *gridDim,
dim3 *blockDim,
size_t *sharedMem,
hipStream_t *stream) {
HIP_INIT_API(gridDim, blockDim, sharedMem, stream);
HIP_INIT_API(NONE, gridDim, blockDim, sharedMem, stream);
ihipExec_t exec;
PlatformState::instance().popExec(exec);
@@ -466,7 +466,7 @@ extern "C" hipError_t hipSetupArgument(
size_t size,
size_t offset)
{
HIP_INIT_API(arg, size, offset);
HIP_INIT_API(NONE, arg, size, offset);
PlatformState::instance().setupArgument(arg, size, offset);
@@ -475,7 +475,7 @@ extern "C" hipError_t hipSetupArgument(
extern "C" hipError_t hipLaunchByPtr(const void *hostFunction)
{
HIP_INIT_API(hostFunction);
HIP_INIT_API(NONE, hostFunction);
int deviceId = ihipGetDevice();
hipFunction_t func = PlatformState::instance().getFunc(hostFunction, deviceId);
@@ -506,7 +506,7 @@ extern "C" hipError_t hipLaunchKernel(const void *hostFunction,
size_t sharedMemBytes,
hipStream_t stream)
{
HIP_INIT_API(hostFunction, gridDim, blockDim, args, sharedMemBytes,
HIP_INIT_API(NONE, hostFunction, gridDim, blockDim, args, sharedMemBytes,
stream);
int deviceId = ihipGetDevice();
@@ -567,7 +567,7 @@ hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor(int* numBlocks,
int blockSize,
size_t dynamicSMemSize)
{
HIP_INIT_API(f, blockSize, dynamicSMemSize);
HIP_INIT_API(NONE, f, blockSize, dynamicSMemSize);
int deviceId = ihipGetDevice();
hipFunction_t func = PlatformState::instance().getFunc(f, deviceId);
if (func == nullptr) {
projects/hip/api/hip/hip_prof_api.h
+255
View File
@@ -0,0 +1,255 @@
/*
Copyright (c) 2019 - 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.
*/
#ifndef HIP_SRC_HIP_PROF_API_H
#define HIP_SRC_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;
enabled_ = true;
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;
enabled_ = true;
cb_release(id);
} else {
ret = false;
}
return ret;
}
void set_enabled(const bool& enabled) {
enabled_ = enabled;
}
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);
}
inline bool is_enabled() const {
return enabled_;
}
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_;
bool enabled_;
};
#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) \
api_callbacks_spawner_t<HIP_API_ID_##CB_ID> __api_tracer(HIP_API_ID_##CB_ID); \
{ \
hip_api_data_t* api_data = __api_tracer.get_api_data_ptr(); \
if (api_data != NULL) { \
hip_api_data_t& api_data_ref = *api_data; \
INIT_CB_ARGS_DATA(CB_ID, api_data_ref); \
__api_tracer.call(); \
} \
}
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) :
api_data_(NULL)
{
if (!is_enabled()) return;
if (cid_ >= HIP_API_ID_NUMBER) {
fprintf(stderr, "HIP %s bad id %d\n", __FUNCTION__, cid_);
abort();
}
callbacks_table.sem_sync(cid_);
hip_act_callback_t act = entry(cid_).act;
if (act != NULL) api_data_ = (hip_api_data_t*) act(cid_, NULL, NULL, NULL);
}
void call() {
hip_api_callback_t fun = entry(cid_).fun;
void* arg = entry(cid_).arg;
if (fun != NULL) fun(HIP_DOMAIN_ID, cid_, api_data_, arg);
}
~api_callbacks_spawner_t() {
if (api_data_ == NULL) return;
hip_act_callback_t act = entry(cid_).act;
void* a_arg = entry(cid_).a_arg;
hip_api_callback_t fun = entry(cid_).fun;
void* arg = entry(cid_).arg;
if (act != NULL) act(cid_, NULL, api_data_, a_arg);
if (fun != NULL) fun(HIP_DOMAIN_ID, cid_, api_data_, arg);
callbacks_table.sem_release(cid_);
}
hip_api_data_t* get_api_data_ptr() {
return api_data_;
}
bool is_enabled() const {
return callbacks_table.is_enabled();
}
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_;
};
template <>
class api_callbacks_spawner_t<HIP_API_ID_NUMBER> {
public:
api_callbacks_spawner_t(const hip_api_id_t& cid) {}
void call() {}
hip_api_data_t* get_api_data_ptr() { return NULL; }
bool is_enabled() const { return false; }
};
#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 false; }
bool set_callback(uint32_t id, fun_t fun, void* arg) { return false; }
};
#endif
#endif // HIP_SRC_HIP_PROF_API_H
projects/hip/api/hip/hip_profile.cpp
+2 -2
View File
@@ -25,7 +25,7 @@ THE SOFTWARE.
#include "hip_internal.hpp"
hipError_t hipProfilerStart() {
HIP_INIT_API();
HIP_INIT_API(hipProfilerStart);
assert(0 && "Unimplemented");
@@ -34,7 +34,7 @@ hipError_t hipProfilerStart() {
hipError_t hipProfilerStop() {
HIP_INIT_API();
HIP_INIT_API(hipProfilerStop);
assert(0 && "Unimplemented");
projects/hip/api/hip/hip_stream.cpp
+10 -10
View File
@@ -120,19 +120,19 @@ static hipError_t ihipStreamCreate(hipStream_t *stream, unsigned int flags, amd:
}
hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) {
HIP_INIT_API(stream, flags);
HIP_INIT_API(hipStreamCreateWithFlags, stream, flags);
HIP_RETURN(ihipStreamCreate(stream, flags, amd::CommandQueue::Priority::Normal));
}
hipError_t hipStreamCreate(hipStream_t *stream) {
HIP_INIT_API(stream);
HIP_INIT_API(hipStreamCreate, stream);
HIP_RETURN(ihipStreamCreate(stream, hipStreamDefault, amd::CommandQueue::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);
if (priority > static_cast<int>(amd::CommandQueue::Priority::High)) {
priority = static_cast<int>(amd::CommandQueue::Priority::High);
@@ -144,7 +144,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 != nullptr) {
*leastPriority = static_cast<int>(amd::CommandQueue::Priority::Normal);
@@ -157,7 +157,7 @@ hipError_t hipDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPrio
}
hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags) {
HIP_INIT_API(stream, flags);
HIP_INIT_API(hipStreamGetFlags, stream, flags);
hip::Stream* hStream = reinterpret_cast<hip::Stream*>(stream);
@@ -171,7 +171,7 @@ hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags) {
}
hipError_t hipStreamSynchronize(hipStream_t stream) {
HIP_INIT_API(stream);
HIP_INIT_API(hipStreamSynchronize, stream);
amd::HostQueue* hostQueue = hip::getQueue(stream);
hostQueue->finish();
@@ -180,7 +180,7 @@ hipError_t hipStreamSynchronize(hipStream_t stream) {
}
hipError_t hipStreamDestroy(hipStream_t stream) {
HIP_INIT_API(stream);
HIP_INIT_API(hipStreamDestroy, stream);
if (stream == nullptr) {
HIP_RETURN(hipErrorInvalidResourceHandle);
@@ -199,7 +199,7 @@ hipError_t hipStreamDestroy(hipStream_t stream) {
}
hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags) {
HIP_INIT_API(stream, event, flags);
HIP_INIT_API(hipStreamWaitEvent, stream, event, flags);
amd::HostQueue* queue;
@@ -219,7 +219,7 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
}
hipError_t hipStreamQuery(hipStream_t stream) {
HIP_INIT_API(stream);
HIP_INIT_API(hipStreamQuery, stream);
amd::HostQueue* hostQueue;
if (stream == nullptr) {
@@ -242,7 +242,7 @@ hipError_t hipStreamQuery(hipStream_t stream) {
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);
amd::HostQueue* hostQueue = reinterpret_cast<hip::Stream*>
(stream)->asHostQueue();
projects/hip/api/hip/hip_surface.cpp
+2 -2
View File
@@ -47,7 +47,7 @@ using namespace hip;
hipError_t hipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject,
const hipResourceDesc* pResDesc) {
HIP_INIT_API(pSurfObject, pResDesc);
HIP_INIT_API(NONE, pSurfObject, pResDesc);
hipSurface* pSurface = new hipSurface(pResDesc);
assert(pSurface != nullptr);
@@ -82,7 +82,7 @@ hipError_t hipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject,
hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject) {
HIP_INIT_API(surfaceObject);
HIP_INIT_API(NONE, surfaceObject);
amd::ScopedLock lock(surfaceLock);
hipSurface* pSurface = surfaceHash[surfaceObject];
projects/hip/api/hip/hip_texture.cpp
+25 -25
View File
@@ -215,7 +215,7 @@ hip::TextureObject* ihipCreateTextureObject(const hipResourceDesc& resDesc, amd:
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(NONE, pTexObject, pResDesc, pTexDesc, pResViewDesc);
amd::Device* device = hip::getCurrentContext()->devices()[0];
@@ -307,7 +307,7 @@ void ihipDestroyTextureObject(hip::TextureObject* texture) {
}
hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject) {
HIP_INIT_API(textureObject);
HIP_INIT_API(NONE, textureObject);
hip::TextureObject* texture = reinterpret_cast<hip::TextureObject*>(textureObject);
@@ -318,7 +318,7 @@ hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject) {
hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDesc,
hipTextureObject_t textureObject) {
HIP_INIT_API(pResDesc, textureObject);
HIP_INIT_API(NONE, pResDesc, textureObject);
hip::TextureObject* texture = reinterpret_cast<hip::TextureObject*>(textureObject);
@@ -331,7 +331,7 @@ hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDesc,
hipError_t hipGetTextureObjectResourceViewDesc(hipResourceViewDesc* pResViewDesc,
hipTextureObject_t textureObject) {
HIP_INIT_API(pResViewDesc, textureObject);
HIP_INIT_API(NONE, pResViewDesc, textureObject);
assert(0 && "Unimplemented");
@@ -340,7 +340,7 @@ hipError_t hipGetTextureObjectResourceViewDesc(hipResourceViewDesc* pResViewDesc
hipError_t hipGetTextureObjectTextureDesc(hipTextureDesc* pTexDesc,
hipTextureObject_t textureObject) {
HIP_INIT_API(pTexDesc, textureObject);
HIP_INIT_API(NONE, pTexDesc, textureObject);
assert(0 && "Unimplemented");
@@ -416,7 +416,7 @@ hipError_t ihipBindTexture(cl_mem_object_type type,
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(NONE, offset, tex, devPtr, desc, size);
if (desc == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
@@ -432,14 +432,14 @@ hipError_t hipBindTexture(size_t* offset, textureReference* tex, const void* dev
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(NONE, offset, tex, devPtr, desc, width, height, pitch);
HIP_RETURN(ihipBindTexture(CL_MEM_OBJECT_IMAGE2D, offset, tex, devPtr, desc, width, height, pitch));
}
hipError_t hipBindTextureToArray(textureReference* tex, hipArray_const_t array,
const hipChannelFormatDesc* desc) {
HIP_INIT_API(tex, array, desc);
HIP_INIT_API(NONE, tex, array, desc);
assert(0 && "Unimplemented");
@@ -449,7 +449,7 @@ hipError_t hipBindTextureToArray(textureReference* tex, hipArray_const_t array,
hipError_t ihipBindTextureImpl(TlsData* tls, int dim, enum hipTextureReadMode readMode, size_t* offset,
const void* devPtr, const struct hipChannelFormatDesc* desc,
size_t size, textureReference* tex) {
HIP_INIT_API(dim, readMode, offset, devPtr, size, tex);
HIP_INIT_API(NONE, dim, readMode, offset, devPtr, size, tex);
assert(1 == dim);
@@ -460,7 +460,7 @@ hipError_t ihipBindTextureToArrayImpl(TlsData* tls, int dim, enum hipTextureRead
hipArray_const_t array,
const struct hipChannelFormatDesc& desc,
textureReference* tex) {
HIP_INIT_API(dim, readMode, &desc, array, tex);
HIP_INIT_API(NONE, dim, readMode, &desc, array, tex);
cl_mem_object_type clType;
size_t offset = 0;
@@ -483,7 +483,7 @@ hipError_t ihipBindTextureToArrayImpl(TlsData* tls, int dim, enum hipTextureRead
hipError_t hipBindTextureToMipmappedArray(textureReference* tex,
hipMipmappedArray_const_t mipmappedArray,
const hipChannelFormatDesc* desc) {
HIP_INIT_API(tex, mipmappedArray, desc);
HIP_INIT_API(NONE, tex, mipmappedArray, desc);
assert(0 && "Unimplemented");
@@ -491,7 +491,7 @@ hipError_t hipBindTextureToMipmappedArray(textureReference* tex,
}
hipError_t hipUnbindTexture(const textureReference* tex) {
HIP_INIT_API(tex);
HIP_INIT_API(NONE, tex);
ihipDestroyTextureObject(reinterpret_cast<hip::TextureObject*>(tex->textureObject));
@@ -499,7 +499,7 @@ hipError_t hipUnbindTexture(const textureReference* tex) {
}
hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_const_t array) {
HIP_INIT_API(desc, array);
HIP_INIT_API(NONE, desc, array);
if (desc != nullptr) {
*desc = array->desc;
@@ -509,7 +509,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(NONE, offset, tex);
if ((offset == nullptr) || (tex == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
@@ -521,7 +521,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(NONE, tex, symbol);
assert(0 && "Unimplemented");
@@ -529,7 +529,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(NONE, tex, fmt, NumPackedComponents);
if (tex == nullptr) {
HIP_RETURN(hipErrorInvalidImage);
@@ -542,7 +542,7 @@ 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(NONE, tex, flags);
if (tex == nullptr) {
HIP_RETURN(hipErrorInvalidImage);
@@ -554,7 +554,7 @@ hipError_t hipTexRefSetFlags(textureReference* tex, unsigned int flags) {
}
hipError_t hipTexRefSetFilterMode(textureReference* tex, hipTextureFilterMode fm) {
HIP_INIT_API(tex, fm);
HIP_INIT_API(NONE, tex, fm);
if (tex == nullptr) {
HIP_RETURN(hipErrorInvalidImage);
@@ -566,7 +566,7 @@ hipError_t hipTexRefSetFilterMode(textureReference* tex, hipTextureFilterMode fm
}
hipError_t hipTexRefGetAddressMode(hipTextureAddressMode* am, textureReference tex, int dim) {
HIP_INIT_API(am, &tex, dim);
HIP_INIT_API(NONE, am, &tex, dim);
if ((am == nullptr) || (dim >= 3)) {
HIP_RETURN(hipErrorInvalidValue);
@@ -578,7 +578,7 @@ hipError_t hipTexRefGetAddressMode(hipTextureAddressMode* am, textureReference t
}
hipError_t hipTexRefSetAddressMode(textureReference* tex, int dim, hipTextureAddressMode am) {
HIP_INIT_API(tex, dim, am);
HIP_INIT_API(NONE, tex, dim, am);
if (tex == nullptr) {
HIP_RETURN(hipErrorInvalidImage);
@@ -590,7 +590,7 @@ hipError_t hipTexRefSetAddressMode(textureReference* tex, int dim, hipTextureAdd
}
hipError_t hipTexRefGetArray(hipArray_t* array, textureReference tex) {
HIP_INIT_API(array, &tex);
HIP_INIT_API(NONE, array, &tex);
hip::TextureObject* texture = nullptr;
@@ -613,7 +613,7 @@ hipError_t hipTexRefGetArray(hipArray_t* array, textureReference tex) {
}
hipError_t hipTexRefSetArray(textureReference* tex, hipArray_const_t array, unsigned int flags) {
HIP_INIT_API(tex, array, flags);
HIP_INIT_API(NONE, tex, array, flags);
size_t offset = 0;
@@ -626,7 +626,7 @@ hipError_t hipTexRefSetArray(textureReference* tex, hipArray_const_t array, unsi
}
hipError_t hipTexRefGetAddress(hipDeviceptr_t* dev_ptr, textureReference tex) {
HIP_INIT_API(dev_ptr, &tex);
HIP_INIT_API(NONE, dev_ptr, &tex);
hip::TextureObject* texture = nullptr;
device::Memory* dev_mem = nullptr;
@@ -648,7 +648,7 @@ hipError_t hipTexRefGetAddress(hipDeviceptr_t* dev_ptr, textureReference tex) {
hipError_t hipTexRefSetAddress(size_t* offset, textureReference* tex, hipDeviceptr_t devPtr,
size_t size) {
HIP_INIT_API(offset, tex, devPtr, size);
HIP_INIT_API(NONE, offset, tex, devPtr, size);
if (tex == nullptr) {
HIP_RETURN(hipErrorInvalidImage);
@@ -664,7 +664,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(NONE, tex, desc, devPtr, pitch);
if (desc == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
projects/hip/api/hip/prof_protocol.h
+90
View File
@@ -0,0 +1,90 @@
/*
Copyright (c) 2019 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.
*/
#ifndef INC_EXT_PROF_PROTOCOL_H_
#define INC_EXT_PROF_PROTOCOL_H_
#include <stdlib.h>
// Traced API domains
typedef enum {
ACTIVITY_DOMAIN_HSA_API = 0, // HSA API domain
ACTIVITY_DOMAIN_HSA_OPS = 1, // HSA async activity domain
ACTIVITY_DOMAIN_HCC_OPS = 2, // HCC async activity domain
ACTIVITY_DOMAIN_HIP_API = 3, // HIP API domain
ACTIVITY_DOMAIN_EXT_API = 4, // External ID domain
ACTIVITY_DOMAIN_ROCTX = 5, // ROCTX domain
ACTIVITY_DOMAIN_NUMBER
} activity_domain_t;
// Extension API opcodes
typedef enum {
ACTIVITY_EXT_OP_MARK = 0,
ACTIVITY_EXT_OP_EXTERN_ID = 1
} activity_ext_op_t;
// API calback type
typedef void (*activity_rtapi_callback_t)(uint32_t domain, uint32_t cid, const void* data, void* arg);
typedef uint32_t activity_kind_t;
typedef uint32_t activity_op_t;
// API callback phase
typedef enum {
ACTIVITY_API_PHASE_ENTER = 0,
ACTIVITY_API_PHASE_EXIT = 1
} activity_api_phase_t;
// Trace record types
// Correlation id
typedef uint64_t activity_correlation_id_t;
// Activity record type
struct activity_record_t {
uint32_t domain; // activity domain id
activity_kind_t kind; // activity kind
activity_op_t op; // activity op
activity_correlation_id_t correlation_id; // activity ID
uint64_t begin_ns; // host begin timestamp
uint64_t end_ns; // host end timestamp
union {
struct {
int device_id; // device id
uint64_t queue_id; // queue id
};
struct {
uint32_t process_id; // device id
uint32_t thread_id; // thread id
};
struct {
activity_correlation_id_t external_id; // external correlatino id
};
};
size_t bytes; // data size bytes
};
// Activity sync calback type
typedef void* (*activity_sync_callback_t)(uint32_t cid, activity_record_t* record, const void* data, void* arg);
// Activity async calback type
typedef void (*activity_id_callback_t)(activity_correlation_id_t id);
typedef void (*activity_async_callback_t)(uint32_t op, void* record, void* arg);
#endif // INC_EXT_PROF_PROTOCOL_H_