Skip to content

Commit

Permalink
More work
Browse files Browse the repository at this point in the history
  • Loading branch information
@sagarwalTT
sagarwalTT committed Nov 14, 2024
1 parent a662843 commit 953bd0b
Show file tree
Hide file tree
Showing 7 changed files with 154 additions and 275 deletions.
1 change: 0 additions & 1 deletion tests/tt_metal/tt_metal/common/command_queue_fixture.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -8,7 +8,6 @@
#include "common/env_lib.hpp"
#include "gtest/gtest.h"
#include "dispatch_fixture.hpp"
#include "trace_fixture.hpp"
#include "hostdevcommon/common_runtime_address_map.h"
#include "hostdevcommon/common_values.hpp"
#include "impl/buffers/circular_buffer_types.hpp"
Expand Down
32 changes: 32 additions & 0 deletions tests/tt_metal/tt_metal/common/device_fixture.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -85,6 +85,38 @@ class DeviceSingleCardFixture : public ::testing::Test {

class DeviceSingleCardBufferFixture : public DeviceSingleCardFixture {};

class SingleDeviceTraceFixture : public ::testing::Test {
protected:
Device* device_;
tt::ARCH arch_;

void Setup(const size_t buffer_size, const uint8_t num_hw_cqs = 1) {
auto slow_dispatch = getenv("TT_METAL_SLOW_DISPATCH_MODE");
if (slow_dispatch) {
tt::log_info(
tt::LogTest, "This suite can only be run with fast dispatch or TT_METAL_SLOW_DISPATCH_MODE unset");
GTEST_SKIP();
}
if (num_hw_cqs > 1) {
// Running multi-CQ test. User must set this explicitly.
auto num_cqs = getenv("TT_METAL_GTEST_NUM_HW_CQS");
if (num_cqs == nullptr or strcmp(num_cqs, "2")) {
TT_THROW("This suite must be run with TT_METAL_GTEST_NUM_HW_CQS=2");
GTEST_SKIP();
}
}
this->arch_ = tt::get_arch_from_string(tt::test_utils::get_umd_arch_name());
const int device_id = 0;
this->device_ = tt::tt_metal::CreateDevice(device_id, num_hw_cqs, 0, buffer_size);
}

void TearDown() override {
if (!getenv("TT_METAL_SLOW_DISPATCH_MODE")) {
tt::tt_metal::CloseDevice(this->device_);
}
}
};

class BlackholeSingleCardFixture : public DeviceSingleCardFixture {
protected:
void SetUp() override {
Expand Down
42 changes: 0 additions & 42 deletions tests/tt_metal/tt_metal/common/trace_fixture.hpp
View file

This file was deleted.

4 changes: 0 additions & 4 deletions tests/tt_metal/tt_metal/dispatch/dispatch_program/CMakeLists.txt
View file
Original file line number Diff line number Diff line change
@@ -1,13 +1,9 @@
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
<<<<<<< HEAD:tests/tt_metal/tt_metal/dispatch_program/CMakeLists.txt
${CMAKE_CURRENT_SOURCE_DIR}/test_sub_device.cpp
=======
CACHE INTERNAL
""
>>>>>>> #0: More test structuring:tests/tt_metal/tt_metal/dispatch/dispatch_program/CMakeLists.txt
)

add_executable(unit_tests_dispatch_program ${UNIT_TESTS_DISPATCH_PROGRAM_SRC})
Expand Down
116 changes: 2 additions & 114 deletions ...etal/dispatch_program/test_sub_device.cpp → ...atch/dispatch_program/test_sub_device.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -17,119 +17,7 @@
#include "test_utils/stimulus.hpp"
#include "command_queue_fixture.hpp"
#include "command_queue_test_utils.hpp"

std::tuple<Program, CoreCoord, std::unique_ptr<GlobalSemaphore>> create_single_sync_program(Device *device, SubDevice sub_device) {
auto syncer_coord = sub_device.cores(HalProgrammableCoreType::TENSIX).ranges().at(0).start_coord;
auto syncer_core = CoreRangeSet(CoreRange(syncer_coord, syncer_coord));
auto global_sem = CreateGlobalSemaphore(device, sub_device.cores(HalProgrammableCoreType::TENSIX), INVALID);

Program syncer_program = CreateProgram();
auto syncer_kernel = CreateKernel(
syncer_program,
"tests/tt_metal/tt_metal/test_kernels/misc/sub_device/syncer.cpp",
syncer_core,
DataMovementConfig{
.processor = DataMovementProcessor::RISCV_0,
.noc = NOC::RISCV_0_default});
std::array<uint32_t, 1> syncer_rt_args = {global_sem->address()};
SetRuntimeArgs(syncer_program, syncer_kernel, syncer_core, syncer_rt_args);
return {std::move(syncer_program), std::move(syncer_coord), std::move(global_sem)};
}

std::tuple<Program, Program, Program, std::unique_ptr<GlobalSemaphore>> create_basic_sync_program(Device *device, const SubDevice& sub_device_1, const SubDevice& sub_device_2) {
auto waiter_coord = sub_device_2.cores(HalProgrammableCoreType::TENSIX).ranges().at(0).start_coord;
auto waiter_core = CoreRangeSet(CoreRange(waiter_coord, waiter_coord));
auto waiter_core_physical = device->worker_core_from_logical_core(waiter_coord);
auto incrementer_cores = sub_device_1.cores(HalProgrammableCoreType::TENSIX);
auto syncer_coord = incrementer_cores.ranges().back().end_coord;
auto syncer_core = CoreRangeSet(CoreRange(syncer_coord, syncer_coord));
auto syncer_core_physical = device->worker_core_from_logical_core(syncer_coord);
auto all_cores = waiter_core.merge(incrementer_cores).merge(syncer_core);
auto global_sem = CreateGlobalSemaphore(device, all_cores, INVALID);

Program waiter_program = CreateProgram();
auto waiter_kernel = CreateKernel(
waiter_program,
"tests/tt_metal/tt_metal/test_kernels/misc/sub_device/persistent_waiter.cpp",
waiter_core,
DataMovementConfig{
.processor = DataMovementProcessor::RISCV_0,
.noc = NOC::RISCV_0_default});
std::array<uint32_t, 4> waiter_rt_args = {global_sem->address(), incrementer_cores.num_cores(), syncer_core_physical.x, syncer_core_physical.y};
SetRuntimeArgs(waiter_program, waiter_kernel, waiter_core, waiter_rt_args);

Program syncer_program = CreateProgram();
auto syncer_kernel = CreateKernel(
syncer_program,
"tests/tt_metal/tt_metal/test_kernels/misc/sub_device/syncer.cpp",
syncer_core,
DataMovementConfig{
.processor = DataMovementProcessor::RISCV_0,
.noc = NOC::RISCV_0_default});
std::array<uint32_t, 1> syncer_rt_args = {global_sem->address()};
SetRuntimeArgs(syncer_program, syncer_kernel, syncer_core, syncer_rt_args);

Program incrementer_program = CreateProgram();
auto incrementer_kernel = CreateKernel(
incrementer_program,
"tests/tt_metal/tt_metal/test_kernels/misc/sub_device/incrementer.cpp",
incrementer_cores,
DataMovementConfig{
.processor = DataMovementProcessor::RISCV_1,
.noc = NOC::RISCV_1_default});
std::array<uint32_t, 3> incrementer_rt_args = {global_sem->address(), waiter_core_physical.x, waiter_core_physical.y};
SetRuntimeArgs(incrementer_program, incrementer_kernel, incrementer_cores, incrementer_rt_args);
return {std::move(waiter_program), std::move(syncer_program), std::move(incrementer_program), std::move(global_sem)};
}

std::tuple<Program, Program, Program, std::unique_ptr<GlobalSemaphore>> create_basic_eth_sync_program(Device *device, const SubDevice& sub_device_1, const SubDevice& sub_device_2) {
auto waiter_coord = sub_device_2.cores(HalProgrammableCoreType::ACTIVE_ETH).ranges().at(0).start_coord;
auto waiter_core = CoreRangeSet(CoreRange(waiter_coord, waiter_coord));
auto waiter_core_physical = device->ethernet_core_from_logical_core(waiter_coord);
auto tensix_waiter_coord = sub_device_2.cores(HalProgrammableCoreType::TENSIX).ranges().at(0).start_coord;
auto tensix_waiter_core = CoreRangeSet(CoreRange(tensix_waiter_coord, tensix_waiter_coord));
auto tensix_waiter_core_physical = device->worker_core_from_logical_core(tensix_waiter_coord);
auto incrementer_cores = sub_device_1.cores(HalProgrammableCoreType::TENSIX);
auto syncer_coord = incrementer_cores.ranges().back().end_coord;
auto syncer_core = CoreRangeSet(CoreRange(syncer_coord, syncer_coord));
auto syncer_core_physical = device->worker_core_from_logical_core(syncer_coord);
auto all_cores = tensix_waiter_core.merge(incrementer_cores).merge(syncer_core);
auto global_sem = CreateGlobalSemaphore(device, all_cores, INVALID);

Program waiter_program = CreateProgram();
auto waiter_kernel = CreateKernel(
waiter_program,
"tests/tt_metal/tt_metal/test_kernels/misc/sub_device/persistent_remote_waiter.cpp",
waiter_core,
EthernetConfig{
.noc = NOC::RISCV_0_default,
.processor = DataMovementProcessor::RISCV_0});
std::array<uint32_t, 7> waiter_rt_args = {global_sem->address(), incrementer_cores.num_cores(), syncer_core_physical.x, syncer_core_physical.y, tensix_waiter_core_physical.x, tensix_waiter_core_physical.y, eth_l1_mem::address_map::ERISC_L1_UNRESERVED_BASE};
SetRuntimeArgs(waiter_program, waiter_kernel, waiter_core, waiter_rt_args);

Program syncer_program = CreateProgram();
auto syncer_kernel = CreateKernel(
syncer_program,
"tests/tt_metal/tt_metal/test_kernels/misc/sub_device/syncer.cpp",
syncer_core,
DataMovementConfig{
.processor = DataMovementProcessor::RISCV_0,
.noc = NOC::RISCV_0_default});
std::array<uint32_t, 1> syncer_rt_args = {global_sem->address()};
SetRuntimeArgs(syncer_program, syncer_kernel, syncer_core, syncer_rt_args);

Program incrementer_program = CreateProgram();
auto incrementer_kernel = CreateKernel(
incrementer_program,
"tests/tt_metal/tt_metal/test_kernels/misc/sub_device/incrementer.cpp",
incrementer_cores,
DataMovementConfig{
.processor = DataMovementProcessor::RISCV_1,
.noc = NOC::RISCV_1_default});
std::array<uint32_t, 3> incrementer_rt_args = {global_sem->address(), tensix_waiter_core_physical.x, tensix_waiter_core_physical.y};
SetRuntimeArgs(incrementer_program, incrementer_kernel, incrementer_cores, incrementer_rt_args);
return {std::move(waiter_program), std::move(syncer_program), std::move(incrementer_program), std::move(global_sem)};
}
#include "sub_device_test_utils.hpp"

TEST_F(CommandQueueSingleCardFixture, TensixTestSubDeviceAllocations) {
uint32_t local_l1_size = 3200;
Expand Down Expand Up @@ -326,4 +214,4 @@ TEST_F(CommandQueueSingleCardFixture, TensixActiveEthTestSubDeviceBasicEthProgra
}
Synchronize(device);
}
}
}
116 changes: 2 additions & 114 deletions ..._metal/dispatch_trace/test_sub_device.cpp → ...spatch/dispatch_trace/test_sub_device.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -17,119 +17,7 @@
#include "test_utils/stimulus.hpp"
#include "command_queue_fixture.hpp"
#include "command_queue_test_utils.hpp"

std::tuple<Program, CoreCoord, std::unique_ptr<GlobalSemaphore>> create_single_sync_program(Device *device, SubDevice sub_device) {
auto syncer_coord = sub_device.cores(HalProgrammableCoreType::TENSIX).ranges().at(0).start_coord;
auto syncer_core = CoreRangeSet(CoreRange(syncer_coord, syncer_coord));
auto global_sem = CreateGlobalSemaphore(device, sub_device.cores(HalProgrammableCoreType::TENSIX), INVALID);

Program syncer_program = CreateProgram();
auto syncer_kernel = CreateKernel(
syncer_program,
"tests/tt_metal/tt_metal/test_kernels/misc/sub_device/syncer.cpp",
syncer_core,
DataMovementConfig{
.processor = DataMovementProcessor::RISCV_0,
.noc = NOC::RISCV_0_default});
std::array<uint32_t, 1> syncer_rt_args = {global_sem->address()};
SetRuntimeArgs(syncer_program, syncer_kernel, syncer_core, syncer_rt_args);
return {std::move(syncer_program), std::move(syncer_coord), std::move(global_sem)};
}

std::tuple<Program, Program, Program, std::unique_ptr<GlobalSemaphore>> create_basic_sync_program(Device *device, const SubDevice& sub_device_1, const SubDevice& sub_device_2) {
auto waiter_coord = sub_device_2.cores(HalProgrammableCoreType::TENSIX).ranges().at(0).start_coord;
auto waiter_core = CoreRangeSet(CoreRange(waiter_coord, waiter_coord));
auto waiter_core_physical = device->worker_core_from_logical_core(waiter_coord);
auto incrementer_cores = sub_device_1.cores(HalProgrammableCoreType::TENSIX);
auto syncer_coord = incrementer_cores.ranges().back().end_coord;
auto syncer_core = CoreRangeSet(CoreRange(syncer_coord, syncer_coord));
auto syncer_core_physical = device->worker_core_from_logical_core(syncer_coord);
auto all_cores = waiter_core.merge(incrementer_cores).merge(syncer_core);
auto global_sem = CreateGlobalSemaphore(device, all_cores, INVALID);

Program waiter_program = CreateProgram();
auto waiter_kernel = CreateKernel(
waiter_program,
"tests/tt_metal/tt_metal/test_kernels/misc/sub_device/persistent_waiter.cpp",
waiter_core,
DataMovementConfig{
.processor = DataMovementProcessor::RISCV_0,
.noc = NOC::RISCV_0_default});
std::array<uint32_t, 4> waiter_rt_args = {global_sem->address(), incrementer_cores.num_cores(), syncer_core_physical.x, syncer_core_physical.y};
SetRuntimeArgs(waiter_program, waiter_kernel, waiter_core, waiter_rt_args);

Program syncer_program = CreateProgram();
auto syncer_kernel = CreateKernel(
syncer_program,
"tests/tt_metal/tt_metal/test_kernels/misc/sub_device/syncer.cpp",
syncer_core,
DataMovementConfig{
.processor = DataMovementProcessor::RISCV_0,
.noc = NOC::RISCV_0_default});
std::array<uint32_t, 1> syncer_rt_args = {global_sem->address()};
SetRuntimeArgs(syncer_program, syncer_kernel, syncer_core, syncer_rt_args);

Program incrementer_program = CreateProgram();
auto incrementer_kernel = CreateKernel(
incrementer_program,
"tests/tt_metal/tt_metal/test_kernels/misc/sub_device/incrementer.cpp",
incrementer_cores,
DataMovementConfig{
.processor = DataMovementProcessor::RISCV_1,
.noc = NOC::RISCV_1_default});
std::array<uint32_t, 3> incrementer_rt_args = {global_sem->address(), waiter_core_physical.x, waiter_core_physical.y};
SetRuntimeArgs(incrementer_program, incrementer_kernel, incrementer_cores, incrementer_rt_args);
return {std::move(waiter_program), std::move(syncer_program), std::move(incrementer_program), std::move(global_sem)};
}

std::tuple<Program, Program, Program, std::unique_ptr<GlobalSemaphore>> create_basic_eth_sync_program(Device *device, const SubDevice& sub_device_1, const SubDevice& sub_device_2) {
auto waiter_coord = sub_device_2.cores(HalProgrammableCoreType::ACTIVE_ETH).ranges().at(0).start_coord;
auto waiter_core = CoreRangeSet(CoreRange(waiter_coord, waiter_coord));
auto waiter_core_physical = device->ethernet_core_from_logical_core(waiter_coord);
auto tensix_waiter_coord = sub_device_2.cores(HalProgrammableCoreType::TENSIX).ranges().at(0).start_coord;
auto tensix_waiter_core = CoreRangeSet(CoreRange(tensix_waiter_coord, tensix_waiter_coord));
auto tensix_waiter_core_physical = device->worker_core_from_logical_core(tensix_waiter_coord);
auto incrementer_cores = sub_device_1.cores(HalProgrammableCoreType::TENSIX);
auto syncer_coord = incrementer_cores.ranges().back().end_coord;
auto syncer_core = CoreRangeSet(CoreRange(syncer_coord, syncer_coord));
auto syncer_core_physical = device->worker_core_from_logical_core(syncer_coord);
auto all_cores = tensix_waiter_core.merge(incrementer_cores).merge(syncer_core);
auto global_sem = CreateGlobalSemaphore(device, all_cores, INVALID);

Program waiter_program = CreateProgram();
auto waiter_kernel = CreateKernel(
waiter_program,
"tests/tt_metal/tt_metal/test_kernels/misc/sub_device/persistent_remote_waiter.cpp",
waiter_core,
EthernetConfig{
.noc = NOC::RISCV_0_default,
.processor = DataMovementProcessor::RISCV_0});
std::array<uint32_t, 7> waiter_rt_args = {global_sem->address(), incrementer_cores.num_cores(), syncer_core_physical.x, syncer_core_physical.y, tensix_waiter_core_physical.x, tensix_waiter_core_physical.y, eth_l1_mem::address_map::ERISC_L1_UNRESERVED_BASE};
SetRuntimeArgs(waiter_program, waiter_kernel, waiter_core, waiter_rt_args);

Program syncer_program = CreateProgram();
auto syncer_kernel = CreateKernel(
syncer_program,
"tests/tt_metal/tt_metal/test_kernels/misc/sub_device/syncer.cpp",
syncer_core,
DataMovementConfig{
.processor = DataMovementProcessor::RISCV_0,
.noc = NOC::RISCV_0_default});
std::array<uint32_t, 1> syncer_rt_args = {global_sem->address()};
SetRuntimeArgs(syncer_program, syncer_kernel, syncer_core, syncer_rt_args);

Program incrementer_program = CreateProgram();
auto incrementer_kernel = CreateKernel(
incrementer_program,
"tests/tt_metal/tt_metal/test_kernels/misc/sub_device/incrementer.cpp",
incrementer_cores,
DataMovementConfig{
.processor = DataMovementProcessor::RISCV_1,
.noc = NOC::RISCV_1_default});
std::array<uint32_t, 3> incrementer_rt_args = {global_sem->address(), tensix_waiter_core_physical.x, tensix_waiter_core_physical.y};
SetRuntimeArgs(incrementer_program, incrementer_kernel, incrementer_cores, incrementer_rt_args);
return {std::move(waiter_program), std::move(syncer_program), std::move(incrementer_program), std::move(global_sem)};
}
#include "sub_device_test_utils.hpp"

TEST_F(CommandQueueSingleCardTraceFixture, TensixTestSubDeviceTraceBasicPrograms) {
SubDevice sub_device_1(std::array{CoreRangeSet(CoreRange({0, 0}, {2, 2}))});
Expand Down Expand Up @@ -379,4 +267,4 @@ TEST_F(CommandQueueSingleCardTraceFixture, TensixTestSubDeviceIllegalOperations)
device->remove_sub_device_manager(sub_device_manager_1);
EXPECT_THROW(device->load_sub_device_manager(sub_device_manager_1), std::exception);
}
}
}
Loading

0 comments on commit 953bd0b

Please sign in to comment.