Skip to content

GitLab

Commit 93d77eaf authored by Aleksander Dudek's avatar Aleksander Dudek
Browse files

[CK_TILE] Add EnvLogging - wrap gemm kernel error messages

parent 12142253
Hide whitespace changes
Inline Side-by-side
Showing with 70 additions and 30 deletions
include/ck_tile/ops/gemm/kernel/gemm_kernel.hpp
View file @ 93d77eaf
...@@ -167,17 +167,15 @@ struct GemmKernel ...@@ -167,17 +167,15 @@ struct GemmKernel
CK_TILE_HOST static bool IsSupportedArgument(const GemmKernelArgs& kargs) CK_TILE_HOST static bool IsSupportedArgument(const GemmKernelArgs& kargs)
{ {
if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING)))
{
std::cout << "Testing CK_TILE_ENV logging... " << __FILE__ << ":" << __LINE__
<< ", in function: " << __func__ << std::endl;
}
if constexpr(EpiloguePipeline::template GetVectorSizeC<CDataType>() % 2 != 0 && if constexpr(EpiloguePipeline::template GetVectorSizeC<CDataType>() % 2 != 0 &&
is_any_of<CDataType, fp16_t, bf16_t>::value) is_any_of<CDataType, fp16_t, bf16_t>::value)
{ {
if(kargs.k_batch != 1) if(kargs.k_batch != 1)
{ {
std::cerr << "Conditions not met for Kbatch >1 !" << std::endl; if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING)))
{
std::cerr << "Conditions not met for Kbatch >1 !" << std::endl;
}
return false; return false;
} }
} }
...@@ -186,14 +184,21 @@ struct GemmKernel ...@@ -186,14 +184,21 @@ struct GemmKernel
{ {
if(kargs.K % TilePartitioner::KPerBlock != 0 && GemmPipeline::kPadK == false) if(kargs.K % TilePartitioner::KPerBlock != 0 && GemmPipeline::kPadK == false)
{ {
std::cerr << "Can't support K that is not a multiple of KPerBlock" if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING)))
" without padding!" {
<< std::endl; std::cerr << "Can't support K that is not a multiple of KPerBlock"
" without padding!"
<< std::endl;
}
return false; return false;
} }
if(kargs.K % GemmPipeline::GetVectorSizeA() != 0) if(kargs.K % GemmPipeline::GetVectorSizeA() != 0)
{ {
std::cerr << "K is not a multiple of vector load size for A tensor!" << std::endl; if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING)))
{
std::cerr << "K is not a multiple of vector load size for A tensor!"
<< std::endl;
}
return false; return false;
} }
} }
...@@ -201,14 +206,21 @@ struct GemmKernel ...@@ -201,14 +206,21 @@ struct GemmKernel
{ {
if(kargs.M % TilePartitioner::MPerBlock != 0 && GemmPipeline::kPadM == false) if(kargs.M % TilePartitioner::MPerBlock != 0 && GemmPipeline::kPadM == false)
{ {
std::cerr << "Can't support M that is not a multiple of MPerBlock" if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING)))
" without padding!" {
<< std::endl; std::cerr << "Can't support M that is not a multiple of MPerBlock"
" without padding!"
<< std::endl;
}
return false; return false;
} }
if(kargs.M % GemmPipeline::GetVectorSizeA() != 0) if(kargs.M % GemmPipeline::GetVectorSizeA() != 0)
{ {
std::cerr << "M is not a multiple of vector load size for A tensor!" << std::endl; if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING)))
{
std::cerr << "M is not a multiple of vector load size for A tensor!"
<< std::endl;
}
return false; return false;
} }
} }
...@@ -217,14 +229,21 @@ struct GemmKernel ...@@ -217,14 +229,21 @@ struct GemmKernel
{ {
if(kargs.N % TilePartitioner::NPerBlock != 0 && GemmPipeline::kPadN == false) if(kargs.N % TilePartitioner::NPerBlock != 0 && GemmPipeline::kPadN == false)
{ {
std::cerr << "Can't support N that is not a multiple of NPerBlock" if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING)))
" without padding!" {
<< std::endl; std::cerr << "Can't support N that is not a multiple of NPerBlock"
" without padding!"
<< std::endl;
}
return false; return false;
} }
if(kargs.N % GemmPipeline::GetVectorSizeB() != 0) if(kargs.N % GemmPipeline::GetVectorSizeB() != 0)
{ {
std::cerr << "N is not a multiple of vector load size for B tensor!" << std::endl; if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING)))
{
std::cerr << "N is not a multiple of vector load size for B tensor!"
<< std::endl;
}
return false; return false;
} }
} }
...@@ -232,14 +251,21 @@ struct GemmKernel ...@@ -232,14 +251,21 @@ struct GemmKernel
{ {
if(kargs.K % TilePartitioner::KPerBlock != 0 && GemmPipeline::kPadK == false) if(kargs.K % TilePartitioner::KPerBlock != 0 && GemmPipeline::kPadK == false)
{ {
std::cerr << "Can't support K that is not a multiple of KPerBlock" if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING)))
" without padding!" {
<< std::endl; std::cerr << "Can't support K that is not a multiple of KPerBlock"
" without padding!"
<< std::endl;
}
return false; return false;
} }
if(kargs.K % GemmPipeline::GetVectorSizeB() != 0) if(kargs.K % GemmPipeline::GetVectorSizeB() != 0)
{ {
std::cerr << "K is not a multiple of vector load size for B tensor!" << std::endl; if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING)))
{
std::cerr << "K is not a multiple of vector load size for B tensor!"
<< std::endl;
}
return false; return false;
} }
} }
...@@ -248,14 +274,21 @@ struct GemmKernel ...@@ -248,14 +274,21 @@ struct GemmKernel
{ {
if(kargs.N % TilePartitioner::NPerBlock != 0 && GemmPipeline::kPadN == false) if(kargs.N % TilePartitioner::NPerBlock != 0 && GemmPipeline::kPadN == false)
{ {
std::cerr << "Can't support N that is not a multiple of NPerBlock" if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING)))
" without padding!" {
<< std::endl; std::cerr << "Can't support N that is not a multiple of NPerBlock"
" without padding!"
<< std::endl;
}
return false; return false;
} }
if(kargs.N % EpiloguePipeline::template GetVectorSizeC<CDataType>() != 0) if(kargs.N % EpiloguePipeline::template GetVectorSizeC<CDataType>() != 0)
{ {
std::cerr << "N is not a multiple of vector load size for C tensor!" << std::endl; if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING)))
{
std::cerr << "N is not a multiple of vector load size for C tensor!"
<< std::endl;
}
return false; return false;
} }
} }
...@@ -263,14 +296,21 @@ struct GemmKernel ...@@ -263,14 +296,21 @@ struct GemmKernel
{ {
if(kargs.M % TilePartitioner::MPerBlock != 0 && GemmPipeline::kPadM == false) if(kargs.M % TilePartitioner::MPerBlock != 0 && GemmPipeline::kPadM == false)
{ {
std::cerr << "Can't support M that is not a multiple of MPerBlock" if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING)))
" without padding!" {
<< std::endl; std::cerr << "Can't support M that is not a multiple of MPerBlock"
" without padding!"
<< std::endl;
}
return false; return false;
} }
if(kargs.M % EpiloguePipeline::template GetVectorSizeC<CDataType>() != 0) if(kargs.M % EpiloguePipeline::template GetVectorSizeC<CDataType>() != 0)
{ {
std::cerr << "M is not a multiple of vector load size for C tensor!" << std::endl; if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING)))
{
std::cerr << "M is not a multiple of vector load size for C tensor!"
<< std::endl;
}
return false; return false;
} }
} }
......
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