 		                    BF_STATUS_UNSUPPORTED_STRIDE);
 		BF_ASSERT_EXCEPTION(out->strides[ndim-2] % out->strides[ndim-1] == 0,
 		                    BF_STATUS_UNSUPPORTED_STRIDE);
-		size_t istride =  in->strides[ndim-2] / in->strides[ndim-1];
-		size_t ostride = out->strides[ndim-2] / in->strides[ndim-1];
+		size_t istride =  in->strides[ndim-2] /  in->strides[ndim-1];
+		size_t ostride = out->strides[ndim-2] / out->strides[ndim-1];
 		BF_ASSERT_EXCEPTION( in->strides[in->ndim-2] > 0, BF_STATUS_UNSUPPORTED_STRIDE);
 		BF_ASSERT_EXCEPTION(out->strides[in->ndim-2] > 0, BF_STATUS_UNSUPPORTED_STRIDE);
 		

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

template<typename DType>

			                    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] / out->strides[ndim-1];
		}
		
		BF_ASSERT_EXCEPTION(in->strides[ndim-1] == BF_DTYPE_NBYTE(in->dtype),
