 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>
BFstatus reduce_complex_standard_itype_otype(BFarray const* in,
                                    BFarray const* out,
                                    BFreduce_op    op,
                                    int            axis) {
	BF_ASSERT(in->shape[axis] % out->shape[axis] == 0, BF_STATUS_INVALID_SHAPE);
	long reduce_size    =  in->shape[axis] / out->shape[axis];
	long istride_reduce =  in->strides[axis];


template<int N, typename IType>
void launch_reduce_complex_standard_vector_kernel(IType const* in,
                                                  float*       out,
                                                  int4         shape,
                                                  int4         istrides,
                                                  int4         ostrides,
                                                  BFreduce_op  op,
                                                  cudaStream_t stream) {
	dim3 block(128); // TODO: Tune this
	dim3 grid(std::min((shape.x - 1) / block.x + 1, 65535u),
	          std::min(shape.y, 65535),
