A pair of runtime::ttmetal CQExecutor buffer map improvements (#408)

[kmabeeTT]

• Prevent duplicate Buffers from being created inside CreateBufferCommand handler by checking for existence in buffers umap.
• Change to use buffers.erase() in DeallocateBufferCommand to actually remove the entry from buffers umap. Buffer will still be destroyed because it goes out of scope.
• Neither of these help with the original segfault in this ticket but these were found through visual observation.

[kmabeeTT]

Heads up - @nsmithtt, I pushed commit just now to include a quick-and-dirty hack in ttmetal.h's createBufferFromTensorRef() (the OG source of headaches for this ticket) that avoids the segfault by "correcting" the address stroed in tt-metal's BUFFER_MAP, that I hope can be merged to hold us over until proper fix gets made in metal and intregrated here (which I worry might take a week or more). I made it in such a way that it would be harmless if the metal fix + uplift happens to show up unexpectedly and it's not removed, although I will definitely keep my finger on the pulse and not forget about this. Original failing test added back to CI too.

[rpavlovicTT]

Can you do all of this only if (tensorRef->address() != buffer->address())?
Also, is there a reason to have a map copy?

[nsmithtt]

@rpavlovicTT , it's pretty much guaranteed that tensorRef->address() != buffer->address(). This is a hack that reaches into tt-metal data structure to fix things up and avoid a segfault. We have a proper fix in metal in flight and ideally lands in a week, but depending on metal reviewers could take longer. We just wanted something that could unblock our silicon CI. We'll remove this whole block once the fix in metal lands and we do another metal uplift.

[kmabeeTT]

Yep, temporary hack, and yes it's guaranteed tensorRef->address() != buffer->address() today, the buffer->address() picked by metal will be completely unrelated, which is why it's later (after this hack) set to be tensorRef->address().

For the question about map copy, due to how tt::tt_metal::detail::BUFFER_MAP is implemented, it only returns a copy of it's map and so if I did something like this (without storing a single copy of the map) :

if (auto it = buffer_map.value().find(old_key); it != buffer_map.value().end())

it doesn't compile:

/localdev/kmabee/mlir2/runtime/include/tt/runtime/detail/ttmetal.h:197:19: error: object backing the pointer will be destroyed at the end of the full-expression [-Werror,-Wdangling-gsl]
  197 |     if (auto it = buffer_map.value().find(old_key); it != buffer_map.value().end()) {
      |                   ^~~~~~~~~~~~~~~~~~

[nsmithtt]

CI failing with:

python: /__w/tt-mlir/tt-mlir/runtime/include/tt/runtime/types.h:51: T &tt::runtime::detail::RuntimeCheckedObjectImpl::as(DeviceRuntime) [T = std::vector<tt::tt_metal::Device *>]: Assertion `associatedRuntime == expectedRuntime && "Associated runtime does not match expected runtime of cast"' failed.

[kmabeeTT]

CI failing with:

python: /__w/tt-mlir/tt-mlir/runtime/include/tt/runtime/types.h:51: T &tt::runtime::detail::RuntimeCheckedObjectImpl::as(DeviceRuntime) [T = std::vector<tt::tt_metal::Device *>]: Assertion `associatedRuntime == expectedRuntime && "Associated runtime does not match expected runtime of cast"' failed.

Yeah... From offline chat with @jnie-TT and @tapspatel looks like ttrt-run is mixing ttnn and ttmetal binaries together but it's not being handled properly, I will have to keep test excluded for now I think, and open issue for ttrt change/fix with details. Going to remove test change from this PR now if that helps to approve it.

[tapspatel]

Taking a look at the mixing runtimes issue: #563

[tapspatel]

fix merged: 19c5407

[kmabeeTT]

Since Taps fixed mixing test issue , I'll add passing test back and try again with above fix and merge if good. @nsmithtt any other concerns here before you can approve?

[nsmithtt]

Lgtm!