[FEA] Support for (u)int8 matrix in row_normalize

[enp1s0]

Overflow can occur when T = (u)int8 in the row_normalize function because the accumulator data type is also T.

raft/cpp/include/raft/linalg/detail/normalize.cuh

Lines 38 to 63 in da3b9a9

	RAFT_KERNEL __launch_bounds__(Policy::ThreadsPerBlock)
	coalesced_normalize_thin_kernel(Type* out,
	const Type* in,
	IdxType D,
	IdxType N,
	Type init,
	MainLambda main_op,
	ReduceLambda reduce_op,
	FinalLambda fin_op,
	Type eps)
	{
	IdxType i = threadIdx.y + (Policy::RowsPerBlock * static_cast<IdxType>(blockIdx.x));
	if (i >= N) return;
	Type acc = init;
	for (IdxType j = threadIdx.x; j < D; j += Policy::LogicalWarpSize) {
	Type val = in[j + D * i];
	acc = reduce_op(acc, main_op(val, j));
	}
	acc = raft::logicalWarpReduce<Policy::LogicalWarpSize>(acc, reduce_op);
	acc = fin_op(acc);
	if (acc <= eps) return;
	for (IdxType j = threadIdx.x; j < D; j += Policy::LogicalWarpSize) {
	out[j + D * i] = in[j + D * i] / acc;
	}
	}

Required in this PR: #2287 (comment)

modifications

• Update accumulator data type (e.g. int32_t, float)
  • Template parameter?
• Add scaling method (e.g. [scale](int8_t elm, int32_t norm){return min(INT8_MAX, max(INT8_MIN, elm * scale / norm));})
  • This is required since adjusting the norm to 1 is not appropriate for integer matrices.
  • Overflow avoidance is also required.