Make kernel global
@@ -1997,7 +1997,20 @@ cpdef ndarray concatenate(tup, axis, shape, dtype):
x = array([a.data.ptr for a in tup])
axis_sizes = array([a.shape[axis] for a in tup], 'i')
x_strides = array([a.strides for a in tup], 'i')
- kernel = ElementwiseKernel(
+ _concatenate_kernel(
+ x, axis, axis_sizes, x_strides, array(shape, 'i'), base, ret)
+ else:
+ skip = (slice(None),) * axis
+ i = 0
+ for a in tup:
+ aw = a.shape[axis]
+ ret[skip + (slice(i, i + aw),)] = a
+ i += aw
+
+ return ret
+
+
+cdef _concatenate_kernel = ElementwiseKernel(
'''raw P x, int32 axis, raw int32 axis_sizes, raw int32 x_strides,
raw int32 shape, int32 base''',
'T y',
@@ -2030,17 +2043,6 @@ cpdef ndarray concatenate(tup, axis, shape, dtype):
''',
'cupy_concatenate'
)
- kernel(
- x, axis, axis_sizes, x_strides, array(shape, 'i'), base, ret)
- else:
- skip = (slice(None),) * axis
- i = 0
- for a in tup:
- aw = a.shape[axis]
- ret[skip + (slice(i, i + aw),)] = a
- i += aw
-
- return ret
# -----------------------------------------------------------------------------
