Merge branch 'main' into fix-include

Author: Conless

libkineto/sample_programs/build.sh

@@ -18,5 +18,7 @@ g++ \
   -lpthread \
   -lcuda \
   -lcudart \
+  -lcupti \
+  -lnvperf_host \
   /usr/local/lib/libkineto.a \
   kplay_cu.o

libkineto/sample_programs/kineto_playground.cpp

@@ -6,6 +6,7 @@
  * LICENSE file in the root directory of this source tree.
  */
 
+#include <cuda_runtime.h>
 #include <iostream>
 #include <string>
 
@@ -14,6 +15,20 @@
 // @lint-ignore-every CLANGTIDY facebook-hte-RelativeInclude
 #include "kineto_playground.cuh"
 
+#define CHECK_CUDA(call)               \
+  do {                                 \
+    cudaError_t status = call;         \
+    if (status != cudaSuccess) {       \
+      fprintf(                         \
+          stderr,                      \
+          "CUDA Error at %s:%d: %s\n", \
+          __FILE__,                    \
+          __LINE__,                    \
+          cudaGetErrorString(status)); \
+      exit(1);                         \
+    }                                  \
+  } while (0)
+
 using namespace kineto;
 
 static const std::string kFileName = "/tmp/kineto_playground_trace.json";
@@ -23,6 +38,7 @@ int main() {
   warmup();
 
   // Kineto config
+  libkineto_init(false, true);
 
   // Empty types set defaults to all types
   std::set<libkineto::ActivityType> types;
@@ -38,6 +54,7 @@ int main() {
   profiler.startTrace();
   std::cout << "Start playground" << std::endl;
   playground();
+  CHECK_CUDA(cudaDeviceSynchronize());
 
   std::cout << "Stop Trace" << std::endl;
   auto trace = profiler.stopTrace();

libkineto/sample_programs/kineto_playground.cu

@@ -76,7 +76,7 @@ __global__ void square(float* A, int N) {
 
 void playground(void) {
   // Add your experimental CUDA implementation here.
-  basicMemcpyFromDevice();
+  basicMemcpyToDevice();
   compute();
   basicMemcpyFromDevice();
 }

libkineto/src/CuptiActivityProfiler.cpp

@@ -459,12 +459,9 @@ void CuptiActivityProfiler::processCpuTrace(
               const std::unique_ptr<GenericTraceActivity>>::value,
           "handleActivity is unsafe and relies on the caller to maintain not "
           "only lifetime but also address stability.");
-      if (act->type() == ActivityType::USER_ANNOTATION &&
-          act->duration() <= 0) {
+      if (act->duration() <= 0) {
         act->endTime = captureWindowEndTime_;
         act->addMetadata("finished", "false");
-      } else {
-        act->addMetadata("finished", "true");
       }
       logger.handleActivity(*act);
     }
@@ -581,8 +578,8 @@ inline static bool isBlockListedRuntimeCbid(CUpti_CallbackId cbid) {
   if (cbid == CUPTI_RUNTIME_TRACE_CBID_cudaGetDevice_v3020 ||
       cbid == CUPTI_RUNTIME_TRACE_CBID_cudaSetDevice_v3020 ||
       cbid == CUPTI_RUNTIME_TRACE_CBID_cudaGetLastError_v3020 ||
-      // Support cudaEventRecord and cudaEventSynchronize, revisit if others are
-      // needed
+      // Support cudaEventRecord and cudaEventSynchronize, revisit if others
+      // are needed
       cbid == CUPTI_RUNTIME_TRACE_CBID_cudaEventCreate_v3020 ||
       cbid == CUPTI_RUNTIME_TRACE_CBID_cudaEventCreateWithFlags_v3020 ||
       cbid == CUPTI_RUNTIME_TRACE_CBID_cudaEventDestroy_v3020) {
@@ -1339,20 +1336,21 @@ const time_point<system_clock> CuptiActivityProfiler::performRunLoopStep(
           || cupti_.stopCollection
 #endif // HAS_CUPTI || HAS_ROCTRACER
       ) {
-        // Update runloop state first to prevent further updates to shared state
+        // Update runloop state first to prevent further updates to shared
+        // state
         LOG(INFO) << "Tracing complete.";
         VLOG_IF(1, currentIter >= 0)
             << "This state change was invoked by application's step() call";
 
         // currentIter >= 0 means this is called from the step() api of
-        // the profile in pytorch main thread, it should be executed in another
-        // thread in case pytorch main thread is blocked
+        // the profile in pytorch main thread, it should be executed in
+        // another thread in case pytorch main thread is blocked
         if (currentIter >= 0) {
           // if collectTraceThread_ is already running, there's no need to
           // execute collectTrace twice.
-          // Do not call collectTrace when profilerThread_ is collecting Trace.
-          // Otherwise, libkineto::api().client()->stop will be called twice,
-          // which leads to an unrecoverable ::c10:Error at
+          // Do not call collectTrace when profilerThread_ is collecting
+          // Trace. Otherwise, libkineto::api().client()->stop will be called
+          // twice, which leads to an unrecoverable ::c10:Error at
           // disableProfiler
           if (!collectTraceThread_ && !getCollectTraceState()) {
             std::lock_guard<std::recursive_mutex> guard(mutex_);

libkineto/src/output_json.cpp

@@ -131,6 +131,13 @@ void ChromeTraceLogger::handleTraceStart(
       device_properties);
 
   metadataToJSON(metadata);
+
+  traceOf_ << fmt::format(
+      R"JSON(
+  "displayTimeUnit": "ms",
+  "baseTimeNanoseconds": {},)JSON",
+      ChromeTraceBaseTime::singleton().get());
+
   traceOf_ << R"JSON(
   "traceEvents": [)JSON";
 }
@@ -681,13 +688,9 @@ void ChromeTraceLogger::finalizeTrace(
 #endif // !USE_GOOGLE_LOG
 
   // Putting this here because the last entry MUST not end with a comma.
-
-  traceOf_ << fmt::format(R"JSON(
-  "traceName": "{}",
-  "displayTimeUnit": "ms",
-  "baseTimeNanoseconds": {}
-}})JSON", fileName_, ChromeTraceBaseTime::singleton().get());
-  // clang-format on
+    traceOf_ << fmt::format(R"JSON(
+  "traceName": "{}"
+}})JSON", fileName_);
 
   traceOf_.close();
   // On some systems, rename() fails if the destination file exists.