mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-14 10:09:41 +00:00
* WIP POC of dispatcher
* Dispatcher python workflow setup.
* Dispatcher cleanup and updates.
Further dispatcher cleanup and updates.
Build fixes
Improvements and python to CK example
Improvements to readme
* Fixes to python paths
* Cleaning up code
* Improving dispatcher support for different arch
Fixing typos
* Fix formatting errors
* Cleaning up examples
* Improving codegeneration
* Improving and fixing C++ examples
* Adding conv functionality (fwd,bwd,bwdw) and examples.
* Fixes based on feedback.
* Further fixes based on feedback.
* Adding stress test for autogeneration and autocorrection, and fixing preshuffle bug.
* Another round of improvements based on feedback.
* Trimming out unnecessary code.
* Fixing the multi-D implementation.
* Using gpu verification for gemms and fixing convolutions tflops calculation.
* Fix counter usage issue and arch filtering per ops.
* Adding changelog and other fixes.
* Improve examples and resolve critical bugs.
* Reduce build time for python examples.
* Fixing minor bug.
* Fix compilation error.
* Improve installation instructions for dispatcher.
* Add docker based installation instructions for dispatcher.
* Fixing arch-based filtering to match tile engine.
* Remove dead code and fix arch filtering.
* Minor bugfix.
* Updates after rebase.
* Trimming code.
* Fix copyright headers.
* Consolidate examples, cut down code.
* Minor fixes.
* Improving python examples.
* Update readmes.
* Remove conv functionality.
* Cleanup following conv removable.
[ROCm/composable_kernel commit: 9e049a32a1]
97 lines
1.9 KiB
C++
97 lines
1.9 KiB
C++
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
|
|
// SPDX-License-Identifier: MIT
|
|
|
|
/// Unit tests for Problem using Google Test
|
|
|
|
#include "ck_tile/dispatcher/problem.hpp"
|
|
#include <gtest/gtest.h>
|
|
|
|
using namespace ck_tile::dispatcher;
|
|
|
|
TEST(ProblemTest, DefaultConstruction)
|
|
{
|
|
Problem p;
|
|
EXPECT_EQ(p.M, 0);
|
|
EXPECT_EQ(p.N, 0);
|
|
EXPECT_EQ(p.K, 0);
|
|
EXPECT_EQ(p.k_batch, 1);
|
|
EXPECT_FALSE(p.is_valid());
|
|
}
|
|
|
|
TEST(ProblemTest, ConstructorWithDimensions)
|
|
{
|
|
Problem p(1024, 1024, 1024);
|
|
EXPECT_EQ(p.M, 1024);
|
|
EXPECT_EQ(p.N, 1024);
|
|
EXPECT_EQ(p.K, 1024);
|
|
EXPECT_TRUE(p.is_valid());
|
|
}
|
|
|
|
TEST(ProblemTest, Validation)
|
|
{
|
|
Problem p;
|
|
|
|
// Invalid: all zeros
|
|
p.M = 0;
|
|
p.N = 0;
|
|
p.K = 0;
|
|
EXPECT_FALSE(p.is_valid());
|
|
|
|
// Invalid: negative
|
|
p.M = -1;
|
|
p.N = 1024;
|
|
p.K = 1024;
|
|
EXPECT_FALSE(p.is_valid());
|
|
|
|
// Invalid: zero K
|
|
p.M = 1024;
|
|
p.N = 1024;
|
|
p.K = 0;
|
|
EXPECT_FALSE(p.is_valid());
|
|
|
|
// Valid
|
|
p.M = 1024;
|
|
p.N = 1024;
|
|
p.K = 1024;
|
|
EXPECT_TRUE(p.is_valid());
|
|
|
|
// Invalid k_batch
|
|
p.k_batch = 0;
|
|
EXPECT_FALSE(p.is_valid());
|
|
|
|
p.k_batch = 1;
|
|
EXPECT_TRUE(p.is_valid());
|
|
}
|
|
|
|
TEST(ProblemTest, NumOps)
|
|
{
|
|
Problem p(100, 200, 300);
|
|
|
|
// 2 * M * N * K (multiply-add = 2 ops)
|
|
std::int64_t expected = 2 * 100 * 200 * 300;
|
|
EXPECT_EQ(p.num_ops(), expected);
|
|
}
|
|
|
|
TEST(ProblemTest, Configuration)
|
|
{
|
|
Problem p(1024, 1024, 1024);
|
|
|
|
// Set preferences
|
|
p.prefer_persistent = true;
|
|
p.enable_validation = true;
|
|
p.smem_budget = 65536;
|
|
p.k_batch = 2;
|
|
|
|
EXPECT_TRUE(p.prefer_persistent);
|
|
EXPECT_TRUE(p.enable_validation);
|
|
EXPECT_EQ(p.smem_budget, 65536);
|
|
EXPECT_EQ(p.k_batch, 2);
|
|
}
|
|
|
|
TEST(ProblemTest, LargeDimensions)
|
|
{
|
|
Problem p(1024, 1024, 1024); // Use smaller but still large dimensions
|
|
EXPECT_TRUE(p.is_valid());
|
|
EXPECT_GT(p.num_ops(), 0);
|
|
}
|