[GitHub] [incubator-mxnet] DickJC123 commented on a change in pull request #18622: [WIP] Use RTC for elementwise and broadcast ops

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

DickJC123 commented on a change in pull request #18622:
URL: https://github.com/apache/incubator-mxnet/pull/18622#discussion_r457754438



##########
File path: src/common/cuda/rtc.cc
##########
@@ -0,0 +1,246 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements.  See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership.  The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License.  You may obtain a copy of the License at
+ *
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied.  See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include "mxnet/base.h"
+
+#if MXNET_USE_CUDA
+
+#include <nvrtc.h>
+
+#include <mutex>
+#include <string>
+#include <fstream>
+#include <unordered_map>
+#include <vector>
+
+#include "rtc.h"
+#include "rtc/half-inl.h"
+#include "rtc/util-inl.h"
+#include "rtc/forward_functions-inl.h"
+#include "rtc/backward_functions-inl.h"
+#include "rtc/vectorization-inl.h"
+#include "rtc/special_functions-inl.h"
+#include "rtc/reducer-inl.h"
+#include "utils.h"
+
+
+namespace mxnet {
+namespace common {
+namespace cuda {
+namespace rtc {
+
+std::mutex lock;
+
+namespace util {
+
+std::string to_string(OpReqType req) {
+  switch (req) {
+    case kNullOp:
+      return "OpReqType::kNullOp";
+    case kWriteTo:
+    case kWriteInplace:
+      return "OpReqType::kWriteTo";
+    case kAddTo:
+      return "OpReqType::kAddTo";
+  }
+  LOG(FATAL) << "Unrecognized req.";
+  return "";
+}
+
+}  // namespace util
+
+namespace {
+
+// Obtain compilation log from the program.
+std::string GetCompileLog(nvrtcProgram program) {
+  size_t log_size_including_null;
+  NVRTC_CALL(nvrtcGetProgramLogSize(program, &log_size_including_null));
+  // For most std::string implementations, this is probably 1 char bigger than needed.  OK though.

Review comment:
       I wrote the code awhile back, but I think I see the logic behind it.  Forgive me, but I prefer to think of it by way of a concrete example:
   
   Let's say the compile log had 100 characters. The nvrtcGetProgramLogSize() call would report a size of 101 and the nvrtcGetProgramLog() call would copy 101 characters to the pointer address passed to it.  The 101 characters are the 100 characters of the log plus a null byte.  This approach is targeting a C interface- you're expected to malloc an array of 101 bytes (the length reported to you) and all is well.
   
   Here though, we want the program log ultimately to look like std::string of 100 characters.  However, do we want to allocate the string of size 100, then hope we can copy 101 bytes into it?  I think most implementations of std::string will add 1 to the size you give at string construction time, because they want an efficient implementation of string::c_str().  So it would probably work, but why risk it?  Basically, we allocate the string with length 101, let nvrtcGetProgramLog() copy the 101 bytes into it, then resize the string to length 100 so we don't have the null character messing up any print or other string manipulations we do down the road.




----------------------------------------------------------------
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