 			                    BF_STATUS_UNSUPPORTED_STRIDE);
 			BF_ASSERT_EXCEPTION(out->strides[ndim-3] % out->strides[ndim-1] == 0,
 			                    BF_STATUS_UNSUPPORTED_STRIDE);
-			ibatchstride =  in->strides[ndim-3] / in->strides[ndim-1];
-			obatchstride = out->strides[ndim-3] / in->strides[ndim-1];
+			ibatchstride =  in->strides[ndim-3] /  in->strides[ndim-1];
+			obatchstride = out->strides[ndim-3] / out->strides[ndim-1];
 		}
 		
 		BF_ASSERT_EXCEPTION(in->strides[ndim-1] == BF_DTYPE_NBYTE(in->dtype),

		in  =  &in_flattened;
		BF_ASSERT(in_flattened.ndim == out_flattened.ndim,
		          BF_STATUS_INTERNAL_ERROR);
		BF_ASSERT(in_flattened.ndim == 3, BF_STATUS_UNSUPPORTED_SHAPE);
		BF_ASSERT_EXCEPTION(in_flattened.shape[0] == out_flattened.shape[0],
		                    BF_STATUS_INVALID_SHAPE);

	                &d_out,
	                &ostride,
	                &obatchstride};
	BF_CHECK_CUDA_EXCEPTION(
		cudaLaunchKernel((void*)fdmt_init_kernel<InType,OutType>,
		                 grid, block,
		                 &args[0], 0, stream),
		BF_STATUS_INTERNAL_ERROR);
}

template<typename DType>
