#0: set data movement kernels to build at O2

tenstorrent · Feb 25, 2025 · d15ef7e · d15ef7e
1 parent 663244a
commit d15ef7e
Show file tree

Hide file tree

Showing 4 changed files with 9 additions and 4 deletions.
diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/common.h b/tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/common.h
@@ -563,7 +563,8 @@ void configure_kernel_variant(
             .processor = tt::tt_metal::DataMovementProcessor::RISCV_1,
             .noc = my_noc_index,
             .compile_args = compile_args,
-            .defines = defines});
+            .defines = defines,
+            .opt_level = KernelBuildOptLevel::Os});
 }
 
 // Specific to this test. This test doesn't use Buffers, and for Storage cores in L1 that have 2 banks, they are

diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/test_dispatcher.cpp b/tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/test_dispatcher.cpp
@@ -6,6 +6,7 @@
 #include <functional>
 #include <random>
 
+#include "kernel_types.hpp"
 #include "logger.hpp"
 #include <tt-metalium/host_api.hpp>
 #include <tt-metalium/tt_align.hpp>

diff --git a/tt_metal/api/tt-metalium/kernel_types.hpp b/tt_metal/api/tt-metalium/kernel_types.hpp
@@ -36,7 +36,7 @@ struct DataMovementConfig {
     // This file is then automatically included in the generated compiled kernel files
     std::map<std::string, std::string> defines;
     // Set the compiler and linker optimization level
-    KernelBuildOptLevel opt_level = KernelBuildOptLevel::Os;
+    KernelBuildOptLevel opt_level = KernelBuildOptLevel::O2;
 };
 
 struct ReaderDataMovementConfig : public DataMovementConfig {

diff --git a/tt_metal/impl/dispatch/kernel_config/fd_kernel.cpp b/tt_metal/impl/dispatch/kernel_config/fd_kernel.cpp
@@ -7,6 +7,7 @@
 #include <tt_metal.hpp>
 #include "dprint_server.hpp"
 
+#include "kernel_types.hpp"
 #include "prefetch.hpp"
 #include "dispatch.hpp"
 #include "dispatch_s.hpp"
@@ -134,7 +135,8 @@ void FDKernel::configure_kernel_variant(
                                            : tt::tt_metal::DataMovementProcessor::RISCV_1,
                 .noc = noc_selection_.non_dispatch_noc,
                 .compile_args = compile_args,
-                .defines = defines});
+                .defines = defines,
+                .opt_level = KernelBuildOptLevel::Os});
     } else {
         tt::tt_metal::CreateKernel(
             *program_,
@@ -144,6 +146,7 @@ void FDKernel::configure_kernel_variant(
                 .eth_mode = is_active_eth_core ? Eth::SENDER : Eth::IDLE,
                 .noc = noc_selection_.non_dispatch_noc,
                 .compile_args = compile_args,
-                .defines = defines});
+                .defines = defines,
+                .opt_level = KernelBuildOptLevel::Os});
     }
 }