Commit 7863805d authored by Bartlomiej Wroblewski's avatar Bartlomiej Wroblewski
Browse files

Review: Improve error messages

parent e7836e46
...@@ -123,7 +123,7 @@ struct BlockwiseGemmDpp_k0mk1_k0nk1_m0n0m1n1m2n2 ...@@ -123,7 +123,7 @@ struct BlockwiseGemmDpp_k0mk1_k0nk1_m0n0m1n1m2n2
{ {
static_assert(AK0MK1BlockDesc::IsKnownAtCompileTime() && static_assert(AK0MK1BlockDesc::IsKnownAtCompileTime() &&
BK0NK1BlockDesc::IsKnownAtCompileTime(), BK0NK1BlockDesc::IsKnownAtCompileTime(),
"wrong! Desc should be known at compile-time"); "Wrong! Block descriptors should be known at the time of compilation.");
#if defined(__HIP_DEVICE_COMPILE__) #if defined(__HIP_DEVICE_COMPILE__)
// Host wave size can be different than the device one and this assert could fail for host, // Host wave size can be different than the device one and this assert could fail for host,
...@@ -132,8 +132,10 @@ struct BlockwiseGemmDpp_k0mk1_k0nk1_m0n0m1n1m2n2 ...@@ -132,8 +132,10 @@ struct BlockwiseGemmDpp_k0mk1_k0nk1_m0n0m1n1m2n2
"ThisThreadBlock::GetNumOfThread() != MWaves * NWaves * WaveSize\n"); "ThisThreadBlock::GetNumOfThread() != MWaves * NWaves * WaveSize\n");
#endif #endif
static_assert(MPerBlock % (MPerDpp * MRepeat) == 0 && NPerBlock % (NPerDpp * NRepeat) == 0, static_assert(MPerBlock % (MPerDpp * MRepeat) == 0,
"wrong!"); "Invalid parameters. MPerBlock must be divisible by MPerDpp * MRepeat.");
static_assert(NPerBlock % (NPerDpp * NRepeat) == 0,
"Invalid parameters. NPerBlock must be divisible by NPerDpp * NRepeat.");
} }
__host__ __device__ static constexpr auto GetCThreadDescriptor_M0_N0_M1_N1_M2_N2() __host__ __device__ static constexpr auto GetCThreadDescriptor_M0_N0_M1_N1_M2_N2()
......
...@@ -260,11 +260,14 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_dpp ...@@ -260,11 +260,14 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_dpp
__host__ static constexpr bool CheckValidity(const Problem& problem) __host__ static constexpr bool CheckValidity(const Problem& problem)
{ {
static_assert(is_known_at_compile_time<remove_cv_t<decltype(K1)>>::value, static_assert(is_known_at_compile_time<remove_cv_t<decltype(K1)>>::value,
"wrong! K1 need to be known at compile-time"); "Wrong! K1 must be known at the time of compilation.");
static_assert((MPerBlock % (MPerDpp * MDppPerWave) == 0) && static_assert(
(NPerBlock % (NDppPerWave * NPerDpp)) == 0, MPerBlock % (MPerDpp * MDppPerWave) == 0,
"Invalid tuning param!"); "Invalid tuning parameters! MPerBlock must be divisible by MPerDpp * MDppPerWave.");
static_assert(
NPerBlock % (NPerDpp * NDppPerWave) == 0,
"Invalid tuning parameters! NPerBlock must be divisible by NPerDpp * NDppPerWave.");
if constexpr(!(GemmSpec == tensor_operation::device::GemmSpecialization::MPadding || if constexpr(!(GemmSpec == tensor_operation::device::GemmSpecialization::MPadding ||
GemmSpec == tensor_operation::device::GemmSpecialization::MNPadding || GemmSpec == tensor_operation::device::GemmSpecialization::MNPadding ||
......
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