From 2097038ba8be4501ae392741e05d4700144bc0c1 Mon Sep 17 00:00:00 2001 From: rkebichi <54912798+rkebichi@users.noreply.github.com> Date: Wed, 4 Sep 2019 15:46:01 -0400 Subject: [PATCH 1/4] Update roctx.cpp Added roctxRangePushA and roctxRangePop fcts [ROCm/roctracer commit: 226240ea4039c8b0fa7d51996f6d2ceb5dcc1322] --- projects/roctracer/src/roctx/roctx.cpp | 24 ++++++++++++++++++++++-- 1 file changed, 22 insertions(+), 2 deletions(-) diff --git a/projects/roctracer/src/roctx/roctx.cpp b/projects/roctracer/src/roctx/roctx.cpp index 1799b17910..b7faab91e3 100644 --- a/projects/roctracer/src/roctx/roctx.cpp +++ b/projects/roctracer/src/roctx/roctx.cpp @@ -28,6 +28,7 @@ THE SOFTWARE. #include "inc/ext/prof_protocol.h" #include "util/exception.h" #include "util/logger.h" +#include #define PUBLIC_API __attribute__((visibility("default"))) #define CONSTRUCTOR_API __attribute__((constructor)) @@ -61,6 +62,8 @@ THE SOFTWARE. (void)err; \ return X; +std::stack message_stack; + static inline uint32_t GetPid() { return syscall(__NR_getpid); } static inline uint32_t GetTid() { return syscall(__NR_gettid); } @@ -114,13 +117,30 @@ PUBLIC_API void roctxMarkA(const char* message) { PUBLIC_API int roctxRangePushA(const char* message) { API_METHOD_PREFIX - EXC_ABORT(ROCTX_STATUS_ERROR, "method is not implemented"); + //EXC_ABORT(ROCTX_STATUS_ERROR, "method is not implemented"); + roctx_api_data_t api_data{}; + api_data.args.roctxRangePushA.message = strdup(message); + activity_rtapi_callback_t api_callback_fun = NULL; + void* api_callback_arg = NULL; + roctx::cb_table.get(ROCTX_API_ID_roctxRangePushA, &api_callback_fun, &api_callback_arg); + if (api_callback_fun) api_callback_fun(ACTIVITY_DOMAIN_ROCTX, ROCTX_API_ID_roctxRangePushA, &api_data, api_callback_arg); + message_stack.push(strdup(message)); API_METHOD_SUFFIX } PUBLIC_API int roctxRangePop() { API_METHOD_PREFIX - EXC_ABORT(ROCTX_STATUS_ERROR, "method is not implemented"); + //EXC_ABORT(ROCTX_STATUS_ERROR, "method is not implemented"); + roctx_api_data_t api_data{}; + activity_rtapi_callback_t api_callback_fun = NULL; + void* api_callback_arg = NULL; + roctx::cb_table.get(ROCTX_API_ID_roctxRangePop, &api_callback_fun, &api_callback_arg); + if (api_callback_fun) api_callback_fun(ACTIVITY_DOMAIN_ROCTX, ROCTX_API_ID_roctxRangePop, &api_data, api_callback_arg); + if (message_stack.empty()) { + EXC_ABORT(ROCTX_STATUS_ERROR, "Pop from empty stack!"); + } else { + message_stack.pop(); + } API_METHOD_SUFFIX } From 85b8b91e56141299f84cb016b309b918bcf85cb6 Mon Sep 17 00:00:00 2001 From: rkebichi <54912798+rkebichi@users.noreply.github.com> Date: Wed, 4 Sep 2019 15:56:46 -0400 Subject: [PATCH 2/4] Update MatrixTranspose.cpp [ROCm/roctracer commit: c6a035577fd87a97b9721fe2943d14a239b06e60] --- .../MatrixTranspose_test/MatrixTranspose.cpp | 22 ++++++++++++++++++- 1 file changed, 21 insertions(+), 1 deletion(-) diff --git a/projects/roctracer/test/MatrixTranspose_test/MatrixTranspose.cpp b/projects/roctracer/test/MatrixTranspose_test/MatrixTranspose.cpp index dffb0eb984..b519801940 100644 --- a/projects/roctracer/test/MatrixTranspose_test/MatrixTranspose.cpp +++ b/projects/roctracer/test/MatrixTranspose_test/MatrixTranspose.cpp @@ -25,6 +25,9 @@ THE SOFTWARE. // hip header file #include +// roctx header file +#include + #ifndef ITERATIONS # define ITERATIONS 100 #endif @@ -97,14 +100,23 @@ int main() { // Memory transfer from host to device hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice); + roctxMarkA("before hipLaunchKernel"); + roctxRangePushA("hipLaunchKernel"); + //roctxRangePushA("before hipLaunchKernel"); + // roctxRangePushA("before hipLaunchKernel"); // Lauching kernel from host hipLaunchKernel(matrixTranspose, dim3(WIDTH / THREADS_PER_BLOCK_X, WIDTH / THREADS_PER_BLOCK_Y), dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, gpuTransposeMatrix, gpuMatrix, WIDTH); + //roctxRangePop(); + roctxMarkA("after hipLaunchKernel"); // Memory transfer from device to host + roctxRangePushA("hipMemcpy"); hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost); - + roctxRangePop(); // for "hipMemcpy" + roctxRangePop(); // for "hipLaunchKernel" + // CPU MatrixTranspose computation matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH); @@ -143,6 +155,7 @@ int main() { #if 1 #include #include +#include // Macro to check ROC-tracer calls status #define ROCTRACER_CALL(call) \ @@ -162,6 +175,13 @@ void api_callback( void* arg) { (void)arg; + + if (domain == ACTIVITY_DOMAIN_ROCTX) { + const roctx_api_data_t* data = reinterpret_cast(callback_data); + fprintf(stdout, "ROCTX: \"%s\"\n", data->args.message); + return; + } + const hip_api_data_t* data = reinterpret_cast(callback_data); fprintf(stdout, "<%s id(%u)\tcorrelation_id(%lu) %s> ", roctracer_op_string(ACTIVITY_DOMAIN_HIP_API, cid, 0), From 0bb02f32dee193e0e9f5c03d21f72cb763873226 Mon Sep 17 00:00:00 2001 From: rkebichi <54912798+rkebichi@users.noreply.github.com> Date: Thu, 5 Sep 2019 10:35:28 -0400 Subject: [PATCH 3/4] Update roctx.cpp [ROCm/roctracer commit: 6e1fbda1c772dc4086973786efbf78fa74976ec0] --- projects/roctracer/src/roctx/roctx.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/projects/roctracer/src/roctx/roctx.cpp b/projects/roctracer/src/roctx/roctx.cpp index b7faab91e3..78af443ffe 100644 --- a/projects/roctracer/src/roctx/roctx.cpp +++ b/projects/roctracer/src/roctx/roctx.cpp @@ -62,7 +62,7 @@ THE SOFTWARE. (void)err; \ return X; -std::stack message_stack; +static thread_local std::stack message_stack; static inline uint32_t GetPid() { return syscall(__NR_getpid); } static inline uint32_t GetTid() { return syscall(__NR_gettid); } @@ -117,7 +117,6 @@ PUBLIC_API void roctxMarkA(const char* message) { PUBLIC_API int roctxRangePushA(const char* message) { API_METHOD_PREFIX - //EXC_ABORT(ROCTX_STATUS_ERROR, "method is not implemented"); roctx_api_data_t api_data{}; api_data.args.roctxRangePushA.message = strdup(message); activity_rtapi_callback_t api_callback_fun = NULL; @@ -125,12 +124,12 @@ PUBLIC_API int roctxRangePushA(const char* message) { roctx::cb_table.get(ROCTX_API_ID_roctxRangePushA, &api_callback_fun, &api_callback_arg); if (api_callback_fun) api_callback_fun(ACTIVITY_DOMAIN_ROCTX, ROCTX_API_ID_roctxRangePushA, &api_data, api_callback_arg); message_stack.push(strdup(message)); - API_METHOD_SUFFIX + API_METHOD_CATCH(-1); + return 0; } PUBLIC_API int roctxRangePop() { API_METHOD_PREFIX - //EXC_ABORT(ROCTX_STATUS_ERROR, "method is not implemented"); roctx_api_data_t api_data{}; activity_rtapi_callback_t api_callback_fun = NULL; void* api_callback_arg = NULL; @@ -141,7 +140,8 @@ PUBLIC_API int roctxRangePop() { } else { message_stack.pop(); } - API_METHOD_SUFFIX + API_METHOD_CATCH(-1) + return 0; } } // extern "C" From 548e9464b88772e1307a483adc0bccdb65f75312 Mon Sep 17 00:00:00 2001 From: rkebichi <54912798+rkebichi@users.noreply.github.com> Date: Thu, 5 Sep 2019 10:36:55 -0400 Subject: [PATCH 4/4] Update MatrixTranspose.cpp [ROCm/roctracer commit: 7a46f29e964e0361f50883ec3246f88687608075] --- .../roctracer/test/MatrixTranspose_test/MatrixTranspose.cpp | 3 --- 1 file changed, 3 deletions(-) diff --git a/projects/roctracer/test/MatrixTranspose_test/MatrixTranspose.cpp b/projects/roctracer/test/MatrixTranspose_test/MatrixTranspose.cpp index b519801940..e76802f85d 100644 --- a/projects/roctracer/test/MatrixTranspose_test/MatrixTranspose.cpp +++ b/projects/roctracer/test/MatrixTranspose_test/MatrixTranspose.cpp @@ -102,13 +102,10 @@ int main() { roctxMarkA("before hipLaunchKernel"); roctxRangePushA("hipLaunchKernel"); - //roctxRangePushA("before hipLaunchKernel"); - // roctxRangePushA("before hipLaunchKernel"); // Lauching kernel from host hipLaunchKernel(matrixTranspose, dim3(WIDTH / THREADS_PER_BLOCK_X, WIDTH / THREADS_PER_BLOCK_Y), dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, gpuTransposeMatrix, gpuMatrix, WIDTH); - //roctxRangePop(); roctxMarkA("after hipLaunchKernel"); // Memory transfer from device to host