Convert CUDA JIT to use nvrtc instead of nvvm by pavanky · Pull Request #1836 · arrayfire/arrayfire

pavanky · 2017-06-19T03:51:14Z

No description provided.

umar456

Looks great. I added a couple of minor comments.

umar456 · 2017-06-19T23:55:38Z

src/backend/cuda/Array.cpp

    {
-        Node_ptr node(reinterpret_cast<Node *>(new BufferNode<T>(irname<T>(), afShortName<T>())));
-        return node;
+        return Node_ptr(reinterpret_cast<Node *>(new BufferNode<T>(getFullName<T>(),


I am pretty sure you can just do this:

return Node_ptr(new BufferNode<T>(getFullName<T>(), shortname<T>(true)));

umar456 · 2017-06-19T23:56:18Z

src/backend/cuda/CMakeLists.txt

@@ -3,10 +3,6 @@ CMAKE_MINIMUM_REQUIRED(VERSION 2.8)
 FIND_PACKAGE(CUDA 7.0 REQUIRED)

 INCLUDE(CLKernelToH)


We should change the name of this guy at some point.

I am not doing it now :(

umar456 · 2017-06-20T00:05:16Z

src/backend/cuda/JIT/BinaryNode.hpp

-            kerStream << std::setw(2) << std::setfill('0') << std::hex << ids.child_ids[0];
-            kerStream << std::setw(2) << std::setfill('0') << std::hex << ids.child_ids[1];
-            kerStream << std::setw(2) << std::setfill('0') << std::hex << ids.id << std::dec;
+            // Make the dec representation of enum part of the Kernel name


These aren't templated so it would be better if they were implemented in the cpp file.

I'll do another pass about reorganizing the JIT nodes at a later point (including trying to use the same code for CUDA and OpenCL JIT). Can we leave this be for now?

umar456 · 2017-06-20T00:24:37Z

src/backend/cuda/JIT/BufferNode.hpp

        std::shared_ptr<T> m_data;
        Param<T> m_param;
        unsigned m_bytes;
+        bool m_linear_buffer;


Since bool is most likely smaller than once_flag it would be better to place it after the m_set_data_flag variable. It would be good to look at the other variables while you are at it.

umar456 · 2017-06-20T00:45:02Z

src/backend/cuda/JIT/BufferNode.hpp


+        bool isBuffer() { return true; }
+
+        ~BufferNode()


Is the destructor necessary? You aren't really doing anything in here.

umar456 · 2017-06-20T01:06:51Z

src/backend/cuda/jit.cpp

-        "br i1 %cmp0, label %core, label %end\n";
+    const char *includeFileStr = jit_cuh;
+
+    const char paramTStr[] =""


Since we are using C++11 now, raw string literals would be easier to read.

umar456 · 2017-06-20T01:17:58Z

src/backend/cuda/jit.cpp

-    } else {
-        for (int i = 0; i < 4; i++) args.push_back((void *)(strides + i));
-        for (int i = 0; i < 4; i++) args.push_back((void *)(dims + i));
+        args.push_back((void *)&outputs[i]);


Looks like you know the size of the args vector. Could yo call the reserve function?

We don't know the size of the args vector. It depends on what each type of the node is.

umar456 · 2017-06-20T01:19:32Z

src/backend/cuda/jit.cpp

+    NVRTC_CHECK(nvrtcCreateProgram(&prog, jit_ker.c_str(), ker_name, 0, NULL, NULL));
+    NVRTC_CHECK(nvrtcCompileProgram(prog, 0, NULL));
+    NVRTC_CHECK(nvrtcGetPTXSize(prog, ptx_size));
+    unique_ptr<char []> ptx(new char[*ptx_size]);


Can you use vector instead?

umar456 · 2017-06-20T01:21:32Z

src/backend/cuda/jit.cpp

-        kerStream << triple32;
+        outParamStream << "Param<" << full_nodes[id]->getTypeStr() << "> out" << id << ", \n";
+        // Generate code to write the output
+        outWriteStream << "out" << id << ".ptr[idx] = " << "val" << id << ";\n";


outWriteStream << "out" << id << ".ptr[idx] = val" << id << ";\n";

umar456 · 2017-06-20T01:23:02Z

src/backend/cuda/jit.cpp

 #include <stdexcept>
 #include <thread>
 #include <vector>
+#include <kernel_headers/jit.hpp>


Could you separate the standard headers from the others?

umar456 · 2017-06-20T04:01:16Z

Windows builds were failing because of a Windows macro definition. Fix #1839 should resolve this issue. Please rebase and test again.

FloopCZ

That is a lot of great code.

Most of my review comments are of the OCD type, so feel free to ignore if you don't suffer from this terrible disorder. ;-)

FloopCZ · 2017-06-20T06:58:57Z

src/backend/cuda/jit.cpp

+        unique_ptr<char []> log(new char[logSize]); \
+        char *logptr = log.get();                   \
+        nvrtcGetProgramLog(prog, logptr);           \
+        logptr[logSize] = '\x0';                    \


I believe this write is one-past-end.

FloopCZ · 2017-06-20T07:07:45Z

src/backend/cuda/types.hpp

-template<typename T, bool binary> const std::string cuMangledName(const char *fn);
-template<typename T > const char *afShortName(bool caps = true);
-template<typename T > const char *irname();
+template<typename T > const char *shortname(bool caps = true);


OCD - extra space after typename T

FloopCZ · 2017-06-20T07:11:47Z

src/backend/cuda/cast.hpp

+    }
+};
+
+


extra newline here and L65 and L60
(Nitpicking, but since the code looks so good and you already edit the line, it deserves to be perfect... ;-))

FloopCZ · 2017-06-20T07:20:52Z

src/backend/cuda/JIT/BufferNode.hpp

+            if (is_linear) {
+                kerStream << idx_str << " = idx;\n";
+            } else {
+                std::string info_str = std::string("in") + std::to_string(id);;


double semicolon

FloopCZ · 2017-06-20T07:22:38Z

src/backend/cuda/JIT/BufferNode.hpp

@@ -158,21 +123,6 @@ namespace JIT
            bytes += m_bytes;
            return;


extra return statement

FloopCZ · 2017-06-20T07:26:32Z

src/backend/cuda/binary.hpp

+            return "__c"#fn"f";                 \
+        }                                       \
+    };                                          \
+                                                \


OCD - newline in here, but not after any of the BinOps above

FloopCZ · 2017-06-20T07:30:53Z

src/backend/cuda/types.cpp

-#define INSTANTIATE(T)                                                  \
-    template const std::string cuMangledName<T, false>(const char *fn); \
-    template const std::string cuMangledName<T, true>(const char *fn);  \
+    template<typename T > const char *shortname(bool caps) { return caps ?  "Q" : "q"; }


OCD - extra space after typename T (makes it not aligned)

FloopCZ · 2017-06-20T07:33:39Z

src/backend/cuda/kernel/jit.cuh

+    double rhs_x = inv_rhs_abs * rhs.x;
+    double rhs_y = inv_rhs_abs * rhs.y;
+    cdouble out = {lhs.x * rhs_x + lhs.y * rhs_y,
+                  lhs.y * rhs_x - lhs.x * rhs_y};


OCD - prefix with one extra space to be aligned as L106

FloopCZ · 2017-06-20T07:34:33Z

src/backend/opencl/kernel/jit.cl

+    double rhs_x = inv_rhs_abs * rhs.x;
+    double rhs_y = inv_rhs_abs * rhs.y;
+    double2 out = {lhs.x * rhs_x + lhs.y * rhs_y,
+                  lhs.y * rhs_x - lhs.x * rhs_y};


OCD - prefix with one space to make it aligned as L75

FloopCZ · 2017-06-20T07:37:19Z

src/backend/cuda/jit.cpp

+    // Common CUDA code
+    // This part of the code does not change with the kernel.
+
+    static const char *kernelVoid =  "extern \"C\" __global__ void\n";


OCD - extra space after equal sign

pavanky · 2017-06-20T12:19:13Z

@umar456 @FloopCZ Addressed all your feedback.

pavanky force-pushed the nvrtc branch from 4e108c6 to 7a6c49b Compare June 19, 2017 18:21

pavanky changed the title ~~[WIP]: Convert CUDA JIT to use nvrtc instead of nvvm~~ Convert CUDA JIT to use nvrtc instead of nvvm Jun 19, 2017

pavanky added 2 commits June 19, 2017 11:48

Convert CUDA JIT to use nvrtc instead of nvvm

8d6f981

Fixing underflow/overflow issues with complex math in opencl jit

225a3e0

pavanky force-pushed the nvrtc branch from 7a6c49b to 225a3e0 Compare June 19, 2017 19:32

umar456 requested changes Jun 20, 2017

View reviewed changes

FloopCZ reviewed Jun 20, 2017

View reviewed changes

Style changes for JIT in CUDA and OpenCL backends

81d74af

umar456 approved these changes Jun 20, 2017

View reviewed changes

Use std::string instead of char * for storing text of the jit.cuh file

bd5667a

umar456 merged commit b19c02a into arrayfire:devel Jun 20, 2017

pavanky deleted the nvrtc branch June 21, 2017 01:42

pavanky added this to the v3.5.0 milestone Jun 21, 2017

pavanky mentioned this pull request Jun 21, 2017

Release notes for v3.5.0 #1822

Merged

pavanky mentioned this pull request Aug 2, 2017

Use nvrtc for CUDA JIT #1400

Closed

9prady9 mentioned this pull request Dec 8, 2017

Update unified backend doc #2008

Closed

		@@ -3,10 +3,6 @@ CMAKE_MINIMUM_REQUIRED(VERSION 2.8)
		FIND_PACKAGE(CUDA 7.0 REQUIRED)

		INCLUDE(CLKernelToH)

Conversation

pavanky commented Jun 19, 2017 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

umar456 left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

umar456 commented Jun 20, 2017

Uh oh!

FloopCZ left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

pavanky commented Jun 20, 2017

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

pavanky commented Jun 19, 2017 •

edited

Loading