Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

[Metal] Add codegen/runtime support for print() #1310

Merged
merged 5 commits into from
Jun 26, 2020
Merged

[Metal] Add codegen/runtime support for print() #1310

merged 5 commits into from
Jun 26, 2020

Conversation

k-ye
Copy link
Member

@k-ye k-ye commented Jun 23, 2020

This PR completes the support for print() on Metal. Again, it's similar to OpenGL's implementation. I added a new Metal buffer to hold PrintMsg. We will need a sync when print is used.

Related issue = #1281

[Click here for the format server]

k-ye
k-ye commented Jun 23, 2020
View reviewed changes
tests/python/test_print.py
def test_print_matrix():
x = ti.Matrix(2, 3, dt=ti.f32, shape=())
y = ti.Vector(3, dt=ti.f32, shape=3)

@ti.kernel
def func(k: ti.f32):
x[None][0, 0] = -1.0
Copy link
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Without this, x and y are just all zero.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Btw, we may able to assert the print result after #1303?

Copy link
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yep, good point. I think we can compare pop_python_print_buffer with expected strs.

@codecov
Copy link

codecov bot commented Jun 23, 2020 •
edited
Loading

Codecov Report

❗ No coverage uploaded for pull request base (master@37976e7). Click here to learn what that means.
The diff coverage is n/a.

Impacted file tree graph

@@            Coverage Diff            @@
##             master    #1310   +/-   ##
=========================================
  Coverage          ?   66.40%           
=========================================
  Files             ?       37           
  Lines             ?     5155           
  Branches          ?      939           
=========================================
  Hits              ?     3423           
  Misses            ?     1568           
  Partials          ?      164

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update 37976e7...7651019. Read the comment docs.

archibate
archibate reviewed Jun 24, 2020
View reviewed changes
Copy link
Collaborator

@archibate archibate left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks, didn't read all yet.

taichi/backends/metal/codegen_metal.cpp
@@ -1043,6 +1086,11 @@ class KernelCodegen : public IRVisitor {
return kernel_name + "_func";
}

void mark_print_used() {
TI_ASSERT(current_kernel_attribs_ != nullptr);
current_kernel_attribs_->uses_print = true;
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

nit: Use a UsedFeatures structure like OpenGL does.

Copy link
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Good point.. This is a sloppy design in Metal. I don't have a data structure that records the attributes of a taichi kernel. All i have is a vector of Metal kernels', see

taichi/taichi/backends/metal/kernel_manager.h

Line 45 in 37976e7

const std::vector<KernelAttributes> &kernels_attribs,
.

I'd like to take this opportunity to clean up this flaw in another PR. Once I have a TaichiKernelAttributes, i can also add a UsedFeatures there. Then we don't even need to iterate over all the metal kernels to figure out if we need to flush the print buffers or not. (I tried adding something, but feel like it's too noisy for this one.) WDYT?

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes, ff2 do it iapr, you may also implement a operator|=(UsedFeatures, UsedFeatures) for recording all used feature among kernels :)

tests/python/test_print.py
def test_print_matrix():
x = ti.Matrix(2, 3, dt=ti.f32, shape=())
y = ti.Vector(3, dt=ti.f32, shape=3)

@ti.kernel
def func(k: ti.f32):
x[None][0, 0] = -1.0
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Btw, we may able to assert the print result after #1303?

@k-ye k-ye requested a review from archibate June 25, 2020 11:04
@k-ye
Copy link
Member Author

k-ye commented Jun 25, 2020 •
edited
Loading

Note that code format presubmit failed, which is intentional. I guess the clang formatter is never gonna play well with these .metal.h files. The suggested code style added nested indentation for each kernel, and looks terrible 😞

+        [[maybe_unused]] inline int mtl_compute_num_print_msg_typemasks(
+            int num_entries) {
+          return (num_entries + kMetalNumPrintMsgTypePerI32 - 1) /
+                 kMetalNumPrintMsgTypePerI32;
+        }
+
+            [[maybe_unused]] inline int mtl_compute_print_msg_bytes(
+                int num_entries) {
+              // See PrintMsg's layout for how this is computed.
+              const int sz =
+                  sizeof(int32_t) *
+                  (1 + mtl_compute_num_print_msg_typemasks(num_entries) +
+                   num_entries);
+              return sz;
+            }

yuanming-hu
yuanming-hu approved these changes Jun 26, 2020
View reviewed changes
Copy link
Member

@yuanming-hu yuanming-hu left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thank you so much! This will help a lot of Metal users to debug their programs. The changes LGTM.

TH3CHARLie
TH3CHARLie approved these changes Jun 26, 2020
View reviewed changes
Copy link
Collaborator

@TH3CHARLie TH3CHARLie left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM! Verified locally with LLVM 10.0.0, Xcode 11.5

@k-ye
Copy link
Member Author

k-ye commented Jun 26, 2020

Merging so that i can further address #1310 (comment)

@k-ye k-ye merged commit 477b1e1 into taichi-dev:master Jun 26, 2020
@k-ye k-ye deleted the print2 branch June 26, 2020 10:47
@k-ye k-ye mentioned this pull request Jun 26, 2020
Merged
@yuanming-hu yuanming-hu mentioned this pull request Jun 28, 2020
Merged
@archibate archibate mentioned this pull request Jul 2, 2020
Closed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers

@TH3CHARLie TH3CHARLie TH3CHARLie approved these changes

@yuanming-hu yuanming-hu yuanming-hu approved these changes

@taichi-gardener taichi-gardener Awaiting requested review from taichi-gardener

@archibate archibate Awaiting requested review from archibate

Assignees
No one assigned
Labels
None yet
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.
4 participants
@k-ye @yuanming-hu @archibate @TH3CHARLie