#4984: Move CB configs into kernel config ring buffer

tests/tt_metal/tt_metal/test_kernels/dataflow/unit_tests/command_queue/random_program.cpp

@@ -39,7 +39,11 @@ void MAIN  {
     constexpr volatile uint32_t page_size = get_compile_time_arg_val(7);
 
     for (uint32_t i = 0; i < num_cbs; i++) {
-        uint32_t cb_val = reinterpret_cast<volatile tt_l1_ptr uint32_t*>(CIRCULAR_BUFFER_CONFIG_BASE + i * 16)[3];
+        tt_l1_ptr mailboxes_t* const mailboxes = (tt_l1_ptr mailboxes_t*)(MEM_MAILBOX_BASE);
+        uint32_t kernel_config_base = mailboxes->launch.kernel_config.kernel_config_base;
+        uint32_t tt_l1_ptr *cb_l1_base = (uint32_t tt_l1_ptr *)(kernel_config_base +
+            mailboxes->launch.kernel_config.cb_offset);
+        uint32_t cb_val = reinterpret_cast<volatile tt_l1_ptr uint32_t*>(cb_l1_base + i * 4)[3];
         uint32_t expected = ((i + 1) * page_size) >> 4;
         if (cb_val != expected) {
             DPRINT << "Problem with CB idx: " << i << " Expected: " << expected << " Got: " << cb_val << ENDL();

tests/tt_metal/tt_metal/unit_tests/circular_buffer/test_CircularBuffer_allocation.cpp

@@ -25,7 +25,7 @@ void validate_cb_address(Program &program, Device *device, const CoreRangeSet &c
             for (auto y = core_range.start_coord.y; y <= core_range.end_coord.y; y++) {
                 CoreCoord core_coord(x, y);
                 tt::tt_metal::detail::ReadFromDeviceL1(
-                    device, core_coord, CIRCULAR_BUFFER_CONFIG_BASE, cb_config_buffer_size, cb_config_vector);
+                    device, core_coord, program.get_cb_base_addr(device, core_coord, CoreType::WORKER), cb_config_buffer_size, cb_config_vector);
 
                 std::map<uint8_t, uint32_t> address_per_buffer_index = core_to_address_per_buffer_index.at(core_coord);
 
@@ -284,6 +284,7 @@ TEST_F(DeviceFixture, TestUpdateCircularBufferAddress) {
 
 TEST_F(DeviceFixture, TestUpdateCircularBufferPageSize) {
   for (unsigned int id = 0; id < num_devices_; id++) {
+    Device *device = this->devices_.at(id);
     Program program;
     CBConfig cb_config;
     CoreCoord core0(0, 0);
@@ -306,7 +307,7 @@ TEST_F(DeviceFixture, TestUpdateCircularBufferPageSize) {
         expected_cb_addr += cb_config.page_size;
     }
 
-    detail::LaunchProgram(this->devices_.at(id), program);
+    detail::LaunchProgram(device, program);
 
     vector<uint32_t> cb_config_vector;
     uint32_t cb_config_buffer_size = NUM_CIRCULAR_BUFFERS * UINT32_WORDS_PER_CIRCULAR_BUFFER_CONFIG * sizeof(uint32_t);
@@ -316,7 +317,7 @@ TEST_F(DeviceFixture, TestUpdateCircularBufferPageSize) {
             for (auto y = core_range.start_coord.y; y <= core_range.end_coord.y; y++) {
                 CoreCoord core_coord(x, y);
                 tt::tt_metal::detail::ReadFromDeviceL1(
-                    this->devices_.at(id), core_coord, CIRCULAR_BUFFER_CONFIG_BASE, cb_config_buffer_size, cb_config_vector);
+                    device, core_coord, program.get_cb_base_addr(device, core_coord, CoreType::WORKER), cb_config_buffer_size, cb_config_vector);
 
                 std::map<uint8_t, uint32_t> address_per_buffer_index = golden_addresses_per_core.at(core_coord);
                 std::map<uint8_t, uint32_t> num_pages_per_buffer_index = golden_num_pages_per_core.at(core_coord);
@@ -333,15 +334,15 @@ TEST_F(DeviceFixture, TestUpdateCircularBufferPageSize) {
     UpdateCircularBufferPageSize(program, cb_ids[1], 1, cb_config.page_size / 2);
     golden_num_pages_per_core[core0][1] = 2;
 
-    detail::LaunchProgram(this->devices_.at(id), program);
+    detail::LaunchProgram(device, program);
 
     // addresses should not be changed
     for (const CoreRange &core_range : cr_set.ranges()) {
         for (auto x = core_range.start_coord.x; x <= core_range.end_coord.x; x++) {
             for (auto y = core_range.start_coord.y; y <= core_range.end_coord.y; y++) {
                 CoreCoord core_coord(x, y);
                 tt::tt_metal::detail::ReadFromDeviceL1(
-                    this->devices_.at(id), core_coord, CIRCULAR_BUFFER_CONFIG_BASE, cb_config_buffer_size, cb_config_vector);
+                    device, core_coord, program.get_cb_base_addr(device, core_coord, CoreType::WORKER), cb_config_buffer_size, cb_config_vector);
 
                 std::map<uint8_t, uint32_t> address_per_buffer_index = golden_addresses_per_core.at(core_coord);
                 std::map<uint8_t, uint32_t> num_pages_per_buffer_index = golden_num_pages_per_core.at(core_coord);

tests/tt_metal/tt_metal/unit_tests/circular_buffer/test_CircularBuffer_creation.cpp

@@ -28,7 +28,7 @@ bool test_cb_config_written_to_core(Program &program, Device *device, const Core
                 for (auto y = core_range.start_coord.y; y <= core_range.end_coord.y; y++) {
                     CoreCoord core_coord(x, y);
                     tt::tt_metal::detail::ReadFromDeviceL1(
-                        device, core_coord, CIRCULAR_BUFFER_CONFIG_BASE, cb_config_buffer_size, cb_config_vector);
+                        device, core_coord, program.get_sem_base_addr(device, core_coord, CoreType::WORKER), cb_config_buffer_size, cb_config_vector);
 
                     for (const auto &[buffer_index, golden_cb_config] : cb_config_per_buffer_index) {
                         auto base_index = UINT32_WORDS_PER_CIRCULAR_BUFFER_CONFIG * buffer_index;

tests/tt_metal/tt_metal/unit_tests_fast_dispatch/command_queue/test_EnqueueProgram.cpp

@@ -76,7 +76,7 @@ std::vector<CBHandle> initialize_dummy_circular_buffers(Program& program, const
     return cb_handles;
 }
 
-bool cb_config_successful(Device* device, const DummyProgramMultiCBConfig & program_config){
+bool cb_config_successful(Device* device, Program &program, const DummyProgramMultiCBConfig & program_config){
     bool pass = true;
 
     // Need to use old APIs to read since we cannot allocate a buffer in the reserved space we're trying
@@ -86,7 +86,9 @@ bool cb_config_successful(Device* device, const DummyProgramMultiCBConfig & prog
 
     for (const CoreRange& core_range : program_config.cr_set.ranges()) {
         for (const CoreCoord& core_coord : core_range) {
-            tt::tt_metal::detail::ReadFromDeviceL1(device, core_coord, CIRCULAR_BUFFER_CONFIG_BASE, cb_config_buffer_size, cb_config_vector);
+            tt::tt_metal::detail::ReadFromDeviceL1(device, core_coord,
+                program.get_sem_base_addr(device, core_coord, CoreType::WORKER),
+                cb_config_buffer_size, cb_config_vector);
 
             uint32_t cb_addr = L1_UNRESERVED_BASE;
             for (uint32_t i = 0; i < program_config.cb_config_vector.size(); i++) {
@@ -115,7 +117,7 @@ bool test_dummy_EnqueueProgram_with_cbs(Device* device, CommandQueue& cq, DummyP
     EnqueueProgram(cq, program, is_blocking_op);
     Finish(cq);
 
-    return cb_config_successful(device, program_config);
+    return cb_config_successful(device, program, program_config);
 }
 
 bool test_dummy_EnqueueProgram_with_cbs_update_size(Device* device, CommandQueue& cq, const DummyProgramMultiCBConfig& program_config) {
@@ -126,7 +128,7 @@ bool test_dummy_EnqueueProgram_with_cbs_update_size(Device* device, CommandQueue
     EnqueueProgram(cq, program, false);
     Finish(cq);
 
-    const bool is_cb_config_before_update_successful = cb_config_successful(device, program_config);
+    const bool is_cb_config_before_update_successful = cb_config_successful(device, program, program_config);
 
     DummyProgramMultiCBConfig program_config_2 = program_config;
     for (uint32_t cb_id = 0; cb_id < program_config.cb_config_vector.size(); cb_id++) {
@@ -139,7 +141,7 @@ bool test_dummy_EnqueueProgram_with_cbs_update_size(Device* device, CommandQueue
     EnqueueProgram(cq, program, false);
     Finish(cq);
 
-    const bool is_cb_config_after_update_successful = cb_config_successful(device, program_config_2);
+    const bool is_cb_config_after_update_successful = cb_config_successful(device, program, program_config_2);
     return is_cb_config_before_update_successful && is_cb_config_after_update_successful;
 }
 
@@ -776,7 +778,8 @@ TEST_F(CommandQueueSingleCardFixture, TestMultiCBSharedAddressSpaceSentSingleCor
         vector<uint32_t> cb_config_vector;
 
         tt::tt_metal::detail::ReadFromDeviceL1(
-            device, core_coord, CIRCULAR_BUFFER_CONFIG_BASE, cb_config_buffer_size, cb_config_vector);
+            device, core_coord,
+            program.get_cb_base_addr(device, core_coord, CoreType::WORKER), cb_config_buffer_size, cb_config_vector);
         uint32_t cb_addr = L1_UNRESERVED_BASE;
         uint32_t intermediate_index = intermediate_cb * sizeof(uint32_t);
 

tt_metal/hostdevcommon/common_runtime_address_map.h

@@ -51,26 +51,19 @@ constexpr static std::uint32_t PROFILER_RISC_COUNT = 5;
 static_assert (PROFILER_FULL_HOST_BUFFER_SIZE_PER_RISC > PROFILER_L1_BUFFER_SIZE);
 
 // Kernel config buffer is WIP
-// Will eventually move CBs/Sems and likely kernel bins into this buffer
-// Size is presently based on the old size of the RTAs (large enough to hold 1 set)
+// Size is presently based on the old sizes of the RTAs + CB config + Sems
 // plus some extra space freed up in the mem map
 constexpr static std::uint32_t L1_KERNEL_CONFIG_BASE = PROFILER_L1_END_ADDRESS;
-constexpr static std::uint32_t L1_KERNEL_CONFIG_SIZE = 4 * 1024 + 256 + 128;
+constexpr static std::uint32_t L1_KERNEL_CONFIG_SIZE = 4 * 1024 + 256 + 128 + 512;
 
 constexpr static std::uint32_t IDLE_ERISC_L1_KERNEL_CONFIG_BASE = 32 * 1024;
 
-// config for 32 L1 buffers is at addr BUFFER_CONFIG_BASE
-// 12 bytes for each buffer: (addr, size, size_in_tiles)
-// addr and size are in 16B words (byte address >> 4)
-// this is a total of 32 * 3 * 4 = 384B
-constexpr static std::uint32_t CIRCULAR_BUFFER_CONFIG_BASE = L1_KERNEL_CONFIG_BASE + L1_KERNEL_CONFIG_SIZE;
 constexpr static std::uint32_t NUM_CIRCULAR_BUFFERS = 32;
 constexpr static std::uint32_t UINT32_WORDS_PER_CIRCULAR_BUFFER_CONFIG = 4;
-constexpr static std::uint32_t CIRCULAR_BUFFER_CONFIG_SIZE = NUM_CIRCULAR_BUFFERS * UINT32_WORDS_PER_CIRCULAR_BUFFER_CONFIG * sizeof(uint32_t);
 
 constexpr static std::uint32_t PROFILER_L1_CONTROL_VECTOR_SIZE = 32;
 constexpr static std::uint32_t PROFILER_L1_CONTROL_BUFFER_SIZE = PROFILER_L1_CONTROL_VECTOR_SIZE * sizeof(uint32_t);
-constexpr static std::uint32_t PROFILER_L1_BUFFER_CONTROL = CIRCULAR_BUFFER_CONFIG_BASE + CIRCULAR_BUFFER_CONFIG_SIZE;
+constexpr static std::uint32_t PROFILER_L1_BUFFER_CONTROL = L1_KERNEL_CONFIG_BASE + L1_KERNEL_CONFIG_SIZE;
 
 constexpr static std::uint32_t L1_UNRESERVED_BASE = ((PROFILER_L1_BUFFER_CONTROL + PROFILER_L1_CONTROL_BUFFER_SIZE - 1) | (DRAM_ALIGNMENT - 1)) + 1;
 

tt_metal/hw/firmware/src/brisc.cc

@@ -384,12 +384,16 @@ int main() {
 
             noc_index = mailboxes->launch.kernel_config.brisc_noc_id;
 
-            setup_cb_read_write_interfaces(0, num_cbs_to_early_init, true, true, false);
+            uint32_t kernel_config_base = mailboxes->launch.kernel_config.kernel_config_base;
+
+            uint32_t tt_l1_ptr *cb_l1_base = (uint32_t tt_l1_ptr *)(kernel_config_base +
+                mailboxes->launch.kernel_config.cb_offset);
+            setup_cb_read_write_interfaces(cb_l1_base, 0, num_cbs_to_early_init, true, true, false);
+
             finish_ncrisc_copy_and_run(enables);
 
             // Run the BRISC kernel
             DEBUG_STATUS("R");
-            uint32_t kernel_config_base = mailboxes->launch.kernel_config.kernel_config_base;
             rta_l1_base = (uint32_t tt_l1_ptr *)(kernel_config_base +
                 mailboxes->launch.kernel_config.mem_map[DISPATCH_CLASS_TENSIX_DM0].rta_offset);
             crta_l1_base = (uint32_t tt_l1_ptr *)(kernel_config_base +
@@ -398,7 +402,7 @@ int main() {
                 mailboxes->launch.kernel_config.sem_offset);
 
             if (enables & DISPATCH_CLASS_MASK_TENSIX_ENABLE_DM0) {
-                setup_cb_read_write_interfaces(num_cbs_to_early_init, mailboxes->launch.kernel_config.max_cb_index, true, true, false);
+                setup_cb_read_write_interfaces(cb_l1_base, num_cbs_to_early_init, mailboxes->launch.kernel_config.max_cb_index, true, true, false);
                 kernel_init();
                 RECORD_STACK_USAGE();
             } else {

tt_metal/hw/firmware/src/idle_erisc.cc

@@ -112,13 +112,15 @@ int main() {
 
             noc_index = mailboxes->launch.kernel_config.brisc_noc_id;
 
-            setup_cb_read_write_interfaces(0, mailboxes->launch.kernel_config.max_cb_index, true, true, false);
+            uint32_t kernel_config_base = mailboxes->launch.kernel_config.kernel_config_base;
+
+            uint32_t tt_l1_ptr *cb_l1_base = (uint32_t tt_l1_ptr *)(kernel_config_base +
+                mailboxes->launch.kernel_config.cb_offset);
+            setup_cb_read_write_interfaces(cb_l1_base, 0, mailboxes->launch.kernel_config.max_cb_index, true, true, false);
 
             // Run the ERISC kernel
             DEBUG_STATUS("R");
             //if (mailboxes->launch.enable_brisc) {
-                //UC FIXME: do i need this?
-                uint32_t kernel_config_base = mailboxes->launch.kernel_config.kernel_config_base;
                 rta_l1_base = (uint32_t tt_l1_ptr *)(kernel_config_base +
                     mailboxes->launch.kernel_config.mem_map[DISPATCH_CLASS_ETH_DM0].rta_offset);
                 crta_l1_base = (uint32_t tt_l1_ptr *)(kernel_config_base +

tt_metal/hw/firmware/src/ncrisc.cc

@@ -96,9 +96,13 @@ int main(int argc, char *argv[]) {
         notify_brisc_and_wait();
         DeviceZoneScopedMainN("NCRISC-FW");
 
-        setup_cb_read_write_interfaces(0, mailboxes->launch.kernel_config.max_cb_index, true, true, false);
-
         uint32_t kernel_config_base = mailboxes->launch.kernel_config.kernel_config_base;
+
+        uint32_t tt_l1_ptr *cb_l1_base = (uint32_t tt_l1_ptr *)(kernel_config_base +
+            mailboxes->launch.kernel_config.cb_offset);
+
+        setup_cb_read_write_interfaces(cb_l1_base, 0, mailboxes->launch.kernel_config.max_cb_index, true, true, false);
+
         rta_l1_base = (uint32_t tt_l1_ptr *)(kernel_config_base +
             mailboxes->launch.kernel_config.mem_map[DISPATCH_CLASS_TENSIX_DM1].rta_offset);
         crta_l1_base = (uint32_t tt_l1_ptr *)(kernel_config_base +

tt_metal/hw/firmware/src/trisc.cc

@@ -58,6 +58,7 @@ volatile tt_l1_ptr uint32_t l1_buffer[16] __attribute__((section("l1_data"))) __
 __attribute__((used));
 
 #if !defined(UCK_CHLKC_MATH)
+uint32_t tt_l1_ptr *cb_l1_base __attribute__((used));
 CBInterface cb_interface[NUM_CIRCULAR_BUFFERS] __attribute__((used));
 #endif
 
@@ -107,15 +108,18 @@ int main(int argc, char *argv[]) {
         while (*trisc_run != RUN_SYNC_MSG_GO);
         DeviceZoneScopedMainN("TRISC-FW");
 
+        uint32_t kernel_config_base = mailboxes->launch.kernel_config.kernel_config_base;
+
 #if !defined(UCK_CHLKC_MATH)
-        setup_cb_read_write_interfaces(0, mailboxes->launch.kernel_config.max_cb_index, cb_init_read, cb_init_write, cb_init_write);
+        uint32_t tt_l1_ptr *cb_l1_base = (uint32_t tt_l1_ptr *)(kernel_config_base +
+            mailboxes->launch.kernel_config.cb_offset);
+        setup_cb_read_write_interfaces(cb_l1_base, 0, mailboxes->launch.kernel_config.max_cb_index, cb_init_read, cb_init_write, cb_init_write);
 #if defined(UCK_CHLKC_UNPACK)
         // Hack workaround for issue #11591
         for (volatile uint32_t xxx = 0; xxx < 100; xxx++);
 #endif
 #endif
 
-        uint32_t kernel_config_base = mailboxes->launch.kernel_config.kernel_config_base;
         rta_l1_base = (uint32_t tt_l1_ptr *)(kernel_config_base +
             mailboxes->launch.kernel_config.mem_map[DISPATCH_CLASS_TENSIX_COMPUTE].rta_offset);
         crta_l1_base = (uint32_t tt_l1_ptr *)(kernel_config_base +

tt_metal/hw/inc/blackhole/eth_l1_address_map.h

@@ -28,7 +28,9 @@ struct address_map {
   static constexpr std::int32_t DATA_BUFFER_SIZE_ETH = 4 * 1024;
   static constexpr std::int32_t DATA_BUFFER_SIZE_NOC = 16 * 1024;
   static constexpr std::int32_t DATA_BUFFER_SIZE = 24 * 1024;
-  static constexpr std::int32_t ERISC_L1_KERNEL_CONFIG_SIZE = 96 * 4;
+  // Kernel config buffer is WIP
+  // Size is presently based on the old sizes of the RTAs + CB config + Sems
+  static constexpr std::int32_t ERISC_L1_KERNEL_CONFIG_SIZE = 96 * 4 + 8 * 16;
 
   // Base addresses
   static constexpr std::int32_t FIRMWARE_BASE = 0x9040;
@@ -58,6 +60,8 @@ struct address_map {
   // erisc early exit functionality re-uses mailboxes_t::ncrisc_halt_msg_t::stack_save memory
   static constexpr std::int32_t ERISC_MEM_MAILBOX_STACK_SAVE = ERISC_MEM_MAILBOX_BASE + 4;
 
+  // Kernel config buffer is WIP
+  // Size is presently based on the old sizes of the RTAs + CB config + Sems
   static constexpr std::uint32_t PROFILER_L1_BUFFER_ER = ERISC_MEM_MAILBOX_BASE + 288 + 256 + 16;
   static constexpr std::uint32_t PROFILER_L1_BUFFER_CONTROL = PROFILER_L1_BUFFER_ER + PROFILER_L1_BUFFER_SIZE;
 

tt_metal/hw/inc/circular_buffer.h

@@ -51,8 +51,11 @@ extern CBInterface cb_interface[NUM_CIRCULAR_BUFFERS];
 
 // NCRISC and BRISC setup read and write
 // TRISC sets up read or write
-inline void setup_cb_read_write_interfaces(uint32_t start_cb_index, uint32_t max_cb_index, bool read, bool write, bool init_wr_tile_ptr) {
-    volatile tt_l1_ptr uint32_t* circular_buffer_config_addr = (volatile tt_l1_ptr uint32_t*)(CIRCULAR_BUFFER_CONFIG_BASE) + start_cb_index * UINT32_WORDS_PER_CIRCULAR_BUFFER_CONFIG;
+inline void setup_cb_read_write_interfaces(uint32_t tt_l1_ptr *cb_l1_base, uint32_t start_cb_index, uint32_t max_cb_index, bool read, bool write, bool init_wr_tile_ptr) {
+
+    constexpr uint32_t WORDS_PER_CIRCULAR_BUFFER_CONFIG = 4;
+
+    volatile tt_l1_ptr uint32_t* circular_buffer_config_addr = cb_l1_base + start_cb_index * WORDS_PER_CIRCULAR_BUFFER_CONFIG;
 
     for (uint32_t cb_id = start_cb_index; cb_id < max_cb_index; cb_id++) {
 
@@ -81,6 +84,6 @@ inline void setup_cb_read_write_interfaces(uint32_t start_cb_index, uint32_t max
             cb_interface[cb_id].fifo_wr_tile_ptr = 0;
         }
 
-        circular_buffer_config_addr += UINT32_WORDS_PER_CIRCULAR_BUFFER_CONFIG;
+        circular_buffer_config_addr += WORDS_PER_CIRCULAR_BUFFER_CONFIG;
     }
 }

tt_metal/hw/inc/grayskull/dev_mem_map.h

@@ -54,6 +54,7 @@
 #define MEM_BOOT_CODE_BASE 0
 #define MEM_L1_BARRIER 12
 #define MEM_MAILBOX_BASE 16
+// Magic size must be big enough to hold dev_msgs_t.  static_asserts will fire if this is too small
 #define MEM_MAILBOX_END (MEM_MAILBOX_BASE + 1348)
 #define MEM_IERISC_MAILBOX_BASE 0
 #define MEM_IERISC_MAILBOX_END 0

tt_metal/hw/inc/wormhole/eth_l1_address_map.h

@@ -28,6 +28,8 @@ struct address_map {
   static constexpr std::int32_t DATA_BUFFER_SIZE_ETH = 4 * 1024;
   static constexpr std::int32_t DATA_BUFFER_SIZE_NOC = 16 * 1024;
   static constexpr std::int32_t DATA_BUFFER_SIZE = 24 * 1024;
+  // Kernel config buffer is WIP
+  // Size is presently based on the old sizes of the RTAs + CB config + Sems
   static constexpr std::int32_t ERISC_L1_KERNEL_CONFIG_SIZE = 96 * 4 + 8 * 16;
 
   // Base addresses

tt_metal/impl/dispatch/command_queue.cpp

@@ -1097,6 +1097,7 @@ void EnqueueProgramCommand::assemble_device_commands(
         }
 
         // CB Configs commands
+        index = hal.get_programmable_core_type_index(HalProgrammableCoreType::TENSIX);
         if (num_multicast_cb_sub_cmds > 0) {
             uint32_t curr_sub_cmd_idx = 0;
             cached_program_command_sequence.cb_configs_payloads.reserve(num_multicast_cb_sub_cmds);
@@ -1105,13 +1106,15 @@ void EnqueueProgramCommand::assemble_device_commands(
                 uint32_t write_offset_bytes = program_command_sequence.write_offset_bytes();
                 program_command_sequence.add_dispatch_write_packed<CQDispatchWritePackedMulticastSubCmd>(
                     num_sub_cmds_in_cmd,
-                    CIRCULAR_BUFFER_CONFIG_BASE,
+                    program.get_program_config(index).cb_offset,
                     cb_config_size_bytes,
                     mcast_cb_payload_sizeB,
                     multicast_cb_config_sub_cmds,
                     multicast_cb_config_data,
                     this->packed_write_max_unicast_sub_cmds,
-                    curr_sub_cmd_idx);
+                    curr_sub_cmd_idx,
+                    false,
+                    DISPATCH_WRITE_OFFSET_TENSIX_L1_CONFIG_BASE);
                 for (auto &data_and_size : multicast_cb_config_data) {
                     RecordDispatchData(program, DISPATCH_DATA_CB_CONFIG, data_and_size.second);
                 }