Bootstrap metal runtime #263

nsmithtt · 2024-07-31T21:15:35Z

No description provided.

kmabeeTT · 2024-08-01T18:00:05Z

Thanks, Nick for putting this together! I am taking a look at this...

kmabeeTT

LGTM, mostly questions for my own benefit. The most interesting stuff was in new files runtime.cpp and command_queue.cpp, but I guess we don't support some important stuff like EnqueueReadBuffer/WriteBuffer yet.

kmabeeTT · 2024-08-01T18:34:30Z

test/ttmlir/Silicon/TTNN/simple_ge.mlir

@@ -1,4 +1,4 @@
-// RUN: ttmlir-opt --ttir-layout --ttnn-open-device --convert-ttir-to-ttnn %s  > %t.mlir
+// RUN: ttmlir-opt --ttir-implicit-device --ttir-layout --ttnn-open-device --convert-ttir-to-ttnn %s  > %t.mlir


Just curious, what changed to require this here?

This binds a default ttir device to the enclosing module. The translate command below relies on a device being bound during flatbuffer serialization.

kmabeeTT · 2024-08-01T18:54:10Z

runtime/lib/ttnn/runtime.cpp

@@ -175,6 +177,8 @@ Event submit(Device deviceHandle, Binary executableHandle,
  return Event(nullptr);
 }

-void wait(Event) { throw std::runtime_error("Not implemented"); }
+void wait(Event) {
+  // Not implemented


do we need to get rid of the runtime error that was here?

Yes because otherwise a FE runtime cannot generically use this interface. It's perfectly legal for a backend runtime to implement submit as blocking and wait is a nop. Which is how ttnn runtime works today, but we should fix this and use real ttnn events and nonblocking APIs

kmabeeTT · 2024-08-01T20:07:49Z

include/ttmlir/Target/TTMetal/binary.fbs

 }

 root_type TTMetalBinary;
-file_identifier "TTB0";
-file_extension "ttb";
+file_identifier "TTM0";


What's the 0 suffix here? Don't see it on TTNN/binary.fbs sibling file.

It has to be a 4 character code, I figured it might be useful one day to put the major version here or something? No good reason, it's arbitrary.

kmabeeTT · 2024-08-01T20:11:16Z

include/ttmlir/Target/TTMetal/command.fbs

@@ -3,39 +3,64 @@ include "program.fbs";
 namespace tt.target.metal;


-table DispatchCommand {
+table EnqueueProgramCommand {


I didn't get far enough to see how these are used or understand if this is initial/bringup limitation, but all these new commands seem rather...bare... (?)... it makes me wonder if they need to more closely match host_api functions in terms of number of arguments?

They appear somewhat bare, but most of them hold some sort of Desc or Ref, with an underlying Desc that fans out to quite a bit of metadata.

kmabeeTT · 2024-08-01T20:15:07Z

include/ttmlir/Target/Utils/MLIRToFlatbuffer.h

@@ -29,6 +31,28 @@ inline ::tt::target::OOBVal toFlatbuffer(FlatbufferObjectCache &,
  }
 }

+inline std::uint64_t getElementSizeBytes(DataType dtype) {


Nit, looks like we have same function in runtime/utils.h called dataTypeElementSize(), perhaps they should be commonized?

I hear your point, and I might be swayed one way or the other, but currently isn't any shared code between compiler and runtime except for FBS. It is somewhat deliberate in that it might be good to leave the door open for one day where the flatbuffer runtime grows out of scope of just tt-mlir and becomes its own repo.

kmabeeTT · 2024-08-01T21:10:33Z

runtime/lib/binary.cpp

@@ -102,6 +104,74 @@ std::vector<TensorDesc> getProgramOutputs(Flatbuffer binary,

 } // namespace ttnn

+namespace metal {


Looks like this stuff (from eyeball check) is quite common with ttnn versions above? Do we keep them separate because they will soon diverge?

They look the same, but the flatbuffer generated functions they call are different and flatbuffer controls the binary layout so they really cannot be mixed. This can be a subtle gotcha.

kmabeeTT · 2024-08-01T23:44:34Z

runtime/lib/ttmetal/runtime.cpp

+  return ::tt::target::metal::GetSizePrefixedTTMetalBinary(binary.handle.get());
+}
+
+std::pair<SystemDesc, DeviceIds> getCurrentSystemDesc() {


In general, should we put TODO/comment with words for temporary hacks/missing functionality? Not seeing much of that in our tree.

I agree, but they should link to an issue. getCurrentSystemDesc I think is being cleaned up by Jackson

kmabeeTT · 2024-08-01T23:51:33Z

runtime/lib/ttmetal/runtime.cpp

+  ::tt::target::Dim2d deviceGrid = toFlatbuffer(device->logical_grid_size());
+  std::vector<::flatbuffers::Offset<tt::target::ChipDesc>> chipDescs = {
+      ::tt::target::CreateChipDesc(
+          fbb, toFlatbuffer(device->arch()), &deviceGrid, (1 << 20), 12,


Looks like ttnn is the same, but will these hardcoded (wormhole?) values (1<<20), 12, etc come from some place that is less...hidden... one day?

Yeah it should come from the device l1_size_per_core I think. But we should be sure this doesn't include the fast dispatch reserved space.

kmabeeTT · 2024-08-02T00:23:34Z

runtime/lib/ttmetal/runtime.cpp

+}
+
+void wait(Event event) {
+  Events events = event.as<Events>();


I'm probably missing something here, but is this intentional (converting from what looks like singular Event, to vector of metal Events)?

Yes, notice above that line 156 creates a vector of metal events. This is because there might be many underlying events created and tracked from this program submit (i.e. executing on multi-device / multi-command queue) that all need to be wrapped up in a single event. To the FE runtime it should only care about managing barriers at program submit granularity.

kmabeeTT · 2024-08-02T00:38:37Z

runtime/lib/ttmetal/command_queue.cpp

+    ::tt::target::metal::EventQueryCommand const *command) {
+  auto event = events.at(command->ref()->global_id());
+  (void)::tt::tt_metal::EventQuery(
+      event); // todo, we need flatbuffer support for tracking and doing


Don't quite understand how this all works yet, but in my head having hard time visualizing how EventQuery (returns bool in metal, true if event is completed, false if not completed) can be used here or what makes sense... does the bool value need to be propogated up, and does that even make sense?

At the moment it makes no sense, in order to support this we need to augment flatbuffer:

table BoolRef { global_id: uint32; } table EventQueryCommand { ref: EventRef; out: BoolRef; }

And then this code would change to be:

boolVars[command->out()->global_id()] = ::tt::tt_metal::EventQuery(event);

And then somewhere later in the flatbuffer program it could reference boolVars[command->out()->global_id()] and conditionally do something. But I actually don't have a use case at the moment to leverage this.

In a separate context I could envision a frontend API:

bool eventQuery(Event event);

For which a frontend runtime might want to query if work has finished.

nsmithtt requested review from tapspatel, jnie-TT, kmabeeTT, AleksKnezevic, sdjordjevicTT, rpavlovicTT, mrakitaTT and nobradovictt as code owners July 31, 2024 21:15

nsmithtt force-pushed the nsmith/metal-runtime2 branch 4 times, most recently from 76075c5 to 77ebc82 Compare August 1, 2024 01:09

nsmithtt and others added 6 commits August 1, 2024 12:27

Bootstrap metal runtime

e8c6286

fix tests with --ttir-implicit-device

cb0c23b

fix ttrt build

9d8f12c

fix cmake

ae42694

fix format

c597e10

ttnn.h ignore -Wunknown-warning-option

3f6b21b

nsmithtt force-pushed the nsmith/metal-runtime2 branch from 4bbf261 to 3f6b21b Compare August 1, 2024 17:56

kmabeeTT approved these changes Aug 2, 2024

View reviewed changes

nsmithtt merged commit 8f2bcbb into main Aug 2, 2024
6 checks passed

nsmithtt deleted the nsmith/metal-runtime2 branch August 2, 2024 02:04

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Bootstrap metal runtime #263

Bootstrap metal runtime #263

nsmithtt commented Jul 31, 2024

kmabeeTT commented Aug 1, 2024

kmabeeTT left a comment

kmabeeTT Aug 1, 2024

nsmithtt Aug 2, 2024

kmabeeTT Aug 1, 2024

nsmithtt Aug 2, 2024

kmabeeTT Aug 1, 2024

nsmithtt Aug 2, 2024

kmabeeTT Aug 1, 2024

nsmithtt Aug 2, 2024

kmabeeTT Aug 1, 2024

nsmithtt Aug 2, 2024 •

edited

Loading

kmabeeTT Aug 1, 2024

nsmithtt Aug 2, 2024

kmabeeTT Aug 1, 2024

nsmithtt Aug 2, 2024

kmabeeTT Aug 1, 2024

nsmithtt Aug 2, 2024

kmabeeTT Aug 2, 2024

nsmithtt Aug 2, 2024

kmabeeTT Aug 2, 2024

nsmithtt Aug 2, 2024

		@@ -1,4 +1,4 @@
		// RUN: ttmlir-opt --ttir-layout --ttnn-open-device --convert-ttir-to-ttnn %s > %t.mlir
		// RUN: ttmlir-opt --ttir-implicit-device --ttir-layout --ttnn-open-device --convert-ttir-to-ttnn %s > %t.mlir

		@@ -102,6 +104,74 @@ std::vector<TensorDesc> getProgramOutputs(Flatbuffer binary,

		} // namespace ttnn

		namespace metal {

Bootstrap metal runtime #263

Bootstrap metal runtime #263

Conversation

nsmithtt commented Jul 31, 2024

kmabeeTT commented Aug 1, 2024

kmabeeTT left a 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

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

nsmithtt Aug 2, 2024 • edited Loading

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

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

nsmithtt Aug 2, 2024 •

edited

Loading