                               int            axis) {
     
     switch( op ) {
-        case BF_REDUCE_POWER_MIN:   // Fall-through
-        case BF_REDUCE_POWER_MAX:    // Fall-through
-        case BF_REDUCE_POWER_SUM:    // Fall-through
-        case BF_REDUCE_POWER_MEAN:   // Fall-through
+        case BF_REDUCE_POWER_MIN:  // Fall-through
+        case BF_REDUCE_POWER_MAX:  // Fall-through
+        case BF_REDUCE_POWER_SUM:  // Fall-through
+        case BF_REDUCE_POWER_MEAN: // Fall-through
         case BF_REDUCE_POWER_STDERR: {
             switch( out->dtype ) {
                 case BF_DTYPE_F32:  return reduce_complex_power_itype_otype<IType>(in, out, op, axis);
                 default: BF_FAIL("Unsupported output dtype", BF_STATUS_UNSUPPORTED_DTYPE);
             }
         }
+        case BF_REDUCE_MIN: // Fall-through
+        case BF_REDUCE_MAX: BF_FAIL("Unsupported operation on input dtype", BF_STATUS_UNSUPPORTED);
         default: {
             switch( out->dtype ) {
                 case BF_DTYPE_CF32:  return reduce_complex_standard_itype_otype<IType>(in, out, op, axis);


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<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,
