split_k hit IMA error

[Adnios]

I encountered the following error when try split_k_slices(2) on V100. As #73 (comment) suggested, I use cudaGetErrorString to dump error message.

an illegal memory access was encountered
terminate called after throwing an instance of 'std::runtime_error'
  what():  Failed to copy equality flag to device.
Aborted

cutlass/tools/util/include/cutlass/util/reference/device/tensor_compare.h

Lines 188 to 195 in af332d4

	if (cudaMemcpy(
	device_equal_flag,
	&equal_flag,
	sizeof(int),
	cudaMemcpyHostToDevice) != cudaSuccess) {
	throw std::runtime_error("Failed to copy equality flag to device.");
	}

M=N=K=1776, FP16, split_k_slices=2

./tools/profiler/cutlass_profiler --operation=Gemm --split_k_slices=2 --m=1776 --n=1776 --k=1776 --A=f16:col --B=f16:row --C=f16:col --alpha=1 --beta=0 --cta_m=256 --cta_n=128 --cta_k=32 --stages=2

split_k_slices=1 can pass, but split_k_slices=2 hit error. Decrease mnk size also can pass.

I'm confused why >1776 will hit IMA.

update:

I tried the dev branch and the problem disappeared. It appears that this issue only occurs in cutlass-2.5.