[GitHub] [incubator-mxnet] ptrendx commented on a change in pull request #19364: Make RTC compatible with CUDA enhanced compatibility

[GitBox <gi...@apache.org>]

ptrendx commented on a change in pull request #19364:
URL: https://github.com/apache/incubator-mxnet/pull/19364#discussion_r516236025



##########
File path: src/common/cuda/rtc.cc
##########
@@ -78,15 +80,47 @@ std::string GetCompileLog(nvrtcProgram program) {
 }
 
 // Obtain compilation result (ptx assembly) from the program.
-std::string GetPtx(nvrtcProgram program) {
+std::string GetCompiledCode(nvrtcProgram program, bool use_cubin) {
+#if CUDA_VERSION >= 11010
+  const auto getSize = use_cubin ? nvrtcGetCUBINSize : nvrtcGetPTXSize;
+  const auto getFunc = use_cubin ? nvrtcGetCUBIN : nvrtcGetPTX;
+#else
+  const auto getSize = nvrtcGetPTXSize;
+  const auto getFunc = nvrtcGetPTX;
+#endif
   size_t ptx_size_including_null;
-  NVRTC_CALL(nvrtcGetPTXSize(program, &ptx_size_including_null));
+  NVRTC_CALL(getSize(program, &ptx_size_including_null));
   std::string ptx(ptx_size_including_null - 1, '\0');
   // Room for terminating null character ensured since C++11
-  NVRTC_CALL(nvrtcGetPTX(program, &ptx[0]));
+  NVRTC_CALL(getFunc(program, &ptx[0]));
   return ptx;
 }
 
+std::tuple<bool, std::string> GetArchString(const int sm_arch) {
+#if CUDA_VERSION < 10000
+  constexpr int max_supported_sm_arch = 70;

Review comment:
       :+1:

##########
File path: src/common/rtc.cc
##########
@@ -55,10 +55,30 @@ CudaModule::Chunk::Chunk(
     LOG(FATAL) << err.data();
   }
 
-  size_t ptx_size;
-  NVRTC_CALL(nvrtcGetPTXSize(prog_, &ptx_size));
-  ptx_ = new char[ptx_size];
-  NVRTC_CALL(nvrtcGetPTX(prog_, ptx_));
+  bool use_ptx = true;
+  for (const auto& opt : options) {
+    if (opt.find("sm_") != std::string::npos) {
+      use_ptx = false;
+      break;
+    }
+  }
+
+  if (use_ptx) {
+    size_t ptx_size;
+    NVRTC_CALL(nvrtcGetPTXSize(prog_, &ptx_size));
+    ptx_.resize(ptx_size);
+    NVRTC_CALL(nvrtcGetPTX(prog_, ptx_.data()));
+  } else {
+#if CUDA_VERSION >= 11010
+    size_t cubin_size;
+    NVRTC_CALL(nvrtcGetCUBINSize(prog_, &cubin_size));
+    ptx_.resize(ptx_size);

Review comment:
       :+1:




----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.

For queries about this service, please contact Infrastructure at:
users@infra.apache.org