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

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

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

-            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

        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


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


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

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

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

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

+    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

-        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

 #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

+        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

-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

+    }
+};
+
+


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

+            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

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


extra return statement

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

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


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

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

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

+    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

+    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

+    // 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)

Uh oh!

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