#0: Merge branch 'main' into llama32-vision

.github/workflows/remove-stale-branches.yaml

@@ -2,7 +2,7 @@ name: "[internal] Remove Stale Branches"
 
 on:
   schedule:
-    - cron: "0 0 * * *" # Runs every night at midnight
+    - cron: "0 */6 * * *"   # Runs at midnight, 6AM, noon, 6pm
   workflow_dispatch: # Allows manual trigger
 
 jobs:

.github/workflows/ttnn-run-sweeps.yaml

@@ -14,6 +14,9 @@ on:
           - ccl.line_all_gather
           - ccl.all_gather_n300
           - ccl.all_gather_n300_focused
+          - creation.zeros.zeros
+          - creation.empty.empty
+          - creation.zeros_like.zeros_like
           - eltwise.unary.abs.abs_pytorch2
           - eltwise.unary.relu.relu
           - eltwise.unary.relu.relu_pytorch2
@@ -52,6 +55,8 @@ on:
           - eltwise.unary.expm1.expm1
           - eltwise.unary.tanh.tanh
           - eltwise.unary.tanh.tanh_pytorch2
+          - eltwise.unary.atanh.atanh
+          - eltwise.unary.atan.atan
           - eltwise.unary.sign.sign
           - eltwise.unary.rad2deg.rad2deg
           - eltwise.unary.deg2rad.deg2rad
@@ -98,6 +103,8 @@ on:
           - eltwise.unary.identity.identity
           - eltwise.unary.neg.neg
           - eltwise.unary.sinh.sinh
+          - eltwise.unary.asinh.asinh
+          - eltwise.unary.cosh.cosh
           - eltwise.unary.relu_min.relu_min
           - eltwise.unary.relu_max.relu_max
           - eltwise.unary.softplus.softplus

CODEOWNERS

@@ -120,6 +120,7 @@ tests/sweep_framework/ @xanderchin @jdesousa-TT @sjameelTT
 tests/sweep_framework/sweeps
 tests/sweep_framework/sweeps/eltwise/ @patrickroberts @yan-zaretskiy @eyonland
 tests/sweep_framework/sweeps/conv2d/  @nkpatel-tt @mywoodstock @shwetankTT @sankarmanoj-tt @pavlejosipovic
+tests/sweep_framework/sweeps/data_movement/  @sjameelTT @ntarafdar @jaykru-tt @yugi957
 
 # TTNN Distributed
 ttnn/cpp/ttnn/distributed/ @cfjchu @ayerofieiev-tt @dmakoviichuk-tt

Doxyfile

@@ -499,7 +499,12 @@ NUM_PROC_THREADS       = 1
 # normally produced when WARNINGS is set to YES.
 # The default value is: NO.
 
-EXTRACT_ALL            = NO
+EXTRACT_ALL            = YES
+
+EXTRACT_NAMESPACES     = YES
+
+INLINE_NAMESPACES = YES
+
 
 # If the EXTRACT_PRIVATE tag is set to YES, all private members of a class will
 # be included in the documentation.
@@ -946,7 +951,17 @@ INPUT                  = tt_metal/hw/inc/dataflow_api.h \
                          tt_metal/include/compute_kernel_api/transpose_wh.h \
                          tt_metal/include/compute_kernel_api/untilize.h \
                          tt_metal/include/compute_kernel_api.h \
-                         tt_metal/impl/kernels/kernel_args.hpp
+                         tt_metal/impl/kernels/kernel_args.hpp \
+                         tt_metal/include/tt_metal/metal.hpp \
+                         tt_metal/include/tt_metal/types.hpp \
+                         tt_metal/include/tt_metal/buffer.hpp \
+                         tt_metal/include/tt_metal/command_queue.hpp \
+                         tt_metal/include/tt_metal/device.hpp \
+                         tt_metal/include/tt_metal/event.hpp \
+                         tt_metal/include/tt_metal/kernel.hpp \
+                         tt_metal/include/tt_metal/program.hpp \
+                         tt_metal/include/tt_metal/trace.hpp
+
 
 # This tag can be used to specify the character encoding of the source files
 # that doxygen parses. Internally doxygen uses the UTF-8 encoding. Doxygen uses

docs/source/tt-metalium/tt_metal/apis/host_apis/buffers/AssignGlobalBufferToProgram.rst

@@ -1,5 +1,5 @@
 
 AssignGlobalBufferToProgram
-===========================
+===============================
 
-.. doxygenfunction:: AssignGlobalBufferToProgram(std::shared_ptr<Buffer> buffer, Program& program)
+.. doxygenfunction:: tt::tt_metal::v0::AssignGlobalBufferToProgram(std::shared_ptr<Buffer> buffer, Program& program)

docs/source/tt-metalium/tt_metal/apis/host_apis/buffers/CircularBuffers.rst

@@ -1,12 +1,17 @@
 CircularBuffers
 ================
 
-.. doxygenfunction:: CreateCircularBuffer
+.. doxygenfunction:: tt::tt_metal::v0::CreateCircularBuffer
 
-.. doxygenfunction:: GetCircularBufferConfig
+.. doxygenfunction:: tt::tt_metal::v0::GetCircularBufferConfig
 
-.. doxygenfunction:: UpdateCircularBufferTotalSize
+.. doxygenfunction:: tt::tt_metal::v0::UpdateCircularBufferTotalSize
 
-.. doxygenfunction:: UpdateCircularBufferPageSize
+.. doxygenfunction:: tt::tt_metal::v0::UpdateCircularBufferPageSize
 
-.. doxygenfunction:: UpdateDynamicCircularBufferAddress
+.. doxygenfunction:: tt::tt_metal::v0::UpdateDynamicCircularBufferAddress
+
+
+Version 1
+-------------------------
+.. doxygenfunction:: tt::tt_metal::v1::CreateCircularBuffer

docs/source/tt-metalium/tt_metal/apis/host_apis/buffers/CreateBuffer.rst

@@ -1,5 +1,5 @@
 CreateBuffer
 =================
 
-.. doxygenfunction:: CreateBuffer(const InterleavedBufferConfig & config);
-.. doxygenfunction:: CreateBuffer(const ShardedBufferConfig & config);
+.. doxygenfunction:: tt::tt_metal::v0::CreateBuffer(const InterleavedBufferConfig & config);
+.. doxygenfunction:: tt::tt_metal::v0::CreateBuffer(const ShardedBufferConfig & config);

docs/source/tt-metalium/tt_metal/apis/host_apis/buffers/CreateSemaphore.rst

@@ -1,4 +1,4 @@
 CreateSemaphore
 ================
 
-.. doxygenfunction:: CreateSemaphore(Program &program, const std::variant<CoreRange,CoreRangeSet> &core_spec, uint32_t initial_value, CoreType core_type)
+.. doxygenfunction:: tt::tt_metal::v0::CreateSemaphore(Program &program, const std::variant<CoreRange,CoreRangeSet> &core_spec, uint32_t initial_value, CoreType core_type)

docs/source/tt-metalium/tt_metal/apis/host_apis/buffers/DeallocateBuffer.rst

@@ -1,4 +1,4 @@
 DeallocateBuffer
 =================
 
-.. doxygenfunction:: DeallocateBuffer
+.. doxygenfunction:: tt::tt_metal::v0::DeallocateBuffer

docs/source/tt-metalium/tt_metal/apis/host_apis/command_queue/BeginTraceCapture.rst

@@ -1,4 +1,4 @@
 BeginTraceCapture
 =================
 
-.. doxygenfunction:: BeginTraceCapture
+.. doxygenfunction:: tt::tt_metal::v0::BeginTraceCapture

docs/source/tt-metalium/tt_metal/apis/host_apis/command_queue/EndTraceCapture.rst

@@ -1,4 +1,4 @@
 EndTraceCapture
 ===============
 
-.. doxygenfunction:: EndTraceCapture
+.. doxygenfunction:: tt::tt_metal::v0::EndTraceCapture

docs/source/tt-metalium/tt_metal/apis/host_apis/command_queue/EnqueueProgram.rst

@@ -1,4 +1,4 @@
 EnqueueProgram
 ==============
 
-.. doxygenfunction:: EnqueueProgram(CommandQueue& cq, Program& program, bool blocking)
+.. doxygenfunction:: tt::tt_metal::v0::EnqueueProgram(CommandQueue& cq, Program& program, bool blocking)

docs/source/tt-metalium/tt_metal/apis/host_apis/command_queue/EnqueueReadBuffer.rst

@@ -1,5 +1,5 @@
 EnqueueReadBuffer
 ==================
 
-.. doxygenfunction:: EnqueueReadBuffer(CommandQueue& cq, std::variant<std::reference_wrapper<Buffer>, std::shared_ptr<Buffer> > buffer, std::vector<uint32_t>& dst, bool blocking)
-.. doxygenfunction:: EnqueueReadBuffer(CommandQueue& cq, std::variant<std::reference_wrapper<Buffer>, std::shared_ptr<Buffer> > buffer, void * dst, bool blocking)
+.. doxygenfunction:: tt::tt_metal::v0::EnqueueReadBuffer(CommandQueue& cq, std::variant<std::reference_wrapper<Buffer>, std::shared_ptr<Buffer> > buffer, std::vector<uint32_t>& dst, bool blocking)
+.. doxygenfunction:: tt::tt_metal::v0::EnqueueReadBuffer(CommandQueue& cq, std::variant<std::reference_wrapper<Buffer>, std::shared_ptr<Buffer> > buffer, void * dst, bool blocking)

docs/source/tt-metalium/tt_metal/apis/host_apis/command_queue/EnqueueRecordEvent.rst

@@ -1,4 +1,4 @@
 EnqueueRecordEvent
 ==================
 
-.. doxygenfunction:: EnqueueRecordEvent
+.. doxygenfunction:: tt::tt_metal::v0::EnqueueRecordEvent

docs/source/tt-metalium/tt_metal/apis/host_apis/command_queue/EnqueueTrace.rst

@@ -1,4 +1,4 @@
 EnqueueTrace
 ============
 
-.. doxygenfunction:: EnqueueTrace(CommandQueue &cq, uint32_t trace_id, bool blocking)
+.. doxygenfunction:: tt::tt_metal::v0::EnqueueTrace(CommandQueue &cq, uint32_t trace_id, bool blocking)

docs/source/tt-metalium/tt_metal/apis/host_apis/command_queue/EnqueueWaitForEvent.rst

@@ -1,4 +1,4 @@
 EnqueueWaitForEvent
 ===================
 
-.. doxygenfunction:: EnqueueWaitForEvent
+.. doxygenfunction:: tt::tt_metal::v0::EnqueueWaitForEvent

docs/source/tt-metalium/tt_metal/apis/host_apis/command_queue/EnqueueWriteBuffer.rst

@@ -1,5 +1,5 @@
 EnqueueWriteBuffer
 ==================
 
-.. doxygenfunction:: EnqueueWriteBuffer(CommandQueue& cq, std::variant<std::reference_wrapper<Buffer>, std::shared_ptr<Buffer> > buffer, std::vector<uint32_t>& src, bool blocking)
-.. doxygenfunction:: EnqueueWriteBuffer(CommandQueue& cq, std::variant<std::reference_wrapper<Buffer>, std::shared_ptr<Buffer> > buffer, HostDataType src, bool blocking)
+.. doxygenfunction:: tt::tt_metal::v0::EnqueueWriteBuffer(CommandQueue& cq, std::variant<std::reference_wrapper<Buffer>, std::shared_ptr<Buffer> > buffer, std::vector<uint32_t>& src, bool blocking)
+.. doxygenfunction:: tt::tt_metal::v0::EnqueueWriteBuffer(CommandQueue& cq, std::variant<std::reference_wrapper<Buffer>, std::shared_ptr<Buffer> > buffer, HostDataType src, bool blocking)

docs/source/tt-metalium/tt_metal/apis/host_apis/command_queue/EventQuery.rst

@@ -1,4 +1,4 @@
 EventQuery
 ==========
 
-.. doxygenfunction:: EventQuery
+.. doxygenfunction:: tt::tt_metal::v0::EventQuery

docs/source/tt-metalium/tt_metal/apis/host_apis/command_queue/EventSynchronize.rst

@@ -1,4 +1,4 @@
 EventSynchronize
 ================
 
-.. doxygenfunction:: EventSynchronize
+.. doxygenfunction:: tt::tt_metal::v0::EventSynchronize

docs/source/tt-metalium/tt_metal/apis/host_apis/command_queue/Finish.rst

@@ -3,4 +3,4 @@
 Finish
 ======
 
-.. doxygenfunction:: Finish(CommandQueue& cq)
+.. doxygenfunction:: tt::tt_metal::v0::Finish(CommandQueue& cq)

docs/source/tt-metalium/tt_metal/apis/host_apis/command_queue/ReleaseTrace.rst

@@ -1,4 +1,4 @@
 ReleaseTrace
 ============
 
-.. doxygenfunction:: ReleaseTrace
+.. doxygenfunction:: tt::tt_metal::v0::ReleaseTrace

docs/source/tt-metalium/tt_metal/apis/host_apis/command_queue/ReplayTrace.rst

@@ -1,4 +1,4 @@
 ReplayTrace
 ===========
 
-.. doxygenfunction:: ReplayTrace
+.. doxygenfunction:: tt::tt_metal::v0::ReplayTrace

docs/source/tt-metalium/tt_metal/apis/host_apis/command_queue/Synchronize.rst

@@ -3,4 +3,4 @@
 Synchronize
 ===========
 
-.. doxygenfunction:: Synchronize
+.. doxygenfunction:: tt::tt_metal::v0::Synchronize

docs/source/tt-metalium/tt_metal/apis/host_apis/device_management/CloseDevice.rst

@@ -2,4 +2,4 @@
 CloseDevice
 =============
 
-.. doxygenfunction:: CloseDevice
+.. doxygenfunction:: tt::tt_metal::v0::CloseDevice

docs/source/tt-metalium/tt_metal/apis/host_apis/device_management/CreateDevice.rst

@@ -1,4 +1,4 @@
 CreateDevice
 =============
 
-.. doxygenfunction:: CreateDevice
+.. doxygenfunction:: tt::tt_metal::v0::CreateDevice

docs/source/tt-metalium/tt_metal/apis/host_apis/device_management/QueryDevices.rst

@@ -1,6 +1,6 @@
 QueryDevices
 =============
 
-.. doxygenfunction:: GetNumAvailableDevices
+.. doxygenfunction:: tt::tt_metal::v0::GetNumAvailableDevices
 
-.. doxygenfunction:: GetNumPCIeDevices
+.. doxygenfunction:: tt::tt_metal::v0::GetNumPCIeDevices

docs/source/tt-metalium/tt_metal/apis/host_apis/kernels/CreateKernel.rst

@@ -1,4 +1,4 @@
 CreateKernel
 ====================
 
-.. doxygenfunction:: CreateKernel
+.. doxygenfunction:: tt::tt_metal::v0::CreateKernel

docs/source/tt-metalium/tt_metal/apis/host_apis/kernels/CreateKernelFromString.rst

@@ -1,4 +1,4 @@
 CreateKernelFromString
 =======================
 
-.. doxygenfunction:: CreateKernelFromString
+.. doxygenfunction:: tt::tt_metal::v0::CreateKernelFromString

docs/source/tt-metalium/tt_metal/apis/host_apis/profiler/DumpDeviceProfileResults.rst

@@ -3,4 +3,4 @@
 DumpDeviceProfileResults
 ========================
 
-.. doxygenfunction:: DumpDeviceProfileResults(Device *device, const Program &program);
+.. doxygenfunction:: tt::tt_metal::v0::DumpDeviceProfileResults(Device *device, const Program &program);

docs/source/tt-metalium/tt_metal/apis/host_apis/program/CreateProgram.rst

@@ -1,4 +1,4 @@
 CreateProgram
 ========================
 
-.. doxygenfunction:: CreateProgram()
+.. doxygenfunction:: tt::tt_metal::v0::CreateProgram()

docs/source/tt-metalium/tt_metal/apis/host_apis/runtime_args/runtime_args.rst

@@ -1,18 +1,18 @@
 Runtime Arguments
 ==================
 
-.. doxygenfunction:: SetRuntimeArgs(const Program &program, KernelHandle kernel_id, const std::variant<CoreCoord,CoreRange,CoreRangeSet> &logical_core, stl::Span<const uint32_t> runtime_args)
+.. doxygenfunction:: tt::tt_metal::v0::SetRuntimeArgs(const Program &program, KernelHandle kernel_id, const std::variant<CoreCoord,CoreRange,CoreRangeSet> &logical_core, stl::Span<const uint32_t> runtime_args)
 
-.. doxygenfunction:: SetRuntimeArgs(const Program &program, KernelHandle kernel, const std::vector< CoreCoord > & core_spec, const std::vector< std::vector<uint32_t> > &runtime_args)
+.. doxygenfunction:: tt::tt_metal::v0::SetRuntimeArgs(const Program &program, KernelHandle kernel, const std::vector< CoreCoord > & core_spec, const std::vector< std::vector<uint32_t> > &runtime_args)
 
-.. doxygenfunction:: SetRuntimeArgs(Device* device, const std::shared_ptr<Kernel> kernel, const std::variant<CoreCoord, CoreRange, CoreRangeSet> &core_spec, std::shared_ptr<RuntimeArgs> runtime_args)
+.. doxygenfunction:: tt::tt_metal::v0::SetRuntimeArgs(Device* device, const std::shared_ptr<Kernel> kernel, const std::variant<CoreCoord, CoreRange, CoreRangeSet> &core_spec, std::shared_ptr<RuntimeArgs> runtime_args)
 
-.. doxygenfunction:: SetRuntimeArgs(Device* device, const std::shared_ptr<Kernel> kernel, const std::vector< CoreCoord > & core_spec, const std::vector<std::shared_ptr<RuntimeArgs>> runtime_args)
+.. doxygenfunction:: tt::tt_metal::v0::SetRuntimeArgs(Device* device, const std::shared_ptr<Kernel> kernel, const std::vector< CoreCoord > & core_spec, const std::vector<std::shared_ptr<RuntimeArgs>> runtime_args)
 
-.. doxygenfunction:: GetRuntimeArgs(const Program &program, KernelHandle kernel_id, const CoreCoord &logical_core)
+.. doxygenfunction:: tt::tt_metal::v0::GetRuntimeArgs(const Program &program, KernelHandle kernel_id, const CoreCoord &logical_core)
 
-.. doxygenfunction:: GetRuntimeArgs(const Program &program, KernelHandle kernel_id)
+.. doxygenfunction:: tt::tt_metal::v0::GetRuntimeArgs(const Program &program, KernelHandle kernel_id)
 
-.. doxygenfunction:: SetCommonRuntimeArgs(const Program &program, KernelHandle kernel_id, stl::Span<const uint32_t> runtime_args)
+.. doxygenfunction:: tt::tt_metal::v0::SetCommonRuntimeArgs(const Program &program, KernelHandle kernel_id, stl::Span<const uint32_t> runtime_args)
 
-.. doxygenfunction:: GetCommonRuntimeArgs(const Program &program, KernelHandle kernel_id)
+.. doxygenfunction:: tt::tt_metal::v0::GetCommonRuntimeArgs(const Program &program, KernelHandle kernel_id)

docs/source/tt-metalium/tt_metal/apis/kernel_apis/circular_buffers/cb_pages_available_at_front.rst

@@ -0,0 +1,4 @@
+cb_pages_available_at_front
+===========================
+
+.. doxygenfunction:: cb_pages_available_at_front(int32_t operand, int32_t num_pages)

docs/source/tt-metalium/tt_metal/apis/kernel_apis/circular_buffers/cb_pages_reservable_at_back.rst

@@ -0,0 +1,4 @@
+cb_pages_reservable_at_back
+===========================
+
+.. doxygenfunction:: cb_pages_reservable_at_back(int32_t operand, int32_t num_pages)

docs/source/tt-metalium/tt_metal/apis/kernel_apis/circular_buffers/circular_buffers.rst

@@ -4,7 +4,9 @@ Circular Buffer APIs
 Circular buffers are used for communication between threads of the Tensix core. They act as limited capacity double-ended queues with producers pushing tiles to the back of the queue and consumers popping tiles off the front of the queue.
 
 .. toctree::
+  cb_pages_available_at_front
   cb_wait_front
+  cb_pages_reservable_at_back
   cb_reserve_back
   cb_push_back
   cb_pop_front

docs/source/tt-metalium/tt_metal/examples/eltwise_binary.rst

@@ -37,19 +37,19 @@ We already have set the circular buffers needed for compute data communication.
   constexpr uint32_t num_input_tiles = 2;
   constexpr uint32_t input_cb_size = num_input_tiles * single_tile_size;
   CircularBufferConfig cb_src0_config = CircularBufferConfig(input_cb_size, {{src0_cb_index, tt::DataFormat::Float16_b}}, src0_cb_addr).set_page_size(src0_cb_index, single_tile_size);
-  CBHandle cb_src0 = CreateCircularBuffer(program, core, cb_src0_config);
+  CBHandle cb_src0 = v0::CreateCircularBuffer(program, core, cb_src0_config);
 
   constexpr uint32_t src1_cb_index = CB::c_in1;
   constexpr uint32_t src1_cb_addr = 300 * 1024;
   CircularBufferConfig cb_src1_config = CircularBufferConfig(input_cb_size, {{src1_cb_index, tt::DataFormat::Float16_b}}, src1_cb_addr).set_page_size(src1_cb_index, single_tile_size);
-  CBHandle cb_src1 = CreateCircularBuffer(program, core, cb_src1_config);
+  CBHandle cb_src1 = v0::CreateCircularBuffer(program, core, cb_src1_config);
 
   constexpr uint32_t output_cb_index = CB::c_out0;
   constexpr uint32_t output_cb_addr = 400 * 1024;
   constexpr uint32_t num_output_tiles = 2;
   constexpr uint32_t input_cb_size = num_input_tiles * single_tile_size;
   CircularBufferConfig cb_output_config = CircularBufferConfig(input_cb_size, {{output_cb_index, tt::DataFormat::Float16_b}}, output_cb_addr).set_page_size(output_cb_index, single_tile_size);
-  CBHandle cb_output = CreateCircularBuffer(program, core, cb_output);
+  CBHandle cb_output = v0::CreateCircularBuffer(program, core, cb_output);
 
 We will create two input circular buffers to accommodate our two input tensors,
 and an output one for the result of the eltwise binary operation.

docs/source/tt-metalium/tt_metal/examples/eltwise_sfpu.rst

@@ -34,12 +34,12 @@ compute, and writer engines.
     constexpr uint32_t src0_cb_index = CB::c_in0;
     constexpr uint32_t num_input_tiles = 2;
     CircularBufferConfig cb_src0_config = CircularBufferConfig(num_input_tiles * single_tile_size, {{src0_cb_index, tt::DataFormat::Float16_b}}).set_page_size(src0_cb_index, single_tile_size);
-    CBHandle cb_src0 = tt_metal::CreateCircularBuffer(program, core, cb_src0_config);
+    CBHandle cb_src0 = tt_metal::v0::CreateCircularBuffer(program, core, cb_src0_config);
 
     constexpr uint32_t output_cb_index = CB::c_out0;
     constexpr uint32_t num_output_tiles = 2;
     CircularBufferConfig cb_output_config = CircularBufferConfig(num_output_tiles * single_tile_size, {{output_cb_index, tt::DataFormat::Float16_b}}).set_page_size(output_cb_index, single_tile_size);
-    CBHandle cb_output = tt_metal::CreateCircularBuffer(program, core, cb_output_config);
+    CBHandle cb_output = tt_metal::v0::CreateCircularBuffer(program, core, cb_output_config);
 
 We will create one input circular buffers to accommodate our input tensor,
 and an output one for the result of the eltwise sfpu operation.

docs/source/tt-metalium/tt_metal/examples/matmul_multi_core_optimizations/data_mcast.rst

@@ -111,13 +111,13 @@ Recall in our data reuse example, we created our L1 circular buffers for all the
 
 .. code-block:: cpp
 
-    auto cb_output = tt_metal::CreateCircularBuffer(program, all_cores, cb_output_config);
+    auto cb_output = tt_metal::v0::CreateCircularBuffer(program, all_cores, cb_output_config);
 
 METALIUM also allows us to pass all of our CoreRanges defined above through a ``CoreRangeSet(...)`` function call as the 2nd argument.  Let's do so with the following:
 
 .. code-block:: cpp
 
-    auto cb_output = tt_metal::CreateCircularBuffer(program, CoreRangeSet({all_cores}), cb_output_config);
+    auto cb_output = tt_metal::v0::CreateCircularBuffer(program, CoreRangeSet({all_cores}), cb_output_config);
 
 In fact, you can instantiate circular buffers on any one of these three options: ``const std::variant<CoreCoord, CoreRange, CoreRangeSet>``.  Please refer to the CircularBuffers page for further details.
 

docs/source/tt-metalium/tt_metal/examples/matmul_multi_core_optimizations/data_reuse.rst

@@ -49,7 +49,7 @@ In addition to our double-buffer config, we introduce a third circular buffer de
     CircularBufferConfig cb_output_config = CircularBufferConfig(out_CB_size, output_cb_data_format_spec)
         .set_page_size(output_cb_index, single_tile_size)
         .set_page_size(interm0_cb_index, single_tile_size);
-    auto cb_output = tt_metal::CreateCircularBuffer(program, all_cores, cb_output_config);
+    auto cb_output = tt_metal::v0::CreateCircularBuffer(program, all_cores, cb_output_config);
 
 Stride Kernel Arguments
 -----------------------

docs/source/tt-metalium/tt_metal/examples/matmul_single_core.rst

@@ -151,18 +151,18 @@ double buffering..
     uint32_t num_input_tiles = 2;
     tt_metal::CircularBufferConfig cb_src0_config = tt_metal::CircularBufferConfig(num_input_tiles * single_tile_size, {{src0_cb_index, cb_data_format}})
         .set_page_size(src0_cb_index, single_tile_size);
-    auto cb_src0 = tt_metal::CreateCircularBuffer(program, core, cb_src0_config);
+    auto cb_src0 = tt_metal::v0::CreateCircularBuffer(program, core, cb_src0_config);
 
     uint32_t src1_cb_index = CB::c_in1; // 1
     tt_metal::CircularBufferConfig cb_src1_config = tt_metal::CircularBufferConfig(num_input_tiles * single_tile_size, {{src1_cb_index, cb_data_format}})
         .set_page_size(src1_cb_index, single_tile_size);
-    auto cb_src1 = tt_metal::CreateCircularBuffer(program, core, cb_src1_config);
+    auto cb_src1 = tt_metal::v0::CreateCircularBuffer(program, core, cb_src1_config);
 
     uint32_t output_cb_index = CB::c_out0; // output operands start at index 16
     uint32_t num_output_tiles = 2;
     tt_metal::CircularBufferConfig cb_output_config = tt_metal::CircularBufferConfig(num_output_tiles * single_tile_size, {{output_cb_index, cb_data_format}})
         .set_page_size(output_cb_index, single_tile_size);
-    auto cb_output = tt_metal::CreateCircularBuffer(program, core, cb_output_config);
+    auto cb_output = tt_metal::v0::CreateCircularBuffer(program, core, cb_output_config);
 
 Compile-time kernels arguments
 ------------------------------