reorder kernel args

python/cudf/cudf/core/frame.py

@@ -1562,7 +1562,7 @@ def _apply(self, func, *args):
         # Mask and data column preallocated
         ans_col = cupy.empty(len(self), dtype=retty)
         ans_mask = cudf.core.column.column_empty(len(self), dtype="bool")
-        launch_args = [(ans_col, ans_mask)]
+        launch_args = [(ans_col, ans_mask), len(self)]
         offsets = []
         for col in self._data.values():
             data = col.data
@@ -1574,7 +1574,6 @@ def _apply(self, func, *args):
             offsets.append(col.offset)
         launch_args += offsets
         launch_args += list(args)
-        launch_args.append(len(self))  # size
         kernel.forall(len(self))(*launch_args)
 
         result = cudf.Series(ans_col).set_mask(

python/cudf/cudf/core/udf/pipeline.py

@@ -111,13 +111,13 @@ def construct_signature(df, return_type, args):
     # Tuple of arrays, first the output data array, then the mask
     return_type = Tuple((return_type[::1], boolean[::1]))
     offsets = []
-    sig = [return_type]
+    sig = [return_type, int64]
     for col in df._data.values():
         sig.append(masked_array_type_from_col(col))
         offsets.append(int64)
 
-    # return_type + data,masks + offsets + size + extra args
-    sig = void(*(sig + offsets + [int64] + [typeof(arg) for arg in args]))
+    # return_type, size, data, masks, offsets, extra args
+    sig = void(*(sig + offsets + [typeof(arg) for arg in args]))
 
     return sig
 
@@ -128,7 +128,7 @@ def mask_get(mask, pos):
 
 
 kernel_template = """\
-def _kernel(retval, {input_columns}, {input_offsets}, {extra_args} size):
+def _kernel(retval, size, {input_columns}, {input_offsets}, {extra_args}):
     i = cuda.grid(1)
     ret_data_arr, ret_mask_arr = retval
     if i < size:
@@ -201,11 +201,7 @@ def _kernel(retval, input_col_0, input_col_1, offset_0, offset_1, size):
     # Create argument list for kernel
     input_columns = ", ".join([f"input_col_{i}" for i in range(len(fr._data))])
     input_offsets = ", ".join([f"offset_{i}" for i in range(len(fr._data))])
-    extra_args = (
-        ", ".join([f"extra_arg_{i}" for i in range(len(args))]) + ","
-        if len(args) > 0
-        else ""
-    )
+    extra_args = ", ".join([f"extra_arg_{i}" for i in range(len(args))])
 
     # Generate the initializers for each device function argument
     initializers = []