ReductionKernel broken in versions > 2021.2.11

[prabhuramachandran]

Description

The ReductionKernel is broken with this commit: https://github.com/inducer/pyopencl/commit/fa408abd0bfb4787dac0c5bcc4931d33fd7705e8

Basically, the commit assumes that the first argument to the kernel is an array, this was not the case earlier and this breaks existing code which did not make this assumption.

To Reproduce

Here is a small example that works with 2021.2.11 but fails with anything higher.

import numpy
import pyopencl as cl
import pyopencl.array as array
from pyopencl.reduction import ReductionKernel

ctx = cl.create_some_context()
queue = cl.CommandQueue(ctx)

a = array.arange(queue, 400, dtype=numpy.float32)
b = array.arange(queue, 400, dtype=numpy.float32)

krnl = ReductionKernel(ctx, numpy.float32, neutral="0",
        reduce_expr="a+b", map_expr="z*x[i]*y[i]",
        arguments="float z, __global float *x, __global float *y")

my_dot_prod = krnl(0.1, a, b).get()
print(my_dot_prod)

Expected behavior

I am not sure if the implicit assumption that the first argument is an array is correct, it does not appear to be documented and certainly introduces a regression in code that used to work.

Environment (please complete the following information):

This fails on any OS but I see test failures on MacOS and Linux.

[inducer]

Thanks for the report, see #536.