ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-13 01:36:06 +00:00
Code Issues Packages Projects Releases Wiki Activity

Gemm_c_shuffle (4 layouts) X (fp32 bf16 int8) (#131)

Browse Source 
* [What] Separate fixpoint gemm from gemm example
[Why] let example of gemm_int8 be pure gemm.
[What]
1. Add gemm_requant_relu_requant,
2. Let CDataType be int32 in pure gemm, because no one use int8 CDataType. It is also part of gemm_requant_relu_requant

* Fix path

* Revise cmakelist due to merge develop

* Add gemm fp16 test

* Extract PrepareGemmTensor

* Extract TestGemm

* Add test for different layout

* Add 4 layouts of shuffle version of fp32

* Add 4 layouts of shuffle version of int8

* Add 4 layouts of shuffle version of bf16

* replace all DeviceGemmPtr_ with DeviceGemmNoOpPtr to fit naming convension

* Add test for non-shuffle verstion of gemm

* Fix typo

* Print kernel information

* Add rest of the fp32 kernel to the test

* 1. Add rest of the fp16 device iop.
2. Mark the invalid device operation

Co-authored-by: rocking <chunylai@amd.com>
This commit is contained in:
rocking5566
2022-03-22 04:59:51 +08:00
committed by GitHub
parent b51808d7a5
commit 485ea46a40
24 changed files with 1497 additions and 322 deletions
Download Patch File Download Diff File Expand all files Collapse all files

43
test/include/test_util.hpp
View File

@@ -54,6 +54,49 @@ check_err(const std::vector<T>& out,
return res;
}
bool check_err(const std::vector<_Float16>& out,
const std::vector<_Float16>& ref,
const std::string& msg,
_Float16 rtol = static_cast<_Float16>(1e-3f),
_Float16 atol = static_cast<_Float16>(1e-3f))
{
if(out.size() != ref.size())
{
std::cout << "out.size() != ref.size(), :" << out.size() << " != " << ref.size()
<< std::endl
<< msg << std::endl;
return false;
}
bool res{true};
int err_count = 0;
double err = 0;
double max_err = std::numeric_limits<_Float16>::min();
for(std::size_t i = 0; i < ref.size(); ++i)
{
double out_ = double(out[i]);
double ref_ = double(ref[i]);
err = std::abs(out_ - ref_);
if(err > atol + rtol * std::abs(ref_) || !std::isfinite(out_) || !std::isfinite(ref_))
{
max_err = err > max_err ? err : max_err;
err_count++;
if(err_count < 5)
{
std::cout << std::setw(12) << std::setprecision(7) << "out[" << i << "] != ref["
<< i << "]: " << out_ << "!=" << ref_ << std::endl
<< msg << std::endl;
}
res = false;
}
}
if(!res)
{
std::cout << std::setw(12) << std::setprecision(7) << "max err: " << max_err << std::endl;
}
return res;
}
template <typename T>
typename std::enable_if<std::is_integral<T>::value, bool>::type check_err(
const std::vector<T>& out, const std::vector<T>& ref, const std::string& msg, T = 0, T = 0)