AI/rocm-systems
2
0
Tükrözés 0
Kód Hibajegyek Egyesítési kérések Actions Packages Projektek Kiadások Wiki Tevékenység

Add HIP_TRACE_API=4. Only display memory allocation/free apis.

Forráskód böngészése 
[ROCm/clr commit: 7e7ba5027f]
This commit is contained in:
Ben Sander
2017-01-27 11:21:08 -06:00
szülő 9a097161ac
commit ee05975efa
5 fájl változott, egészen pontosan 39 új sor hozzáadva és 31 régi sor törölve
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/README.md
+1
Fájl megtekintése
@@ -32,6 +32,7 @@ HIP releases are typically of two types. The tag naming convention is different
- [HIP Runtime API (Doxygen)](http://gpuopen-professionalcompute-tools.github.io/HIP)
- [HIP Porting Guide](docs/markdown/hip_porting_guide.md)
- [HIP Porting Driver Guide](docs/markdown/hip_porting_driver_api.md)
- [HIP Profiling and Debugging](docs/markdown/hip_profiling.md)
- [HIP Terminology](docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL)
- [hipify-clang](hipify-clang/README.md)
- [Developer/CONTRIBUTING Info](CONTRIBUTING.md)
projects/clr/hipamd/docs/markdown/hip_profiling.md
+5
Fájl megtekintése
@@ -267,6 +267,11 @@ info: check result
PASSED!
```
HIP_TRACE_API supports multiple levels of debug information:
- 0x1 = print all HIP APIs
- 0x2 = print HIP APIs which initiate GPU kernels, copies, or memsets. Includes hipLaunchKernel, hipMemcpy*, hipMemset*.
- 0x4 = print HIP APIs which allocate or free memory. Includes hipMalloc, hipHostMalloc, hipFree, hipHostFree.
#### Color
Note this trace mode uses colors. "less -r" can handle raw control characters and will display the debug output in proper colors.
projects/clr/hipamd/src/hip_hcc.cpp
+1 -1
Fájl megtekintése
@@ -1435,7 +1435,7 @@ hipStream_t ihipSyncAndResolveStream(hipStream_t stream)
void ihipPrintKernelLaunch(const char *kernelName, const grid_launch_parm *lp, const hipStream_t stream)
{
if ((HIP_TRACE_API & (1<<TRACE_CMD)) || HIP_PROFILE_API || (COMPILE_HIP_DB && HIP_TRACE_API)) {
if ((HIP_TRACE_API & (1<<TRACE_KCMD)) || HIP_PROFILE_API || (COMPILE_HIP_DB & HIP_TRACE_API)) {
std::stringstream os_pre;
std::stringstream os;
os_pre << "<<hip-api tid:";
projects/clr/hipamd/src/hip_hcc_internal.h
+9 -7
Fájl megtekintése
@@ -188,9 +188,10 @@ extern const char *API_COLOR_END;
//---
//HIP Trace modes
#define TRACE_ALL 0 // 0x1
#define TRACE_CMD 1 // 0x2
#define TRACE_MEM 2 // 0x4
#define TRACE_ALL 0 // 0x1
#define TRACE_KCMD 1 // 0x2, kernel command
#define TRACE_MCMD 2 // 0x4, memory command
#define TRACE_MEM 3 // 0x8
//---
@@ -275,12 +276,13 @@ extern void recordApiTrace(std::string *fullStr, const std::string &apiStr);
API_TRACE(0, __VA_ARGS__);
// Like above, but will trace with DB_CMD.
// Replace HIP_INIT_API with this call inside important APIs that launch work on the GPU:
// Like above, but will trace with TRACE_CMD.
// 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_CMD_API(...) \
#define HIP_INIT_SPECIAL_API(tbit, ...) \
HIP_INIT()\
API_TRACE((HIP_TRACE_API&(1<<TRACE_CMD)), __VA_ARGS__);
API_TRACE((HIP_TRACE_API&(1<<tbit)), __VA_ARGS__);
// 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,
projects/clr/hipamd/src/hip_memory.cpp
+23 -23
Fájl megtekintése
@@ -213,7 +213,7 @@ hipError_t hipHostGetDevicePointer(void **devicePointer, void *hostPointer, unsi
hipError_t hipMalloc(void** ptr, size_t sizeBytes)
{
HIP_INIT_API(ptr, sizeBytes);
HIP_INIT_SPECIAL_API((TRACE_MEM), ptr, sizeBytes);
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
@@ -244,7 +244,7 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes)
hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
{
HIP_INIT_CMD_API(ptr, sizeBytes, flags);
HIP_INIT_SPECIAL_API((TRACE_MCMD), ptr, sizeBytes, flags);
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
@@ -308,7 +308,7 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags)
// width in bytes
hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height)
{
HIP_INIT_CMD_API(ptr, pitch, width, height);
HIP_INIT_SPECIAL_API((TRACE_MEM), ptr, pitch, width, height);
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
@@ -349,7 +349,7 @@ hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannel
hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc,
size_t width, size_t height, unsigned int flags)
{
HIP_INIT_CMD_API(array, desc, width, height, flags);
HIP_INIT_SPECIAL_API((TRACE_MEM), array, desc, width, height, flags);
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
@@ -490,7 +490,7 @@ hipError_t hipHostUnregister(void *hostPtr)
hipError_t hipMemcpyToSymbol(const void* symbolName, const void *src, size_t count, size_t offset, hipMemcpyKind kind)
{
HIP_INIT_CMD_API(symbolName, src, count, offset, kind);
HIP_INIT_SPECIAL_API((TRACE_MCMD), symbolName, src, count, offset, kind);
if(symbolName == nullptr)
{
@@ -560,7 +560,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_CMD_API(symbolName, src, count, offset, kind, stream);
HIP_INIT_SPECIAL_API((TRACE_MCMD), symbolName, src, count, offset, kind, stream);
if(symbolName == nullptr)
{
@@ -637,7 +637,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t co
//---
hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind)
{
HIP_INIT_CMD_API(dst, src, sizeBytes, kind);
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, kind);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
@@ -659,7 +659,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_CMD_API(dst, src, sizeBytes);
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
@@ -681,7 +681,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_CMD_API(dst, src, sizeBytes);
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
@@ -703,7 +703,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_CMD_API(dst, src, sizeBytes);
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
@@ -725,7 +725,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_CMD_API(dst, src, sizeBytes);
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
@@ -750,7 +750,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_CMD_API(dst, src, sizeBytes, kind, stream);
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, kind, stream);
return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, kind, stream));
@@ -759,21 +759,21 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp
hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, hipStream_t stream)
{
HIP_INIT_CMD_API(dst, src, sizeBytes, stream);
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, stream);
return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyHostToDevice, stream));
}
hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream)
{
HIP_INIT_CMD_API(dst, src, sizeBytes, stream);
HIP_INIT_SPECIAL_API((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_CMD_API(dst, src, sizeBytes, stream);
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, stream);
return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyDeviceToHost, stream));
}
@@ -782,7 +782,7 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, h
hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch,
size_t width, size_t height, hipMemcpyKind kind) {
HIP_INIT_CMD_API(dst, dpitch, src, spitch, width, height, kind);
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind);
if(width > dpitch || width > spitch)
return ihipLogStatus(hipErrorUnknown);
@@ -826,7 +826,7 @@ hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch,
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_CMD_API(dst, wOffset, hOffset, src, spitch, width, height, kind);
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, wOffset, hOffset, src, spitch, width, height, kind);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
@@ -879,7 +879,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_CMD_API(dst, wOffset, hOffset, src, count, kind);
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, wOffset, hOffset, src, count, kind);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
@@ -938,7 +938,7 @@ ihipMemsetKernel(hipStream_t stream,
// 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_CMD_API(dst, value, sizeBytes, stream);
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, value, sizeBytes, stream);
hipError_t e = hipSuccess;
@@ -988,7 +988,7 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t s
hipError_t hipMemset(void* dst, int value, size_t sizeBytes )
{
HIP_INIT_CMD_API(dst, value, sizeBytes);
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, value, sizeBytes);
hipError_t e = hipSuccess;
@@ -1148,7 +1148,7 @@ hipError_t hipMemPtrGetInfo(void *ptr, size_t *size)
hipError_t hipFree(void* ptr)
{
HIP_INIT_API(ptr);
HIP_INIT_SPECIAL_API((TRACE_MEM), ptr);
hipError_t hipStatus = hipErrorInvalidDevicePointer;
@@ -1176,7 +1176,7 @@ hipError_t hipFree(void* ptr)
hipError_t hipHostFree(void* ptr)
{
HIP_INIT_API(ptr);
HIP_INIT_SPECIAL_API((TRACE_MEM), ptr);
// Synchronize to ensure all work has finished.
ihipGetTlsDefaultCtx()->locked_waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish.
@@ -1210,7 +1210,7 @@ hipError_t hipFreeHost(void* ptr)
hipError_t hipFreeArray(hipArray* array)
{
HIP_INIT_API(array);
HIP_INIT_SPECIAL_API((TRACE_MEM), array);
hipError_t hipStatus = hipErrorInvalidDevicePointer;