Add a dump_ptx flag for deeper bug analysis

This commit was originally used for tracking an illegal memory access at the
PTX level. It is useful independently.

Author: nicolasvasilache

tc/core/cuda/cuda_rtc.cc

@@ -113,7 +113,9 @@ std::unique_ptr<CudaRTCFunction> CudaRTCFunction::Compile(
   res->nvrtc_ptx = std::vector<char>(ptx_size);
   TC_NVRTC_CHECK(nvrtcGetPTX(prog, res->nvrtc_ptx.data()));
   TC_NVRTC_CHECK(nvrtcDestroyProgram(&prog));
-
+  if (FLAGS_dump_ptx) {
+    LOG(INFO) << "PTX:\n" << std::string(res->nvrtc_ptx.data());
+  }
   return res;
 }
 namespace {

tc/core/flags.cc

@@ -36,6 +36,7 @@ DEFINE_bool(
     false,
     "Print debug spew for the tc_mapper like cuda code, mapping options etc");
 DEFINE_bool(dump_cuda, false, "Print the generated source");
+DEFINE_bool(dump_ptx, false, "Dump the generated PTX");
 
 // CPU codegen options
 DEFINE_bool(llvm_dump_before_opt, false, "Print IR before optimization");

tc/core/flags.h

@@ -29,6 +29,7 @@ DECLARE_bool(debug_tc_mapper);
 DECLARE_bool(debug_cuda);
 DECLARE_bool(debug_tuner);
 DECLARE_bool(dump_cuda);
+DECLARE_bool(dump_ptx);
 
 // llvm codegen
 DECLARE_bool(llvm_dump_before_opt);