diff --git a/projects/roctracer/src/roctx/roctx.cpp b/projects/roctracer/src/roctx/roctx.cpp index 1799b17910..78af443ffe 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; +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); } @@ -114,14 +117,31 @@ 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"); - API_METHOD_SUFFIX + 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_CATCH(-1); + return 0; } PUBLIC_API int roctxRangePop() { API_METHOD_PREFIX - EXC_ABORT(ROCTX_STATUS_ERROR, "method is not implemented"); - API_METHOD_SUFFIX + 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_CATCH(-1) + return 0; } } // extern "C" diff --git a/projects/roctracer/test/MatrixTranspose_test/MatrixTranspose.cpp b/projects/roctracer/test/MatrixTranspose_test/MatrixTranspose.cpp index 154c2ade3d..79a3845588 100644 --- a/projects/roctracer/test/MatrixTranspose_test/MatrixTranspose.cpp +++ b/projects/roctracer/test/MatrixTranspose_test/MatrixTranspose.cpp @@ -28,6 +28,9 @@ THE SOFTWARE. // hip header file #include +// roctx header file +#include + #ifndef ITERATIONS # define ITERATIONS 100 #endif @@ -104,20 +107,30 @@ int main() { // Memory transfer from host to device hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice); - + // correlation reagion33 roctracer_activity_push_external_correlation_id(33); + roctxMarkA("before hipLaunchKernel"); + roctxRangePushA("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); - + + roctxMarkA("after hipLaunchKernel"); + // correlation reagion end roctracer_activity_pop_external_correlation_id(NULL); // Memory transfer from device to host + roctxRangePushA("hipMemcpy"); + hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost); + + roctxRangePop(); // for "hipMemcpy" + roctxRangePop(); // for "hipLaunchKernel" // correlation reagion end roctracer_activity_pop_external_correlation_id(); @@ -165,6 +178,7 @@ int main() { #if 1 #include #include +#include // Macro to check ROC-tracer calls status #define ROCTRACER_CALL(call) \ @@ -184,6 +198,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),