#0: More work

tests/tt_metal/tt_metal/api/test_kernel_creation.cpp

@@ -74,3 +74,32 @@ TEST_F(DispatchFixture, TensixIdleEthCreateKernelsOnDispatchCores) {
         }
     }
 }
+
+TEST_F(CompileProgramWithKernelPathEnvVarFixture, TensixKernelUnderMetalRootDir) {
+    const string &kernel_file = "tests/tt_metal/tt_metal/test_kernels/dataflow/reader_unary_push_4.cpp";
+    create_kernel(kernel_file);
+    detail::CompileProgram(this->device_, this->program_);
+}
+
+TEST_F(CompileProgramWithKernelPathEnvVarFixture, TensixKernelUnderKernelRootDir) {
+    const string &orig_kernel_file = "tests/tt_metal/tt_metal/test_kernels/dataflow/reader_unary_push_4.cpp";
+    const string &new_kernel_file = "tests/tt_metal/tt_metal/test_kernels/dataflow/new_kernel.cpp";
+    this->setup_kernel_dir(orig_kernel_file, new_kernel_file);
+    this->create_kernel(new_kernel_file);
+    detail::CompileProgram(this->device_, this->program_);
+    this->cleanup_kernel_dir();
+}
+
+TEST_F(CompileProgramWithKernelPathEnvVarFixture, TensixKernelUnderMetalRootDirAndKernelRootDir) {
+    const string &kernel_file = "tests/tt_metal/tt_metal/test_kernels/dataflow/reader_unary_push_4.cpp";
+    this->setup_kernel_dir(kernel_file, kernel_file);
+    this->create_kernel(kernel_file);
+    detail::CompileProgram(this->device_, this->program_);
+    this->cleanup_kernel_dir();
+}
+
+TEST_F(CompileProgramWithKernelPathEnvVarFixture, TensixNonExistentKernel) {
+    const string &kernel_file = "tests/tt_metal/tt_metal/test_kernels/dataflow/non_existent_kernel.cpp";
+    this->create_kernel(kernel_file);
+    EXPECT_THROW(detail::CompileProgram(this->device_, this->program_), std::exception);
+}

tests/tt_metal/tt_metal/dispatch/dispatch_program/CMakeLists.txt

@@ -1,4 +1,5 @@
 set(UNIT_TESTS_DISPATCH_PROGRAM_SRC
+    ${CMAKE_CURRENT_SOURCE_DIR}/test_dispatch_program_with_kernel_created_from_string.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/test_dispatch.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/test_EnqueueProgram.cpp
     CACHE INTERNAL

tests/tt_metal/tt_metal/dispatch/dispatch_program/test_dispatch_program_with_kernel_created_from_string.cpp

@@ -0,0 +1,102 @@
+// SPDX-FileCopyrightText: © 2024 Tenstorrent Inc.
+//
+// SPDX-License-Identifier: Apache-2.0
+
+#include <gtest/gtest.h>
+
+#include "core_coord.hpp"
+#include "detail/tt_metal.hpp"
+#include "host_api.hpp"
+#include "impl/device/device.hpp"
+#include "impl/kernels/data_types.hpp"
+#include "impl/kernels/kernel_types.hpp"
+#include "impl/program/program.hpp"
+#include "tt_cluster_descriptor_types.h"
+
+#include "tests/tt_metal/tt_metal/common/dispatch_fixture.hpp"
+
+TEST_F(ProgramWithKernelCreatedFromStringFixture, TensixDataMovementKernel) {
+    const CoreRange cores({0, 0}, {1, 1});
+    const string &kernel_src_code = R"(
+    #include "debug/dprint.h"
+    #include "dataflow_api.h"
+
+    void kernel_main() {
+
+        DPRINT_DATA0(DPRINT << "Hello, I am running a void data movement kernel on NOC 0." << ENDL());
+        DPRINT_DATA1(DPRINT << "Hello, I am running a void data movement kernel on NOC 1." << ENDL());
+
+    }
+    )";
+
+    for (Device *device : this->devices_) {
+        Program program = CreateProgram();
+        tt_metal::CreateKernelFromString(
+            program,
+            kernel_src_code,
+            cores,
+            tt_metal::DataMovementConfig{.processor = DataMovementProcessor::RISCV_1, .noc = NOC::RISCV_1_default});
+        this->RunProgram(device, program);
+    };
+}
+
+TEST_F(ProgramWithKernelCreatedFromStringFixture, TensixComputeKernel) {
+    const CoreRange cores({0, 0}, {1, 1});
+    const string &kernel_src_code = R"(
+    #include "debug/dprint.h"
+    #include "compute_kernel_api.h"
+
+    namespace NAMESPACE {
+
+    void MAIN {
+
+        DPRINT_MATH(DPRINT << "Hello, I am running a void compute kernel." << ENDL());
+
+    }
+
+    }
+    )";
+
+    for (Device *device : this->devices_) {
+        Program program = CreateProgram();
+        tt_metal::CreateKernelFromString(
+            program,
+            kernel_src_code,
+            cores,
+            tt_metal::ComputeConfig{
+                .math_fidelity = MathFidelity::HiFi4,
+                .fp32_dest_acc_en = false,
+                .math_approx_mode = false,
+                .compile_args = {}});
+        this->RunProgram(device, program);
+    };
+}
+
+TEST_F(ProgramWithKernelCreatedFromStringFixture, ActiveEthEthernetKernel) {
+    const string &kernel_src_code = R"(
+    #include "debug/dprint.h"
+    #include "dataflow_api.h"
+
+    void kernel_main() {
+
+        DPRINT << "Hello, I am running a void ethernet kernel." << ENDL();
+
+    }
+    )";
+
+    for (Device *device : this->devices_) {
+        const std::unordered_set<CoreCoord> &active_ethernet_cores = device->get_active_ethernet_cores(true);
+        if (active_ethernet_cores.empty()) {
+            const chip_id_t device_id = device->id();
+            log_info(LogTest, "Skipping this test on device {} because it has no active ethernet cores.", device_id);
+            continue;
+        }
+        Program program = CreateProgram();
+        tt_metal::CreateKernelFromString(
+            program,
+            kernel_src_code,
+            *active_ethernet_cores.begin(),
+            tt_metal::EthernetConfig{.noc = tt_metal::NOC::NOC_0});
+        this->RunProgram(device, program);
+    };
+}