[PASS] Add GPU IR verifier #1296

merrymercy · 2018-06-17T20:21:57Z

Add a pass to check whether a cuda ir is valid

merrymercy · 2018-06-17T21:14:58Z

/workspace/src/arithmetic/modular.cc:168:1: fatal error: error writing to /tmp/cc30MweO.s: No space left on device

tqchen · 2018-06-18T02:49:07Z

include/tvm/ir_pass.h

+ * \return valid Whether it is a valid cuda ir
+ *
+ */
+bool VerifyCuda(Stmt stmt,


This is CUDA specific? or should we call it VerifyGPUCode

tqchen · 2018-06-20T17:10:27Z

@merrymercy please act on the comments and fix the ci error

tqchen · 2018-06-20T20:07:59Z

src/pass/verify_gpu_code.cc

+
+class GPUCodeVerifier : public IRVisitor {
+ public:
+  bool verify(tvm::Stmt stmt, int max_shared_memory_per_block, int max_thread_per_block) {


CamelCase for fucnctions

tqchen · 2018-06-20T20:09:08Z

src/pass/verify_gpu_code.cc

+    if (shared_buffers_.count(op->buffer_var.get()) != 0) {
+      int64_t size = op->type.bytes();
+      for (auto dim : op->extents) {
+        size *= dim.as<IntImm>()->value;


op-> constant_allocation_size()

tqchen · 2018-06-20T20:09:55Z

src/pass/verify_gpu_code.cc

+      // record the number of threads in a block
+      std::string name = var.get()->name_hint;
+      if (name == "threadIdx.x" || name == "threadIdx.y" || name == "threadIdx.z") {
+        if (visited_threads_.find(name) == visited_threads_.end()) {


!count(name)

tqchen · 2018-06-20T20:10:10Z

src/pass/verify_gpu_code.cc

+  size_t max_shared_memory_per_block_;
+  size_t max_thread_per_block_;
+
+  bool valid{true};


tqchen · 2018-06-20T20:10:17Z

src/pass/verify_gpu_code.cc

+
+  bool valid{true};
+
+  void reset_() {


If this visitor is only used once, reset is not necessary

reset is needed because there might be several gpu kernels in one Stmt.

tqchen · 2018-06-20T20:10:48Z

src/pass/verify_gpu_code.cc

+    }
+
+    if (op->is_producer) {
+      nest_level_++;


prefer --nest_+level

tqchen · 2018-06-20T20:11:15Z

@eqy can you also do a round of codereview?

eqy · 2018-06-20T20:42:33Z

@tqchen @merrymercy if we also save the number of threads per dimension (x, y, z), perhaps we can also use this to capture
CL_INVALID_WORK_ITEM_SIZE (e.g., in https://www.khronos.org/registry/OpenCL/sdk/2.0/docs/man/xhtml/clEnqueueNDRangeKernel.html)
for OpenCL kernels. But I am not sure if this can happen in CUDA land, so we could also use a separate verifier pass there.

EDIT: It does seem that CUDA devices can have a similar limit, which can read with deviceQuery
e.g., Max dimension size of a thread block (x,y,z): (1024, 1024, 64)

tqchen · 2018-06-20T21:14:40Z

It would be a good idea to pass in as many constraints as possible and allowing defaults to non-constraints. One possible way to do so is to allow pass in of a Map<str, value> so we don't have a lot of positional arguments

antinucleon · 2018-06-21T17:14:48Z

src/pass/verify_gpu_code.cc

+
+  std::unordered_set<const tvm::Variable *> shared_buffers_;
+  std::unordered_set<std::string> visited_threads_;
+  size_t shared_memory_per_block_;


local_memiry_per_block_ is also needed.

merrymercy · 2018-06-22T07:15:07Z

ready for review

eqy · 2018-06-22T08:11:26Z

Is the plan to skip checking threadId/workitem dimensions in this round?

tqchen · 2018-06-22T20:51:20Z

please fix the compiler warning http://mode-gpu.cs.washington.edu:8080/blue/organizations/jenkins/dmlc%2Ftvm/detail/PR-1296/11/pipeline

Currently we set compiler warning as error so built won't pass if there is a warning

merrymercy · 2018-06-23T14:32:57Z

@eqy your comment is addressed
@tqchen no space left on CI

tqchen · 2018-06-23T17:17:50Z

tests/python/unittest/test_pass_verify_gpu_code.py

+"""Test gpu code verifier"""
+import tvm
+
+global valid


always avoid use of global variable, to carry state, you can use a closure to capture a list

tqchen · 2018-06-23T17:31:17Z

The test error likely indicate there is some problem with the current PR when importing runtime only dll.

add cuda verifier

459e35b

merrymercy changed the title ~~[PASS] Add cuda verifier~~ [PASS] Add CUDA IR verifier Jun 17, 2018

fix lint

9d24748

tqchen requested changes Jun 18, 2018

View reviewed changes

tqchen added the status: need update need update based on feedbacks label Jun 18, 2018

merrymercy added 2 commits June 21, 2018 03:19

rename: cuda ir -> gpu code

a573596

fix test

630cc39

tqchen requested changes Jun 20, 2018

View reviewed changes

antinucleon reviewed Jun 21, 2018

View reviewed changes

use Map<str, value> as argument & add local memory usage check

a3aa9db

merrymercy changed the title ~~[PASS] Add CUDA IR verifier~~ [PASS] Add GPU IR verifier Jun 21, 2018

merrymercy added 2 commits June 22, 2018 18:51

check length per dimension

7546c20

add new device query

576ab67

fix warning

6420092

tqchen requested changes Jun 23, 2018

View reviewed changes

fix runtime load

f50f040

merrymercy force-pushed the cuda-verifier branch from aebd3fd to f50f040 Compare June 23, 2018 19:09

remove global variable

7ff9f3d

tqchen approved these changes Jun 23, 2018

View reviewed changes

tqchen merged commit e66996a into apache:master Jun 23, 2018

tqchen added status: accepted and removed status: need update need update based on feedbacks labels Jun 23, 2018

tqchen pushed a commit to tqchen/tvm that referenced this pull request Jul 6, 2018

[PASS] Add GPU IR verifier (apache#1296)

5cdfd5b

merrymercy deleted the cuda-verifier branch July 10, 2018 20:08

mnuyens pushed a commit to mnuyens/tvm that referenced this pull request Jul 10, 2018

[PASS] Add GPU IR verifier (apache#1296)

531bb7c

sergei-mironov pushed a commit to sergei-mironov/tvm that referenced this pull request Aug 8, 2018

[PASS] Add GPU IR verifier (apache#1296)

7db030b

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[PASS] Add GPU IR verifier #1296

[PASS] Add GPU IR verifier #1296

merrymercy commented Jun 17, 2018

merrymercy commented Jun 17, 2018

tqchen Jun 18, 2018

tqchen commented Jun 20, 2018

tqchen Jun 20, 2018

tqchen Jun 20, 2018

tqchen Jun 20, 2018

tqchen Jun 20, 2018

tqchen Jun 20, 2018

tqchen Jun 20, 2018

merrymercy Jun 21, 2018 •

edited

Loading

tqchen Jun 20, 2018

tqchen commented Jun 20, 2018

eqy commented Jun 20, 2018 •

edited

Loading

tqchen commented Jun 20, 2018

antinucleon Jun 21, 2018

merrymercy commented Jun 22, 2018

eqy commented Jun 22, 2018

tqchen commented Jun 22, 2018

merrymercy commented Jun 23, 2018 •

edited

Loading

tqchen Jun 23, 2018

tqchen commented Jun 23, 2018

[PASS] Add GPU IR verifier #1296

[PASS] Add GPU IR verifier #1296

Conversation

merrymercy commented Jun 17, 2018

merrymercy commented Jun 17, 2018

Choose a reason for hiding this comment

tqchen commented Jun 20, 2018

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

merrymercy Jun 21, 2018 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

tqchen commented Jun 20, 2018

eqy commented Jun 20, 2018 • edited Loading

tqchen commented Jun 20, 2018

Choose a reason for hiding this comment

merrymercy commented Jun 22, 2018

eqy commented Jun 22, 2018

tqchen commented Jun 22, 2018

merrymercy commented Jun 23, 2018 • edited Loading

Choose a reason for hiding this comment

tqchen commented Jun 23, 2018

merrymercy Jun 21, 2018 •

edited

Loading

eqy commented Jun 20, 2018 •

edited

Loading

merrymercy commented Jun 23, 2018 •

edited

Loading