diff --git a/.github/workflows/build-tenstorrent.yml b/.github/workflows/build-tenstorrent.yml new file mode 100644 index 00000000..7bf1a2d3 --- /dev/null +++ b/.github/workflows/build-tenstorrent.yml @@ -0,0 +1,15 @@ +name: Build and Test Tenstorrent Backend + +on: + pull_request: + branches: [main, spatter-devel] + schedule: + - cron: '30 8 * * *' + +jobs: + build-tenstorrent: + runs-on: self-hosted + steps: + - uses: actions/checkout@v4 + - name: Run batch file + run: cd tests/misc && chmod +x run-crnch-tenstorrent.sh && sbatch run-crnch-tenstorrent.sh diff --git a/.gitignore b/.gitignore index f1ab5ffc..cf890055 100644 --- a/.gitignore +++ b/.gitignore @@ -9,3 +9,4 @@ src/python/env *.pyc spatter-cuda-test.out json.tar.xz +generated/ diff --git a/AUTHORS b/AUTHORS index c67dcbf4..ece1d0fd 100644 --- a/AUTHORS +++ b/AUTHORS @@ -20,3 +20,6 @@ Vincent Huang James Wood OneAPI backend support + +Roy Cheng + Tenstorrent backend diff --git a/Build.md b/Build.md index d65fdc26..34c2b8a7 100644 --- a/Build.md +++ b/Build.md @@ -28,4 +28,54 @@ * avx_crossplatform * non_avx * `-DUSE_MPI=1` - * `-DUSE_PAPI=1` \ No newline at end of file + * `-DUSE_PAPI=1` + +## Tenstorrent +Blackhole and other Tenstorrent accelerators, via tt-metal. Enable with +`-DUSE_TENSTORRENT=ON`. + +Device kernels are compiled at run time by tt-metal's vendored sfpi toolchain, +so no extra device compiler is needed at build time. The host side needs a C++20 +compiler and `libtt_metal`. + +### Optional arguments +* `-DTT_METAL_LIB_DIR=` - directory holding `libtt_metal.so` +* `-DTT_METAL_INCLUDE_DIRS=` - include roots for the host headers + +### Runtime environment +* `TT_METAL_RUNTIME_ROOT` must point at the directory containing `tt_metal/`, + otherwise tt-metal aborts with "Root Directory is not set" before it opens a + device. +* `TT_VISIBLE_DEVICES` selects which chip(s) to use on a multi-card host. + +### If you installed Tenstorrent support with a pip `ttnn` wheel +The wheel ships `libtt_metal.so` and the device-kernel headers, but **not** the +host API headers, so CMake will find the library and then report the headers +missing. They can be assembled without root and without building tt-metal: + +``` +git clone --depth 1 --branch --filter=blob:none --sparse \ + https://github.com/tenstorrent/tt-metal ~/ttmetal-src +cd ~/ttmetal-src && git sparse-checkout set tt_metal tt_stl +git submodule update --init --depth 1 tt_metal/third_party/umd +``` + +That covers `tt-metalium`, `tt_stl`, `hostdevcommon` and `umd`. The remaining +four are header-only and are fetched by tt-metal's build rather than vendored, +so clone them at the versions pinned in `~/ttmetal-src/third_party/CMakeLists.txt`: +fmt, nlohmann/json, tt-logger and spdlog, plus enchantum. Pass every include +root via `-DTT_METAL_INCLUDE_DIRS`. + +Do **not** run tt-metal's own CMake configure just to obtain these: it pulls +system packages (boost, capnproto, protobuf) that require root, whereas the +header-only clones do not. + +### Notes +* Blackhole has no FP64 fabric. This does not affect Spatter, because gather and + scatter perform no arithmetic; each element is moved as an opaque 8-byte + payload. +* The DRAM allocator aligns pages to 64 bytes, so a buffer with one 8-byte + element per page occupies 8x its logical size on device. Size `-l` accordingly. +* `gather` and `scatter` are implemented. The `multi_*` and atomic variants are + not yet, and report so at run time. + diff --git a/CMakeLists.txt b/CMakeLists.txt index 606d85e1..a235f22a 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -19,6 +19,7 @@ include(pkgs/JSONSupport) include(pkgs/MPISupport) include(pkgs/OpenMPSupport) include(pkgs/CUDASupport) +include(pkgs/TenstorrentSupport) if (APPLE) set(CMAKE_INSTALL_RPATH "@executable_path/../lib") diff --git a/README.md b/README.md index cfb944e2..455e1ad6 100644 --- a/README.md +++ b/README.md @@ -44,7 +44,7 @@ CMake is required to build Spatter. Currently we require CMake 3.25 or newer. To build with CMake from the main source directory, use the following command structure: ``` -cmake -DCMAKE_BUILD_TYPE= -DUSE_=1 -B build_ -S . +cmake -DCMAKE_BUILD_TYPE= -DUSE_=1 -B build_ -S . cd build_ make ``` @@ -63,6 +63,12 @@ For CUDA builds, we normally load CUDA 11/12 using NVHPC: ``` cmake -DUSE_CUDA=1 -B build_cuda -S . ``` + +For Tenstorrent builds (see `Build.md` for the header requirements): + +``` +cmake -DUSE_TENSTORRENT=ON -B build_tenstorrent -S . +``` For a complete list of build options, see [Build.md](Build.md) ## Running Spatter diff --git a/cmake/pkgs/TenstorrentSupport.cmake b/cmake/pkgs/TenstorrentSupport.cmake new file mode 100644 index 00000000..5455a88e --- /dev/null +++ b/cmake/pkgs/TenstorrentSupport.cmake @@ -0,0 +1,79 @@ +# Tenstorrent device kernels are compiled at run time by tt-metal's vendored +# sfpi toolchain, so there is no device compiler to enable here; only the host +# library matters at build time. There is no upstream FindTTMetal module. + +option(USE_TENSTORRENT "Enable support for Tenstorrent accelerators") + +if (USE_TENSTORRENT) + set(TT_METAL_INCLUDE_DIRS "" CACHE STRING + "Semicolon-separated include roots for tt-metal host headers") + set(TT_METAL_LIB_DIR "" CACHE PATH "Directory containing libtt_metal.so") + + find_package(Python3 COMPONENTS Interpreter QUIET) + if (Python3_Interpreter_FOUND) + execute_process( + COMMAND ${Python3_EXECUTABLE} -c + "import ttnn, os; print(os.path.dirname(ttnn.__file__))" + OUTPUT_VARIABLE TTNN_DIR + OUTPUT_STRIP_TRAILING_WHITESPACE + ERROR_QUIET) + if (TTNN_DIR) + message(STATUS "Tenstorrent: found ttnn wheel at ${TTNN_DIR}") + endif() + endif() + + find_library(TT_METAL_LIB + NAMES tt_metal + HINTS ${TT_METAL_LIB_DIR} + ${TTNN_DIR}/build/lib + $ENV{TT_METAL_HOME}/build/lib + $ENV{TT_METAL_RUNTIME_ROOT}/build/lib) + + # Headers include each other as , so the root is the parent + # of tt-metalium/. + find_path(TT_METALIUM_INCLUDE_DIR + NAMES tt-metalium/host_api.hpp + HINTS ${TT_METAL_INCLUDE_DIRS} + $ENV{TT_METAL_HOME}/tt_metal/api + ${TTNN_DIR}/tt_metal/api) + + if (TT_METAL_LIB AND TT_METALIUM_INCLUDE_DIR) + message(STATUS "Found libtt_metal: ${TT_METAL_LIB}") + message(STATUS "Found tt-metalium headers: ${TT_METALIUM_INCLUDE_DIR}") + + set(CMAKE_CXX_STANDARD 20) + set(CMAKE_CXX_STANDARD_REQUIRED ON) + + include_directories(${TT_METALIUM_INCLUDE_DIR}) + if (TT_METAL_INCLUDE_DIRS) + include_directories(${TT_METAL_INCLUDE_DIRS}) + endif() + + set(COMMON_LINK_LIBRARIES ${COMMON_LINK_LIBRARIES} ${TT_METAL_LIB}) + add_definitions(-DUSE_TENSTORRENT) + else() + if (TT_METAL_LIB AND NOT TT_METALIUM_INCLUDE_DIR) + message(STATUS + "Tenstorrent: libtt_metal was found but the host headers were not. " + "A pip `ttnn` wheel ships the library and the device-kernel headers " + "but no host API headers. Point -DTT_METAL_INCLUDE_DIRS at a matching " + "tt-metal checkout plus its header-only dependencies, e.g.\n" + " git clone --depth 1 --branch v0.72.0 --filter=blob:none --sparse \\\n" + " https://github.com/tenstorrent/tt-metal ~/ttmetal-src\n" + " cd ~/ttmetal-src && git sparse-checkout set tt_metal tt_stl\n" + " git submodule update --init --depth 1 tt_metal/third_party/umd\n" + "then clone fmt, nlohmann/json, tt-logger, spdlog and enchantum at the " + "versions pinned in that checkout's third_party/CMakeLists.txt and pass " + "every include root. Do NOT run tt-metal's own configure for this: it " + "pulls system packages (boost, capnproto, protobuf) that need root, " + "whereas the header-only clones do not.") + elseif (NOT TT_METAL_LIB) + message(STATUS + "Tenstorrent: libtt_metal not found. Install tt-metal, or pip install " + "ttnn, or set -DTT_METAL_LIB_DIR.") + endif() + message(FATAL_ERROR + "USE_TENSTORRENT=ON but no usable Tenstorrent installation was found. " + "See the diagnostic above, or configure without -DUSE_TENSTORRENT=ON.") + endif() +endif() diff --git a/src/Spatter/CMakeLists.txt b/src/Spatter/CMakeLists.txt index 84329bee..1789398d 100644 --- a/src/Spatter/CMakeLists.txt +++ b/src/Spatter/CMakeLists.txt @@ -6,8 +6,19 @@ if (USE_CUDA) set(CUDA_INCLUDE_FILES CudaBackend.hh) endif() +if (USE_TENSTORRENT) + add_library(tenstorrent_backend SHARED TenstorrentBackend.cc) + target_link_libraries(tenstorrent_backend PUBLIC ${TT_METAL_LIB}) + target_compile_features(tenstorrent_backend PUBLIC cxx_std_20) + # tt-metal's headers pull in fmt and spdlog; use them header-only so the + # backend does not need those libraries at link time. + target_compile_definitions(tenstorrent_backend PRIVATE FMT_HEADER_ONLY) + set(TENSTORRENT_INCLUDE_FILES TenstorrentBackend.hh) +endif() + set(SPATTER_INCLUDE_FILES ${CUDA_INCLUDE_FILES} + ${TENSTORRENT_INCLUDE_FILES} Configuration.hh Input.hh JSONParser.hh @@ -61,6 +72,10 @@ if (USE_CUDA) set(COMMON_LINK_LIBRARIES ${COMMON_LINK_LIBRARIES} cuda_backend) endif() +if (USE_TENSTORRENT) + set(COMMON_LINK_LIBRARIES ${COMMON_LINK_LIBRARIES} tenstorrent_backend) +endif() + target_link_libraries(Spatter PUBLIC ${COMMON_LINK_LIBRARIES} @@ -93,6 +108,12 @@ if (USE_CUDA) ARCHIVE DESTINATION lib) endif() +if (USE_TENSTORRENT) + install (TARGETS tenstorrent_backend + LIBRARY DESTINATION lib + ARCHIVE DESTINATION lib) +endif() + # Library/Header installation section #set(ConfigPackageLocation lib/cmake/Spatter) diff --git a/src/Spatter/Configuration.cc b/src/Spatter/Configuration.cc index 7c7079a3..6b9f9e48 100644 --- a/src/Spatter/Configuration.cc +++ b/src/Spatter/Configuration.cc @@ -921,4 +921,150 @@ void Configuration::setup() { } #endif +#ifdef USE_TENSTORRENT +Configuration::Configuration(const size_t id, + const std::string name, const std::string kernel, + const aligned_vector &pattern, + const aligned_vector &pattern_gather, + const aligned_vector &pattern_scatter, + aligned_vector &sparse, double *&dev_sparse, size_t &sparse_size, + aligned_vector &sparse_gather, double *&dev_sparse_gather, + size_t &sparse_gather_size, aligned_vector &sparse_scatter, + double *&dev_sparse_scatter, size_t &sparse_scatter_size, + aligned_vector &dense, + aligned_vector> &dense_perthread, double *&dev_dense, + size_t &dense_size, const size_t delta, const size_t delta_gather, + const size_t delta_scatter, const long int seed, const size_t wrap, + const size_t count, const size_t shared_mem, const size_t local_work_size, + const unsigned long nruns, const bool aggregate, const bool atomic, + const unsigned long verbosity) + : ConfigurationBase(id, name, kernel, pattern, pattern_gather, + pattern_scatter, sparse, dev_sparse, sparse_size, sparse_gather, + dev_sparse_gather, sparse_gather_size, sparse_scatter, + dev_sparse_scatter, sparse_scatter_size, dense, dense_perthread, + dev_dense, dense_size, delta, delta_gather, delta_scatter, seed, + wrap, count, shared_mem, local_work_size, 1, nruns, aggregate, atomic, + false, false, verbosity), + dev_pattern(nullptr), dev_pattern_gather(nullptr), + dev_pattern_scatter(nullptr) { + + setup(); +} + +Configuration::~Configuration() { + tt_device_free(dev_pattern); + tt_device_free(dev_pattern_gather); + tt_device_free(dev_pattern_scatter); + + if (dev_sparse) { + tt_device_free(dev_sparse); + dev_sparse = nullptr; + } + if (dev_sparse_gather) { + tt_device_free(dev_sparse_gather); + dev_sparse_gather = nullptr; + } + if (dev_sparse_scatter) { + tt_device_free(dev_sparse_scatter); + dev_sparse_scatter = nullptr; + } + if (dev_dense) { + tt_device_free(dev_dense); + dev_dense = nullptr; + } +} + +int Configuration::run(bool timed, unsigned long run_id) { + return ConfigurationBase::run(timed, run_id); +} + +void Configuration::gather( + bool timed, unsigned long run_id) { + size_t pattern_length = pattern.size(); + +#ifdef USE_MPI + MPI_Barrier(MPI_COMM_WORLD); +#endif + + // The wrapper is synchronous: it ends with Finish() on the mesh command + // queue, so no separate device synchronize is needed here. + float time_ms = tt_gather_wrapper( + dev_pattern, dev_sparse, dev_dense, pattern_length, delta, wrap, count); + + if (timed) + time_seconds[run_id] = ((double)time_ms / 1000.0); +} + +void Configuration::scatter( + bool timed, unsigned long run_id) { + size_t pattern_length = pattern.size(); + +#ifdef USE_MPI + MPI_Barrier(MPI_COMM_WORLD); +#endif + + float time_ms = 0.0; + + if (atomic) + time_ms = tt_scatter_atomic_wrapper( + dev_pattern, dev_sparse, dev_dense, pattern_length, delta, wrap, count); + else + time_ms = tt_scatter_wrapper( + dev_pattern, dev_sparse, dev_dense, pattern_length, delta, wrap, count); + + if (time_ms < 0.0f) { + std::cerr << "Tenstorrent backend: scatter variant not implemented" + << std::endl; + return; + } + + if (timed) + time_seconds[run_id] = ((double)time_ms / 1000.0); +} + +// gather_scatter / multi_gather / multi_scatter are declared so the class is +// concrete, but the corresponding wrappers return a negative sentinel in this +// draft. Neither is exercised by the stream or ustride suites. +void Configuration::gather_scatter( + bool timed, unsigned long run_id) { + (void)timed; + (void)run_id; + std::cerr << "Tenstorrent backend: gather_scatter not implemented" + << std::endl; +} + +void Configuration::multi_gather( + bool timed, unsigned long run_id) { + (void)timed; + (void)run_id; + std::cerr << "Tenstorrent backend: multi_gather not implemented" << std::endl; +} + +void Configuration::multi_scatter( + bool timed, unsigned long run_id) { + (void)timed; + (void)run_id; + std::cerr << "Tenstorrent backend: multi_scatter not implemented" + << std::endl; +} + +void Configuration::setup() { + ConfigurationBase::setup(); + + // One page holds the whole pattern; the kernel reads page 0 once at start. + // Empty pattern vectors stay null -- a zero-byte allocation is an error. + auto upload = [](size_t *&handle, const aligned_vector &p) { + if (p.empty()) + return; + const size_t bytes = p.size() * sizeof(uint32_t); + handle = static_cast(tt_device_alloc(bytes, bytes)); + tt_pattern_upload(handle, p.data(), p.size()); + }; + + upload(dev_pattern, pattern); + upload(dev_pattern_gather, pattern_gather); + upload(dev_pattern_scatter, pattern_scatter); +} +#endif + } // namespace Spatter diff --git a/src/Spatter/Configuration.hh b/src/Spatter/Configuration.hh index 1034463c..bb6ae799 100644 --- a/src/Spatter/Configuration.hh +++ b/src/Spatter/Configuration.hh @@ -42,6 +42,10 @@ inline void gpuAssert( } #endif + +#ifdef USE_TENSTORRENT +#include "TenstorrentBackend.hh" +#endif #include "AlignedAllocator.hh" #include "SpatterTypes.hh" #include "Timer.hh" @@ -242,6 +246,45 @@ public: }; #endif +#ifdef USE_TENSTORRENT +template <> class Configuration : public ConfigurationBase { +public: + Configuration(const size_t id, const std::string name, + const std::string kernel, const aligned_vector &pattern, + const aligned_vector &pattern_gather, + const aligned_vector &pattern_scatter, + aligned_vector &sparse, double *&dev_sparse, size_t &sparse_size, + aligned_vector &sparse_gather, double *&dev_sparse_gather, + size_t &sparse_gather_size, aligned_vector &sparse_scatter, + double *&dev_sparse_scatter, size_t &sparse_scatter_size, + aligned_vector &dense, + aligned_vector> &dense_perthread, + double *&dev_dense, size_t &dense_size, const size_t delta, + const size_t delta_gather, const size_t delta_scatter, + const long int seed, const size_t wrap, const size_t count, + const size_t shared_mem, const size_t local_work_size, + const unsigned long nruns, const bool aggregate, const bool atomic, + const unsigned long verbosity); + + ~Configuration(); + + int run(bool timed, unsigned long run_id); + void gather(bool timed, unsigned long run_id); + void scatter(bool timed, unsigned long run_id); + void gather_scatter(bool timed, unsigned long run_id); + void multi_gather(bool timed, unsigned long run_id); + void multi_scatter(bool timed, unsigned long run_id); + void setup(); + +public: + // Opaque MeshBuffer handles, typed as size_t* to match the wrapper + // signatures inherited from CudaBackend.hh. Never dereferenced on the host. + size_t *dev_pattern; + size_t *dev_pattern_gather; + size_t *dev_pattern_scatter; +}; +#endif + } // namespace Spatter #endif diff --git a/src/Spatter/Input.hh b/src/Spatter/Input.hh index 625a0d0a..7574444f 100644 --- a/src/Spatter/Input.hh +++ b/src/Spatter/Input.hh @@ -365,8 +365,10 @@ int parse_input(const int argc, char **argv, ClArgs &cl) { [](unsigned char c) { return std::tolower(c); }); if ((backend.compare("serial") != 0) && - (backend.compare("openmp") != 0) && (backend.compare("cuda") != 0)) { - std::cerr << "Valid Backends are: serial, openmp, cuda" << std::endl; + (backend.compare("openmp") != 0) && (backend.compare("cuda") != 0) && + (backend.compare("tenstorrent") != 0)) { + std::cerr << "Valid Backends are: serial, openmp, cuda, tenstorrent" + << std::endl; return -1; } if (backend.compare("openmp") == 0) { @@ -379,6 +381,12 @@ int parse_input(const int argc, char **argv, ClArgs &cl) { #ifndef USE_CUDA std::cerr << "FAIL - CUDA Backend is not Enabled" << std::endl; return -1; +#endif + } + if (backend.compare("tenstorrent") == 0) { +#ifndef USE_TENSTORRENT + std::cerr << "FAIL - Tenstorrent Backend is not Enabled" << std::endl; + return -1; #endif } break; @@ -533,6 +541,9 @@ int parse_input(const int argc, char **argv, ClArgs &cl) { #ifdef USE_OPENMP backend = "openmp"; #endif +#ifdef USE_TENSTORRENT + backend = "tenstorrent"; +#endif #ifdef USE_CUDA backend = "cuda"; #endif @@ -683,6 +694,17 @@ int parse_input(const int argc, char **argv, ClArgs &cl) { cl.dense_perthread, cl.dev_dense, cl.dense_size, delta, delta_gather, delta_scatter, seed, wrap, count, shared_mem, local_work_size, nruns, aggregate, atomic, verbosity); +#endif +#ifdef USE_TENSTORRENT + else if (backend.compare("tenstorrent") == 0) + c = std::make_unique>(0, + config_name, kernel, pattern, pattern_gather, pattern_scatter, + cl.sparse, cl.dev_sparse, cl.sparse_size, cl.sparse_gather, + cl.dev_sparse_gather, cl.sparse_gather_size, cl.sparse_scatter, + cl.dev_sparse_scatter, cl.sparse_scatter_size, cl.dense, + cl.dense_perthread, cl.dev_dense, cl.dense_size, delta, delta_gather, + delta_scatter, seed, wrap, count, shared_mem, local_work_size, nruns, + aggregate, atomic, verbosity); #endif else { std::cerr << "Invalid Backend " << backend << std::endl; @@ -790,6 +812,23 @@ int parse_input(const int argc, char **argv, ClArgs &cl) { checkCudaErrors(cudaDeviceSynchronize()); } #endif +#ifdef USE_TENSTORRENT + if (backend.compare("tenstorrent") == 0) { + auto alloc_and_copy = [](double *&dev, const aligned_vector &host) { + if (host.empty()) + return; + const size_t bytes = sizeof(double) * host.size(); + // page size == element size, so element index == DRAM page id + dev = static_cast(tt_device_alloc(bytes, sizeof(double))); + tt_memcpy_h2d(dev, host.data(), bytes); + }; + + alloc_and_copy(cl.dev_sparse, cl.sparse); + alloc_and_copy(cl.dev_sparse_gather, cl.sparse_gather); + alloc_and_copy(cl.dev_sparse_scatter, cl.sparse_scatter); + alloc_and_copy(cl.dev_dense, cl.dense); + } +#endif for (auto const &config : cl.configs) { if (config->aggregate != aggregate) { diff --git a/src/Spatter/JSONParser.cc b/src/Spatter/JSONParser.cc index 54dfd970..27093fa9 100644 --- a/src/Spatter/JSONParser.cc +++ b/src/Spatter/JSONParser.cc @@ -288,6 +288,19 @@ std::unique_ptr JSONParser::operator[]( (*data_json_ptr)[index]["wrap"], (*data_json_ptr)[index]["count"], shared_mem_, (*data_json_ptr)[index]["local-work-size"], (*data_json_ptr)[index]["nruns"], aggregate_, atomic_, verbosity_); +#endif +#ifdef USE_TENSTORRENT + else if (backend_.compare("tenstorrent") == 0) + c = std::make_unique>(index, + (*data_json_ptr)[index]["name"], (*data_json_ptr)[index]["kernel"], + pattern, pattern_gather, pattern_scatter, sparse, dev_sparse, + sparse_size, sparse_gather, dev_sparse_gather, sparse_gather_size, + sparse_scatter, dev_sparse_scatter, sparse_scatter_size, dense, + dense_perthread, dev_dense, dense_size, delta, delta_gather, + delta_scatter, (*data_json_ptr)[index]["seed"], + (*data_json_ptr)[index]["wrap"], (*data_json_ptr)[index]["count"], + shared_mem_, (*data_json_ptr)[index]["local-work-size"], + (*data_json_ptr)[index]["nruns"], aggregate_, atomic_, verbosity_); #endif else { std::cerr << "Invalid Backend " << backend_ << std::endl; diff --git a/src/Spatter/SpatterTypes.hh b/src/Spatter/SpatterTypes.hh index a486514c..52b679c1 100644 --- a/src/Spatter/SpatterTypes.hh +++ b/src/Spatter/SpatterTypes.hh @@ -11,6 +11,7 @@ namespace Spatter { struct Serial {}; struct OpenMP {}; struct CUDA {}; +struct Tenstorrent {}; } // namespace Spatter #endif diff --git a/src/Spatter/TenstorrentBackend.cc b/src/Spatter/TenstorrentBackend.cc new file mode 100644 index 00000000..bc6f94fe --- /dev/null +++ b/src/Spatter/TenstorrentBackend.cc @@ -0,0 +1,438 @@ +// Tenstorrent backend, mirroring CudaBackend.cu. Requires tt-metal. +// +// Blackhole has no FP64 fabric. That does not matter here: gather and scatter +// move bytes and do no arithmetic, so each double is treated as an opaque +// 8-byte payload. + +#include "TenstorrentBackend.hh" + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include + +using namespace tt::tt_metal; +using namespace tt::tt_metal::distributed; + +namespace { + +// One element per DRAM page, so element index == page id. Costs 8x the host +// footprint at the 64 B alignment below; the alternative needs an integer +// divide per element, which distorts what a gather benchmark measures. +constexpr uint32_t kElemBytes = sizeof(double); + +// A DRAM->L1 noc_async_read needs its L1 destination on the DRAM alignment +// boundary, 64 B on Blackhole. Packing slots tighter corrupts them silently. +constexpr uint32_t kSlotBytes = 64; + +// Chunked so the index array need not fit in L1, and so a chunk's reads are +// issued before one barrier rather than one at a time. +constexpr uint32_t kChunk = 256; + +// ---------------------------------------------------------------- device ctx +struct Context { + std::shared_ptr mesh; + CoreCoord grid{0, 0}; + + // Deliberately never destroyed. tt-metal's own singletons have static + // lifetime in libtt_metal, and the teardown order between them and ours is + // unspecified, so releasing a MeshWorkload or MeshBuffer from a static + // destructor can touch an already-dead device. + static Context &instance() { + static Context *ctx = new Context(); + static std::once_flag once; + std::call_once(once, [] { + const int device_id = [] { + const char *e = std::getenv("SPATTER_TT_DEVICE"); + return e ? std::atoi(e) : 0; + }(); + ctx->mesh = MeshDevice::create_unit_mesh(device_id); + if (!ctx->mesh) { + throw std::runtime_error("Tenstorrent: create_unit_mesh failed"); + } + ctx->grid = ctx->mesh->compute_with_storage_grid_size(); + }); + return *ctx; + } + + MeshCommandQueue &cq() { return mesh->mesh_command_queue(); } + uint32_t cores() const { return grid.x * grid.y; } + +private: + Context() = default; +}; + +// Spatter stores device allocations in `double*` / `size_t*` slots. Those slots +// hold MeshBuffer* here; the owning shared_ptr lives in this map. +std::map> ®istry() { + static auto *r = new std::map>(); + return *r; +} +std::mutex ®istry_mutex() { + static std::mutex m; + return m; +} + +std::shared_ptr lookup(const void *handle) { + std::lock_guard lock(registry_mutex()); + auto it = registry().find(const_cast(handle)); + if (it == registry().end()) { + throw std::runtime_error("Tenstorrent: unknown device buffer handle"); + } + return it->second; +} + +// Embedded so the binary is self-contained. tt-metal JIT-compiles these for the +// data-movement RISC-V at first enqueue. +const char *kGatherKernel = R"KERNEL( +#include +#include "dataflow_api.h" + +// dense[i*pattern_length + j] = sparse[pattern[j] + delta*i] +void kernel_main() { + uint32_t sparse_addr = get_arg_val(0); + uint32_t pattern_addr = get_arg_val(1); + uint32_t dense_addr = get_arg_val(2); + uint32_t pattern_len = get_arg_val(3); + uint32_t delta = get_arg_val(4); + uint32_t i_start = get_arg_val(5); + uint32_t i_count = get_arg_val(6); + + constexpr uint32_t ELEM_BYTES = get_compile_time_arg_val(0); + constexpr uint32_t SLOT_BYTES = get_compile_time_arg_val(1); + constexpr uint32_t WRITEBACK = get_compile_time_arg_val(2); + + constexpr uint32_t cb_pattern = 0; + constexpr uint32_t cb_scratch = 1; + + const InterleavedAddrGen gp = { + .bank_base_address = pattern_addr, .page_size = pattern_len * 4}; + const InterleavedAddrGen gs = { + .bank_base_address = sparse_addr, .page_size = ELEM_BYTES}; + const InterleavedAddrGen gd = { + .bank_base_address = dense_addr, .page_size = ELEM_BYTES}; + + cb_reserve_back(cb_pattern, 1); + const uint32_t pat_l1 = get_write_ptr(cb_pattern); + noc_async_read(get_noc_addr(0, gp), pat_l1, pattern_len * 4); + noc_async_read_barrier(); + volatile tt_l1_ptr uint32_t *pat = + reinterpret_cast(pat_l1); + + cb_reserve_back(cb_scratch, 1); + const uint32_t dst_l1 = get_write_ptr(cb_scratch); + + for (uint32_t k = 0; k < i_count; ++k) { + const uint32_t i = i_start + k; + const uint32_t base = delta * i; + + for (uint32_t j = 0; j < pattern_len; ++j) { + noc_async_read(get_noc_addr(pat[j] + base, gs), + dst_l1 + j * SLOT_BYTES, ELEM_BYTES); + } + noc_async_read_barrier(); + + if constexpr (WRITEBACK != 0) { + const uint32_t out = i * pattern_len; + for (uint32_t j = 0; j < pattern_len; ++j) { + noc_async_write(dst_l1 + j * SLOT_BYTES, + get_noc_addr(out + j, gd), ELEM_BYTES); + } + noc_async_write_barrier(); + } + } +} +)KERNEL"; + +// sparse[pattern[j] + delta*i] = dense[j + pattern_length*(i%wrap)] +const char *kScatterKernel = R"KERNEL( +#include +#include "dataflow_api.h" + +void kernel_main() { + uint32_t sparse_addr = get_arg_val(0); + uint32_t pattern_addr = get_arg_val(1); + uint32_t dense_addr = get_arg_val(2); + uint32_t pattern_len = get_arg_val(3); + uint32_t delta = get_arg_val(4); + uint32_t i_start = get_arg_val(5); + uint32_t i_count = get_arg_val(6); + uint32_t wrap = get_arg_val(7); + + constexpr uint32_t ELEM_BYTES = get_compile_time_arg_val(0); + constexpr uint32_t SLOT_BYTES = get_compile_time_arg_val(1); + + constexpr uint32_t cb_pattern = 0; + constexpr uint32_t cb_scratch = 1; + + const InterleavedAddrGen gp = { + .bank_base_address = pattern_addr, .page_size = pattern_len * 4}; + const InterleavedAddrGen gs = { + .bank_base_address = sparse_addr, .page_size = ELEM_BYTES}; + const InterleavedAddrGen gd = { + .bank_base_address = dense_addr, .page_size = ELEM_BYTES}; + + cb_reserve_back(cb_pattern, 1); + const uint32_t pat_l1 = get_write_ptr(cb_pattern); + noc_async_read(get_noc_addr(0, gp), pat_l1, pattern_len * 4); + noc_async_read_barrier(); + volatile tt_l1_ptr uint32_t *pat = + reinterpret_cast(pat_l1); + + cb_reserve_back(cb_scratch, 1); + const uint32_t src_l1 = get_write_ptr(cb_scratch); + + for (uint32_t k = 0; k < i_count; ++k) { + const uint32_t i = i_start + k; + const uint32_t base = delta * i; + const uint32_t in = pattern_len * (i % wrap); + + for (uint32_t j = 0; j < pattern_len; ++j) { + noc_async_read(get_noc_addr(in + j, gd), + src_l1 + j * SLOT_BYTES, ELEM_BYTES); + } + noc_async_read_barrier(); + + for (uint32_t j = 0; j < pattern_len; ++j) { + noc_async_write(src_l1 + j * SLOT_BYTES, + get_noc_addr(pat[j] + base, gs), ELEM_BYTES); + } + noc_async_write_barrier(); + } +} +)KERNEL"; + +// ------------------------------------------------------------- program cache +struct Key { + const char *kernel; + uint32_t sparse, pattern, dense; // device addresses + uint32_t pattern_len, delta, wrap, count; + bool operator<(const Key &o) const { + return std::tie(kernel, sparse, pattern, dense, pattern_len, delta, wrap, count) < + std::tie(o.kernel, o.sparse, o.pattern, o.dense, o.pattern_len, o.delta, + o.wrap, o.count); + } +}; + +std::map> &program_cache() { + static auto *c = new std::map>(); + return *c; +} + +std::shared_ptr get_or_build(const Key &key, bool is_scatter) { + auto it = program_cache().find(key); + if (it != program_cache().end()) { + return it->second; + } + + Context &ctx = Context::instance(); + const uint32_t gx = ctx.grid.x, gy = ctx.grid.y; + const uint32_t ncores = gx * gy; + + Program program = CreateProgram(); + const CoreRange all(CoreCoord{0, 0}, CoreCoord{gx - 1, gy - 1}); + + auto add_cb = [&](uint8_t index, uint32_t bytes) { + CircularBufferConfig cfg(bytes, {{index, tt::DataFormat::UInt32}}); + cfg.set_page_size(index, bytes); + CreateCircularBuffer(program, all, cfg); + }; + add_cb(0, key.pattern_len * 4); // pattern, one page + add_cb(1, key.pattern_len * kSlotBytes); // gather/scatter staging + + DataMovementConfig dm{}; + dm.processor = DataMovementProcessor::RISCV_0; + dm.noc = NOC::RISCV_0_default; + dm.compile_args = is_scatter + ? std::vector{kElemBytes, kSlotBytes} + : std::vector{kElemBytes, kSlotBytes, /*WRITEBACK=*/1u}; + + // The JIT's default -I list stops at tt_metal/hw/inc, so "dataflow_api.h" + // does not resolve without these. + if (const char *root = std::getenv("TT_METAL_RUNTIME_ROOT")) { + const std::string api = std::string(root) + "/tt_metal/hw/inc/api"; + dm.compiler_include_paths = {api, api + "/dataflow", api + "/compute", + api + "/tensor", api + "/debug", api + "/numeric"}; + } + + KernelHandle k = CreateKernelFromString(program, key.kernel, all, dm); + + const uint32_t base = key.count / ncores; + const uint32_t rem = key.count % ncores; + uint32_t start = 0; + for (uint32_t c = 0; c < ncores; ++c) { + const CoreCoord core{c % gx, c / gx}; + const uint32_t n = base + (c < rem ? 1u : 0u); + std::vector rt = {key.sparse, key.pattern, key.dense, + key.pattern_len, key.delta, start, n}; + if (is_scatter) { + rt.push_back(key.wrap); + } + SetRuntimeArgs(program, k, core, rt); + start += n; + } + + auto workload = std::make_shared(); + workload->add_program(MeshCoordinateRange(MeshCoordinate(0, 0), MeshCoordinate(0, 0)), + std::move(program)); + program_cache().emplace(key, workload); + return workload; +} + +uint32_t addr_of(const void *handle) { + return static_cast(lookup(handle)->address()); +} + +float launch(const Key &key, bool is_scatter) { + Context &ctx = Context::instance(); + auto workload = get_or_build(key, is_scatter); // JIT happens here, untimed + + const auto t0 = std::chrono::high_resolution_clock::now(); + EnqueueMeshWorkload(ctx.cq(), *workload, /*blocking=*/false); + Finish(ctx.cq()); + const auto t1 = std::chrono::high_resolution_clock::now(); + + return std::chrono::duration(t1 - t0).count(); +} + +constexpr float kUnimplemented = -1.0f; + +} // namespace + +// ---------------------------------------------------------------- public API +void *tt_device_alloc(size_t bytes, size_t page_size) { + Context &ctx = Context::instance(); + + // A page smaller than the 64 B DRAM alignment still occupies a full granule, + // so an 8-byte-element buffer costs 8x its logical size. Check before asking: + // the standard GPU test suite sizes its patterns for ~1e9 elements, which is + // 8 GB on the host and 64 GB here, and a request that large is better + // refused with a number than handed to the driver. + const size_t pages = (bytes + page_size - 1) / page_size; + const size_t granule = page_size < kSlotBytes + ? kSlotBytes + : (page_size + kSlotBytes - 1) / kSlotBytes * kSlotBytes; + const size_t on_device = pages * granule; + const size_t capacity = static_cast(ctx.mesh->num_dram_channels()) * + ctx.mesh->dram_size_per_channel(); + if (capacity && on_device > capacity * 9 / 10) { + throw std::runtime_error( + "Tenstorrent: allocation of " + std::to_string(on_device >> 20) + + " MiB exceeds device DRAM (" + std::to_string(capacity >> 20) + + " MiB). " + std::to_string(bytes >> 20) + " MiB of " + + std::to_string(page_size) + "-byte elements expands to " + + std::to_string(granule) + " bytes each at the DRAM alignment. " + "Reduce -l, or use a pattern whose elements pack into larger pages."); + } + + DeviceLocalBufferConfig local{}; + local.page_size = page_size; + local.buffer_type = BufferType::DRAM; + ReplicatedBufferConfig global{}; + global.size = bytes; + + auto buf = MeshBuffer::create(global, local, ctx.mesh.get()); + void *handle = buf.get(); + { + std::lock_guard lock(registry_mutex()); + registry().emplace(handle, std::move(buf)); + } + return handle; +} + +void tt_device_free(void *handle) { + if (!handle) { + return; + } + std::lock_guard lock(registry_mutex()); + registry().erase(handle); +} + +void tt_memcpy_h2d(void *handle, const void *src, size_t bytes) { + Context &ctx = Context::instance(); + auto buf = lookup(handle); + std::vector staging(static_cast(src), + static_cast(src) + bytes); + staging.resize(buf->size(), 0); + EnqueueWriteMeshBuffer(ctx.cq(), buf, staging, /*blocking=*/true); +} + +void tt_memcpy_d2h(void *dst, const void *handle, size_t bytes) { + Context &ctx = Context::instance(); + auto buf = lookup(handle); + std::vector staging; + EnqueueReadMeshBuffer(ctx.cq(), staging, buf, /*blocking=*/true); + std::memcpy(dst, staging.data(), bytes); +} + +// The device kernel indexes with uint32; 32 GB of DRAM cannot hold more than +// 2^32 8-byte pages anyway. +void tt_pattern_upload(void *handle, const size_t *pattern, size_t length) { + std::vector narrowed(length); + for (size_t i = 0; i < length; ++i) { + narrowed[i] = static_cast(pattern[i]); + } + tt_memcpy_h2d(handle, narrowed.data(), length * sizeof(uint32_t)); +} + +float tt_gather_wrapper(const size_t *pattern, const double *sparse, + double *dense, const size_t pattern_length, const size_t delta, + const size_t wrap, const size_t count) { + (void)wrap; + Key key{kGatherKernel, addr_of(sparse), addr_of(pattern), addr_of(dense), + static_cast(pattern_length), static_cast(delta), + 1u, static_cast(count)}; + return launch(key, /*is_scatter=*/false); +} + +float tt_scatter_wrapper(const size_t *pattern, double *sparse, + const double *dense, const size_t pattern_length, const size_t delta, + const size_t wrap, const size_t count) { + Key key{kScatterKernel, addr_of(sparse), addr_of(pattern), addr_of(dense), + static_cast(pattern_length), static_cast(delta), + static_cast(wrap), static_cast(count)}; + return launch(key, /*is_scatter=*/true); +} + +// Not implemented: atomics need a NoC atomic path, multi_* a second level of +// indirection. Neither is exercised by the stream or ustride suites. +float tt_scatter_atomic_wrapper(const size_t *, double *, const double *, + const size_t, const size_t, const size_t, const size_t) { + return kUnimplemented; +} +float tt_gather_scatter_wrapper(const size_t *, double *, const size_t *, + const double *, const size_t, const size_t, const size_t, const size_t, + const size_t) { + return kUnimplemented; +} +float tt_gather_scatter_atomic_wrapper(const size_t *, double *, const size_t *, + const double *, const size_t, const size_t, const size_t, const size_t, + const size_t) { + return kUnimplemented; +} +float tt_multi_gather_wrapper(const size_t *, const size_t *, const double *, + double *, const size_t, const size_t, const size_t, const size_t) { + return kUnimplemented; +} +float tt_multi_scatter_wrapper(const size_t *, const size_t *, double *, + const double *, const size_t, const size_t, const size_t, const size_t) { + return kUnimplemented; +} +float tt_multi_scatter_atomic_wrapper(const size_t *, const size_t *, double *, + const double *, const size_t, const size_t, const size_t, const size_t) { + return kUnimplemented; +} diff --git a/src/Spatter/TenstorrentBackend.hh b/src/Spatter/TenstorrentBackend.hh new file mode 100644 index 00000000..6ad9707f --- /dev/null +++ b/src/Spatter/TenstorrentBackend.hh @@ -0,0 +1,60 @@ +#ifndef TENSTORRENT_BACKEND_HH +#define TENSTORRENT_BACKEND_HH + +#include + +// tt-metal has no raw device pointers, so these return an opaque handle that +// Configuration keeps in the pointer slots ConfigurationBase +// provides. Nothing on the host dereferences them. page_size is the DRAM page: +// sizeof(double) for sparse/dense, pattern_length*sizeof(uint32_t) for patterns. +void *tt_device_alloc(size_t bytes, size_t page_size); +void tt_device_free(void *handle); +void tt_memcpy_h2d(void *handle, const void *src, size_t bytes); +void tt_memcpy_d2h(void *dst, const void *handle, size_t bytes); + +// Narrows Spatter's size_t pattern to the uint32 the device kernel uses. +void tt_pattern_upload(void *handle, const size_t *pattern, size_t length); + +// Mirrors CudaBackend.hh. Return value is elapsed milliseconds. +// multi_* and *_atomic are declared but return a negative sentinel. + +float tt_gather_wrapper(const size_t *pattern, const double *sparse, + double *dense, const size_t pattern_length, const size_t delta, + const size_t wrap, const size_t count); + +float tt_scatter_wrapper(const size_t *pattern, double *sparse, + const double *dense, const size_t pattern_length, const size_t delta, + const size_t wrap, const size_t count); + +float tt_scatter_atomic_wrapper(const size_t *pattern, double *sparse, + const double *dense, const size_t pattern_length, const size_t delta, + const size_t wrap, const size_t count); + +float tt_gather_scatter_wrapper(const size_t *pattern_scatter, + double *sparse_scatter, const size_t *pattern_gather, + const double *sparse_gather, const size_t pattern_length, + const size_t delta_scatter, const size_t delta_gather, const size_t wrap, + const size_t count); + +float tt_gather_scatter_atomic_wrapper(const size_t *pattern_scatter, + double *sparse_scatter, const size_t *pattern_gather, + const double *sparse_gather, const size_t pattern_length, + const size_t delta_scatter, const size_t delta_gather, const size_t wrap, + const size_t count); + +float tt_multi_gather_wrapper(const size_t *pattern, + const size_t *pattern_gather, const double *sparse, double *dense, + const size_t pattern_length, const size_t delta, const size_t wrap, + const size_t count); + +float tt_multi_scatter_wrapper(const size_t *pattern, + const size_t *pattern_scatter, double *sparse, const double *dense, + const size_t pattern_length, const size_t delta, const size_t wrap, + const size_t count); + +float tt_multi_scatter_atomic_wrapper(const size_t *pattern, + const size_t *pattern_scatter, double *sparse, const double *dense, + const size_t pattern_length, const size_t delta, const size_t wrap, + const size_t count); + +#endif diff --git a/src/main.cc b/src/main.cc index 4338b8a7..07fe6170 100644 --- a/src/main.cc +++ b/src/main.cc @@ -20,6 +20,8 @@ void print_build_info(Spatter::ClArgs &cl) { std::cout << "OpenMP" << std::endl; else if (cl.backend.compare("cuda") == 0) std::cout << "CUDA" << std::endl; + else if (cl.backend.compare("tenstorrent") == 0) + std::cout << "Tenstorrent" << std::endl; std::cout << "Aggregate Results? "; if (cl.aggregate == true) diff --git a/standard-suite/basic-tests/tenstorrent-stream.json b/standard-suite/basic-tests/tenstorrent-stream.json new file mode 100644 index 00000000..203a45d0 --- /dev/null +++ b/standard-suite/basic-tests/tenstorrent-stream.json @@ -0,0 +1,12 @@ +[ + { + "pattern": "UNIFORM:256:1:NR", + "kernel": "Gather", + "local-work-size": 1024 + }, + { + "pattern": "UNIFORM:256:1:NR", + "kernel": "Scatter", + "local-work-size": 1024 + } +] \ No newline at end of file diff --git a/standard-suite/basic-tests/tenstorrent-ustride.json b/standard-suite/basic-tests/tenstorrent-ustride.json new file mode 100644 index 00000000..46f2f757 --- /dev/null +++ b/standard-suite/basic-tests/tenstorrent-ustride.json @@ -0,0 +1,98 @@ +[ + { + "pattern": "UNIFORM:256:1:NR", + "kernel": "Scatter", + "count": 122070, + "local-work-size": 1024 + }, + { + "pattern": "UNIFORM:256:2:NR", + "kernel": "Scatter", + "count": 61035, + "local-work-size": 1024 + }, + { + "pattern": "UNIFORM:256:4:NR", + "kernel": "Scatter", + "count": 30517, + "local-work-size": 1024 + }, + { + "pattern": "UNIFORM:256:8:NR", + "kernel": "Scatter", + "count": 15258, + "local-work-size": 1024 + }, + { + "pattern": "UNIFORM:256:16:NR", + "kernel": "Scatter", + "count": 7629, + "local-work-size": 1024 + }, + { + "pattern": "UNIFORM:256:32:NR", + "kernel": "Scatter", + "count": 3814, + "local-work-size": 1024 + }, + { + "pattern": "UNIFORM:256:64:NR", + "kernel": "Scatter", + "count": 1907, + "local-work-size": 1024 + }, + { + "pattern": "UNIFORM:256:128:NR", + "kernel": "Scatter", + "count": 953, + "local-work-size": 1024 + }, + { + "pattern": "UNIFORM:256:1:NR", + "kernel": "Gather", + "count": 122070, + "local-work-size": 1024 + }, + { + "pattern": "UNIFORM:256:2:NR", + "kernel": "Gather", + "count": 61035, + "local-work-size": 1024 + }, + { + "pattern": "UNIFORM:256:4:NR", + "kernel": "Gather", + "count": 30517, + "local-work-size": 1024 + }, + { + "pattern": "UNIFORM:256:8:NR", + "kernel": "Gather", + "count": 15258, + "local-work-size": 1024 + }, + { + "pattern": "UNIFORM:256:16:NR", + "kernel": "Gather", + "count": 7629, + "local-work-size": 1024 + }, + { + "pattern": "UNIFORM:256:32:NR", + "kernel": "Gather", + "count": 3814, + "local-work-size": 1024 + }, + { + "pattern": "UNIFORM:256:64:NR", + "kernel": "Gather", + "count": 1907, + "local-work-size": 1024 + }, + { + "pattern": "UNIFORM:256:128:NR", + "kernel": "Gather", + "count": 953, + "local-work-size": 1024 + } +] \ No newline at end of file diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index e9a43cdb..426f334a 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -39,6 +39,13 @@ if (USE_CUDA) parse_run_config_suite_gpu) endif() +if (USE_TENSTORRENT) + set(SPATTER_TESTS ${SPATTER_TESTS} + standard_suite_stream_tenstorrent + standard_suite_uniform_tenstorrent + parse_run_config_suite_tenstorrent) +endif() + foreach (APP ${SPATTER_TESTS}) add_executable(${APP} ${APP}.cc) target_link_libraries(${APP} Spatter) diff --git a/tests/misc/run-crnch-tenstorrent.sh b/tests/misc/run-crnch-tenstorrent.sh new file mode 100755 index 00000000..f66e89c3 --- /dev/null +++ b/tests/misc/run-crnch-tenstorrent.sh @@ -0,0 +1,39 @@ +#!/bin/bash +#SBATCH -Jspatter-ci-tenstorrent # Job name +#SBATCH -N1 --cpus-per-task=4 # Number of nodes and CPUs per node required +#SBATCH --mem-per-cpu=4G # Memory per core +#SBATCH -t 00:30:00 # Duration of the job (Ex: 30 mins) +#SBATCH -p rg-nextgen-hpc # Partition Name +#SBATCH -o /tools/ci-reports/spatter-tenstorrent-test-%j.out # Combined output and error messages file +#SBATCH --gres r5accel:p150a:2 # Request both Blackhole cards on the node +#SBATCH -W # Do not exit until the submitted job terminates. + +# Copied from run-crnch-cuda.sh. Four Tenstorrent-specific notes: +# +# 1. The GRES is r5accel:p150a, NOT gpu -- asking for gpu on this host gets an +# L40S, which is a different device in the same chassis. +# 2. Both cards are requested even though the tests use one. With a :1 +# allocation the cgroup blocks /dev/tenstorrent/1, and UMD's enumeration +# walks every card on the host and fails on the blocked one. TT_VISIBLE_DEVICES +# below is what actually restricts execution to a single chip. +# 3. TT_METAL_RUNTIME_ROOT must point at the directory containing tt_metal/, or +# tt-metal aborts with "Root Directory is not set" before opening a device. +# 4. -DUSE_TENSTORRENT=ON is the flag the CMake actually consumes. (run-crnch-cuda.sh +# passes -DBACKEND=cuda -DCOMPILER=nvcc, but BACKEND and COMPILER are not read +# anywhere in the CMake on main -- see the note in the PR description.) + +cd $GITHUB_WORKSPACE +hostname + +# This line allows the GH runner to use the module command on the targeted node +source /tools/misc/.read_profile + +# tt-metal host library + Python-side runtime root come from the ttnn install. +source ~/ttnn-env/bin/activate +export TT_METAL_RUNTIME_ROOT=$(python3 -c 'import ttnn, os; print(os.path.dirname(ttnn.__file__))') +export TT_VISIBLE_DEVICES=0 + +cmake -DUSE_TENSTORRENT=ON -B build_tenstorrent_workflow -S . +make -C build_tenstorrent_workflow +cd build_tenstorrent_workflow +make test diff --git a/tests/parse_run_config_suite_tenstorrent.cc b/tests/parse_run_config_suite_tenstorrent.cc new file mode 100644 index 00000000..04e9f4be --- /dev/null +++ b/tests/parse_run_config_suite_tenstorrent.cc @@ -0,0 +1,168 @@ +#include +#include +#include + +#include "Spatter/Configuration.hh" +#include "Spatter/Input.hh" + +int parse_check(int argc_, char **argv_, Spatter::ClArgs &cl) { + if (Spatter::parse_input(argc_, argv_, cl) != 0) { + std::cerr << "Parse Input Failed" << std::endl; + return EXIT_FAILURE; + } + + if (cl.configs.size() != 1) { + std::cerr + << "Test failure on Concurrent Pattern: Expected number of runs to " + "be 1, actually was " + << cl.configs.size() << std::endl; + return EXIT_FAILURE; + } + + if (cl.configs[0] == NULL) { + std::cerr + << "Test failure on Concurrent Pattern: Failed to create or allocate " + "a ConfigurationBase object" + << std::endl; + return EXIT_FAILURE; + } + + return EXIT_SUCCESS; +} + +int z_tests(int argc_, char **argv_) { + asprintf(&argv_[2], "-z100"); + + Spatter::ClArgs cl1; + if (parse_check(argc_, argv_, cl1) == EXIT_FAILURE) + return EXIT_FAILURE; + + free(argv_[2]); + + if (cl1.configs[0]->local_work_size != 100) { + std::cerr << "Test failure on Run_Config Suite: -z with argument 100 had " + "incorrect value of " + << cl1.configs[0]->local_work_size << "." << std::endl; + return EXIT_FAILURE; + } + + asprintf(&argv_[2], "-z500"); + + Spatter::ClArgs cl2; + if (parse_check(argc_, argv_, cl2) == EXIT_FAILURE) + return EXIT_FAILURE; + + free(argv_[2]); + + if (cl2.configs[0]->local_work_size != 500) { + std::cerr << "Test failure on Run_Config Suite: -z with argument 500 had " + "incorrect value of " + << cl2.configs[0]->local_work_size << "." << std::endl; + return EXIT_FAILURE; + } + + asprintf(&argv_[2], "--local-work-size=1000"); + + Spatter::ClArgs cl3; + if (parse_check(argc_, argv_, cl3) == EXIT_FAILURE) + return EXIT_FAILURE; + + free(argv_[2]); + + if (cl3.configs[0]->local_work_size != 1000) { + std::cerr << "Test failure on Run_Config Suite: -z with argument 1000 had " + "incorrect value of " + << cl3.configs[0]->local_work_size << "." << std::endl; + return EXIT_FAILURE; + } + + return EXIT_SUCCESS; +} + +int m_tests(int argc_, char **argv_) { + asprintf(&argv_[2], "-m100"); + + Spatter::ClArgs cl1; + if (parse_check(argc_, argv_, cl1) == EXIT_FAILURE) + return EXIT_FAILURE; + + free(argv_[2]); + + if (cl1.configs[0]->shmem != 100) { + std::cerr << "Test Failure on Run_Config Suite: -m with argument 100 had " + "incorrect value of " + << cl1.configs[0]->shmem << "." << std::endl; + return EXIT_FAILURE; + } + + asprintf(&argv_[2], "-m500"); + + Spatter::ClArgs cl2; + if (parse_check(argc_, argv_, cl2) == EXIT_FAILURE) + return EXIT_FAILURE; + + free(argv_[2]); + + if (cl2.configs[0]->shmem != 500) { + std::cerr << "Test failure on Run_Config Suite: -m with argument 500 had " + "incorrect value of " + << cl2.configs[0]->shmem << "." << std::endl; + return EXIT_FAILURE; + } + + asprintf(&argv_[2], "--shared-mem=1000"); + + Spatter::ClArgs cl3; + if (parse_check(argc_, argv_, cl3) == EXIT_FAILURE) + return EXIT_FAILURE; + + free(argv_[2]); + + if (cl3.configs[0]->shmem != 1000) { + std::cerr << "Test failure on Run_Config Suite: -m with argument 1000 had " + "incorrect value of " + << cl3.configs[0]->shmem << "." << std::endl; + return EXIT_FAILURE; + } + + return EXIT_SUCCESS; +} + +int main(int argc, char **argv) { + (void)argc; + (void)argv; + + int argc_ = 4; + char **argv_ = (char **)malloc(sizeof(char *) * argc_); + + int ret; + ret = asprintf(&argv_[0], "./spatter"); + if (ret == -1) + return EXIT_FAILURE; + + ret = asprintf(&argv_[1], "-p1,2,3,4"); + if (ret == -1) + return EXIT_FAILURE; + + // Unlike parse_run_config_suite_gpu.cc, which leaves the backend at its + // default, this passes -b so the Tenstorrent Configuration is actually + // constructed. The CI job that runs it holds a card. + ret = asprintf(&argv_[3], "-btenstorrent"); + if (ret == -1) + return EXIT_FAILURE; + + // local-work-size z + if (z_tests(argc_, argv_) != EXIT_SUCCESS) + return EXIT_FAILURE; + + // shared-mem m + if (m_tests(argc_, argv_) != EXIT_SUCCESS) + return EXIT_FAILURE; + + free(argv_[0]); + free(argv_[1]); + free(argv_[3]); + free(argv_); + + return EXIT_SUCCESS; +} diff --git a/tests/standard_suite_stream_tenstorrent.cc b/tests/standard_suite_stream_tenstorrent.cc new file mode 100644 index 00000000..d8327028 --- /dev/null +++ b/tests/standard_suite_stream_tenstorrent.cc @@ -0,0 +1,27 @@ +#include +#include + +int tenstorrent_stream_test() { + char *command; + + int ret = asprintf(&command, + "../spatter -b tenstorrent -f " + "../../standard-suite/basic-tests/tenstorrent-stream.json"); + if (ret == -1 || system(command) != EXIT_SUCCESS) { + std::cerr << "Test failure on " << command << std::endl; + return EXIT_FAILURE; + } + + free(command); + return EXIT_SUCCESS; +} + +int main(int argc, char **argv) { + (void)argc; + (void)argv; + + if (tenstorrent_stream_test() != EXIT_SUCCESS) + return EXIT_FAILURE; + + return EXIT_SUCCESS; +} diff --git a/tests/standard_suite_uniform_tenstorrent.cc b/tests/standard_suite_uniform_tenstorrent.cc new file mode 100644 index 00000000..b74e3cf4 --- /dev/null +++ b/tests/standard_suite_uniform_tenstorrent.cc @@ -0,0 +1,27 @@ +#include +#include + +int tenstorrent_ustride_test() { + char *command; + + int ret = asprintf(&command, + "../spatter -b tenstorrent -f " + "../../standard-suite/basic-tests/tenstorrent-ustride.json"); + if (ret == -1 || system(command) != EXIT_SUCCESS) { + std::cerr << "Test failure on " << command << std::endl; + return EXIT_FAILURE; + } + + free(command); + return EXIT_SUCCESS; +} + +int main(int argc, char **argv) { + (void)argc; + (void)argv; + + if (tenstorrent_ustride_test() != EXIT_SUCCESS) + return EXIT_FAILURE; + + return EXIT_SUCCESS; +}