Give recast_layout and the flagged downcast a diagnostic

[GrigoryEvko]

Summary

recast_layout(Swizzle) takes a fourth branch when neither side of the type ratio is 1, for example an 8-bit type against a 6-bit type. That branch passes layout, which resolves to the cute::layout overload set at include/cute/layout.hpp:558, and the call to upcast then has no viable candidate.

  // include/cute/swizzle_layout.hpp:455
- return downcast<scale::den>(upcast<scale::num>(layout));
+ static_assert(dependent_false<Swizzle<B,M,S>>,
+               "recast_layout of a Swizzle needs a power-of-two ratio between the two type sizes.");

  auto t = make_tensor(make_smem_ptr<cutlass::float_e2m3_t>(nullptr),
                       composition(Swizzle<3,3,3>{}, smem_ptr_flag_bits<6>{}, layout));
  as_position_independent_swizzle_tensor(t);

  before   73 diagnostic lines, ending at
           error: no matching function for call to
           'upcast<cute::R<6, 8>::num>(<unresolved overloaded function type>)'
           note: there are 6 candidates
  after    14 diagnostic lines, ending at
           error: static assertion failed: recast_layout of a Swizzle needs a
           power-of-two ratio between the two type sizes.

downcast<N> of a flagged layout divides the flag bit count by N with an integer division. smem_ptr_flag is smem_ptr_flag_bits<1>, so every GMMA *_Atom_Bits carries a flag of one bit and any N above 1 takes it to zero.

  // include/cute/pointer_flagged.hpp:83
+ static_assert(B % N == 0, "downcast<N> of a flagged layout needs N to divide the flag bit count.");

  GMMA::Layout_K_SW128_Atom_Bits
  atom          Sw<3,4,3> o smem_ptr[1b](unset) o (_8,_1024):(_1024,_1)
  downcast<2>   Sw<3,4,3> o smem_ptr[0b](unset) o (_8,_2048):(_2048,_1)
  downcast<8>   Sw<3,4,3> o smem_ptr[0b](unset) o (_8,_8192):(_8192,_1)
  after         rejected at compile time