 template<int N, typename IType>
 __global__
 void reduce_complex_standard_vector_kernel(IType const* __restrict__ in,
-                                  float*       __restrict__ out,
-                                  int4         shape,
-                                  int4         istrides,
-                                  int4         ostrides,
-                                  BFreduce_op  op) {
+                                           float*       __restrict__ out,
+                                           int4         shape,
+                                           int4         istrides,
+                                           int4         ostrides,
+                                           BFreduce_op  op) {
 	int4 i0 = make_int4(threadIdx.x + blockIdx.x * blockDim.x,
 	                    threadIdx.y + blockIdx.y * blockDim.y,
 	                    threadIdx.z + blockIdx.z * blockDim.z,

template<typename IType>
__global__
void reduce_complex_standard_loop_kernel(IType const* __restrict__ in,
                                float*       __restrict__ out,
                                int4         shape,
                                int4         istrides,
                                int4         ostrides,
                                BFreduce_op  op) {
	int4 i0 = make_int4(threadIdx.x + blockIdx.x * blockDim.x,
	                    threadIdx.y + blockIdx.y * blockDim.y,
	                    threadIdx.z + blockIdx.z * blockDim.z,

	} else if( use_vec16_kernel ) {
		BF_TRY_RETURN(
			launch_reduce_vector_kernel<16>((IType*)in->data, (OType*)out->data,
			                                shape, istrides, ostrides,
			                                op, g_cuda_stream));
	} else {
		BF_TRY_RETURN(
			launch_reduce_loop_kernel((IType*)in->data, (OType*)out->data,
