Skip to content

GitLab

Unverified Commit 44a85546 authored by lvhan028's avatar lvhan028 Committed by GitHub
Browse files

Fix build test error and move turbmind csrc test cases to `tests/csrc` (#188) 

* fix build tests failure

* move src test cases to tests/csrc
parent 5545bbc5
Hide whitespace changes
Inline Side-by-side
Showing with 0 additions and 154 deletions
tests/unittests/test_activation.cu deleted 100644 → 0
View file @ 5545bbc5
#include <iostream> // snprintf
#include <string> // std::string
#include <vector> // std::vector
#include "src/turbomind/kernels/activation_kernels.h"
#include "src/turbomind/utils/cuda_utils.h"
#include "src/turbomind/utils/memory_utils.h"
#include "src/turbomind/utils/logger.h"
#include "unittest_utils.h"
using namespace turbomind;
#define PRINT_LIMIT 16
#define EPSILON (1e-20)
#define EPSILON_FP16 (1e-10)
struct TestCase {
std::string name;
size_t m;
size_t n;
size_t ite;
std::string toString()
{
char buf[100];
snprintf(buf, sizeof(buf), "TestCase[name=%s, m=%ld, n=%ld]", name.c_str(), m, n);
return buf;
}
void print()
{
TM_LOG_INFO(toString());
}
};
template<typename T>
void testActivationKernel(TestCase tc)
{
const int m = tc.m;
const int n = tc.n;
cudaStream_t stream;
cudaStreamCreate(&stream);
T *output_baseline, *output_opt1, *bias;
deviceMalloc(&output_baseline, m * n);
deviceMalloc(&output_opt1, m * n);
deviceMalloc(&bias, n);
cudaD2Dcpy(output_opt1, output_baseline, m * n);
invokeGenericActivation<GeluActivation>(output_baseline,
(const T*) bias,
(const T*) nullptr,
(const T*) nullptr,
(const int*) nullptr,
(const T*) nullptr,
m,
n,
0,
(const float*) nullptr,
(const float*) nullptr,
stream);
invokeAddBiasGeluV2(output_opt1, bias, (const int*) nullptr, (const T*) nullptr, m, n, stream);
bool passed = checkResult(tc.name, output_baseline, output_opt1, m * n, true, true);
FT_CHECK(passed);
const int ite = tc.ite;
CudaTimer cuda_timer_baseline(stream);
// warmup
for (int i = 0; i < ite; i++) {
invokeGenericActivation<GeluActivation>(output_baseline,
(const T*) bias,
(const T*) nullptr,
(const T*) nullptr,
(const int*) nullptr,
(const T*) nullptr,
m,
n,
0,
(const float*) nullptr,
(const float*) nullptr,
stream);
}
cuda_timer_baseline.start();
for (int i = 0; i < ite; i++) {
invokeGenericActivation<GeluActivation>(output_baseline,
(const T*) bias,
(const T*) nullptr,
(const T*) nullptr,
(const int*) nullptr,
(const T*) nullptr,
m,
n,
0,
(const float*) nullptr,
(const float*) nullptr,
stream);
}
float total_time_baseline = cuda_timer_baseline.stop();
CudaTimer cuda_timer_opt(stream);
// warmup
for (int i = 0; i < ite; i++) {
invokeAddBiasGeluV2(output_baseline, bias, (const int*) nullptr, (const T*) nullptr, m, n, stream);
}
cuda_timer_opt.start();
for (int i = 0; i < ite; i++) {
invokeAddBiasGeluV2(output_baseline, bias, (const int*) nullptr, (const T*) nullptr, m, n, stream);
}
float total_time_opt = cuda_timer_opt.stop();
TM_LOG_INFO("%s baseline_time: %f us, opt_time: %f us, speedup: %f (ite: %d)",
tc.toString().c_str(),
total_time_baseline / ite * 1000.f,
total_time_opt / ite * 1000.f,
total_time_baseline / total_time_opt,
ite);
deviceFree(output_baseline);
deviceFree(output_opt1);
deviceFree(bias);
}
int main()
{
printf("[INFO] Device: %s \n", getDeviceName().c_str());
std::vector<TestCase> test_cases{
// TC: name / m / n
TestCase{"addBiasGelu", 32, 1024, 1000},
TestCase{"addBiasGelu", 128, 1024, 1000},
TestCase{"addBiasGelu", 2048, 1024, 1000},
TestCase{"addBiasGelu", 32, 3072, 1000},
TestCase{"addBiasGelu", 128, 3072, 1000},
TestCase{"addBiasGelu", 2048, 3072, 1000},
TestCase{"addBiasGelu", 32, 4096, 1000},
TestCase{"addBiasGelu", 128, 4096, 1000},
TestCase{"addBiasGelu", 2048, 4096, 1000},
TestCase{"addBiasGelu", 32, 8192, 1000},
TestCase{"addBiasGelu", 128, 8192, 1000},
TestCase{"addBiasGelu", 2048, 8192, 1000},
TestCase{"addBiasGelu", 32, 49152, 1000},
TestCase{"addBiasGelu", 128, 49152, 1000},
TestCase{"addBiasGelu", 2048, 49152, 1000},
TestCase{"addBiasGelu", 32, 81920, 1000},
TestCase{"addBiasGelu", 128, 81920, 1000},
TestCase{"addBiasGelu", 2048, 81920, 1000},
};
for (auto& tc : test_cases) {
// testActivationKernel<float>(tc);
testActivationKernel<half>(tc);
}
TM_LOG_INFO("testActivationKernel done");
return 0;
}
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment