roctracer extension APIs: mark, external id push/pop

test/MatrixTranspose_test/MatrixTranspose.cpp

@@ -22,6 +22,9 @@ THE SOFTWARE.
 
 #include <iostream>
 
+// roctracer extension API
+#include <inc/roctracer_ext.h>
+
 // hip header file
 #include <hip/hip_runtime.h>
 
@@ -94,17 +97,31 @@ int main() {
         hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
         hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
     
+        // correlation reagion32
+        roctracer_activity_push_external_correlation_id(31);
+        // correlation reagion32
+        roctracer_activity_push_external_correlation_id(32);
+
         // Memory transfer from host to device
         hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice);
     
+        // correlation reagion33
+        roctracer_activity_push_external_correlation_id(33);
+
         // 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);
     
+        // correlation reagion end
+        roctracer_activity_pop_external_correlation_id(NULL);
+
         // Memory transfer from device to host
         hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost);
     
+        // correlation reagion end
+        roctracer_activity_pop_external_correlation_id();
+
         // CPU MatrixTranspose computation
         matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH);
     
@@ -126,6 +143,11 @@ int main() {
         hipFree(gpuMatrix);
         hipFree(gpuTransposeMatrix);
     
+        // correlation reagion end
+        roctracer_activity_pop_external_correlation_id();
+        // correlation reagion end
+        roctracer_activity_pop_external_correlation_id();
+
         // free the resources on host side
         free(Matrix);
         free(TransposeMatrix);
@@ -231,11 +253,15 @@ void activity_callback(const char* begin, const char* end, void* arg) {
         record->device_id,
         record->queue_id
       );
+      if (record->op == hc::HSA_OP_ID_COPY) fprintf(stdout, " bytes(0x%zx)", record->bytes);
+    } else if (record->domain == ACTIVITY_DOMAIN_EXT_API) {
+      fprintf(stdout, " external_id(%lu)",
+        record->external_id
+      );
     } else {
       fprintf(stderr, "Bad domain %d\n", record->domain);
       abort();
     }
-    if (record->op == hc::HSA_OP_ID_COPY) fprintf(stdout, " bytes(0x%zx)", record->bytes);
     fprintf(stdout, "\n");
     fflush(stdout);
     ROCTRACER_CALL(roctracer_next_record(record, &record));

test/tool/tracer_tool.cpp

@@ -240,7 +240,7 @@ void hip_api_flush_cb(hip_api_trace_entry_t* entry) {
   const timestamp_t end_timestamp = entry->end;
   std::ostringstream oss;                                                                        \
 
-  const char* str = (domain < ACTIVITY_DOMAIN_NUMBER) ? roctracer_op_string(domain, cid, 0) : strdup("MARK");
+  const char* str = (domain == ACTIVITY_DOMAIN_EXT_API) ? roctracer_op_string(domain, cid, 0) : strdup("MARK");
   oss << std::dec <<
     begin_timestamp << ":" << end_timestamp << " " << entry->pid << ":" << entry->tid << " " << str;