Overview of User Defined Functions with cuDF¶
Like many tabular data processing APIs, cuDF provides a range of composable, DataFrame style operators. While out of the box functions are flexible and useful, it is sometimes necessary to write custom code, or user-defined functions (UDFs), that can be applied to rows, columns, and other groupings of the cells making up the DataFrame.
In conjunction with the broader GPU PyData ecosystem, cuDF provides interfaces to run UDFs on a variety of data structures. Currently, we can only execute UDFs on numeric and Boolean typed data (support for strings is being planned). This guide covers writing and executing UDFs on the following data structures:
Series
DataFrame
Rolling Windows Series
Groupby DataFrames
CuPy NDArrays
Numba DeviceNDArrays
It also demonstrates cuDF’s default null handling behavior, and how to write UDFs that can interact with null values in a limited fashion.
Overview¶
When cuDF executes a UDF, it gets just-in-time (JIT) compiled into a CUDA kernel (either explicitly or implicitly) and is run on the GPU. Exploring CUDA and GPU architecture in-depth is out of scope for this guide. At a high level:
Compute is spread across multiple “blocks”, which have access to both global memory and their own block local memory
Within each block, many “threads” operate independently and simultaneously access their block-specific shared memory with low latency
This guide covers APIs that automatically handle dividing columns into chunks and assigning them into different GPU blocks for parallel computation (see apply_chunks or the numba CUDA JIT API if you need to control this yourself).
Series UDFs¶
You can execute UDFs on Series in two ways:
Writing a standard Python function and using
applymap
Writing a Numba kernel and using Numba’s
forall
syntax
Using applymap
is simpler, but writing a Numba kernel offers the flexibility to build more complex functions (we’ll be writing only simple kernels in this guide).
Let’s start by importing a few libraries and creating a DataFrame of several Series.
[1]:
import numpy as np
import cudf
from cudf.datasets import randomdata
df = randomdata(nrows=10, dtypes={'a':float, 'b':bool, 'c':str}, seed=12)
df.head()
[1]:
a | b | c | |
---|---|---|---|
0 | -0.691674 | True | Dan |
1 | 0.480099 | False | Bob |
2 | -0.473370 | True | Xavier |
3 | 0.067479 | True | Alice |
4 | -0.970850 | False | Sarah |
Next, we’ll define a basic Python function and call it as a UDF with applymap
.
[2]:
def udf(x):
if x > 0:
return x + 5
else:
return x - 5
[3]:
df['a'].applymap(udf)
[3]:
0 -5.691674
1 5.480099
2 -5.473370
3 5.067479
4 -5.970850
5 5.837494
6 5.801430
7 -5.933157
8 5.913899
9 -5.725581
Name: a, dtype: float64
That’s all there is to it. For more complex UDFs, though, we’d want to write an actual Numba kernel.
For more complex logic (for instance, accessing values from multiple input columns or rows, you’ll need to use a more complex API. There are several types. First we’ll cover writing and running a Numba JITed CUDA kernel.
The easiest way to write a Numba kernel is to use cuda.grid(1)
to manage our thread indices, and then leverage Numba’s forall
method to configure the kernel for us. Below, define a basic multiplication kernel as an example and use @cuda.jit
to compile it.
[4]:
from numba import cuda
@cuda.jit
def multiply(in_col, out_col, multiplier):
i = cuda.grid(1)
if i < in_col.size: # boundary guard
out_col[i] = in_col[i] * multiplier
This kernel will take an input array, multiply it by a configurable value (supplied at runtime), and store the result in an output array. Notice that we wrapped our logic in an if
statement. Because we can launch more threads than the size of our array, we need to make sure that we don’t use threads with an index that would be out of bounds. Leaving this out can result in undefined behavior.
To execute our kernel, we just need to pre-allocate an output array and leverage the forall
method mentioned above. First, we create a Series of all 0.0
in our DataFrame, since we want float64
output. Next, we run the kernel with forall
. forall
requires us to specify our desired number of tasks, so we’ll supply in the length of our Series (which we store in size
). The **cuda_array_interface** is
what allows us to directly call our Numba kernel on our Series.
[5]:
size = len(df['a'])
df['e'] = 0.0
multiply.forall(size)(df['a'], df['e'], 10.0)
After calling our kernel, our DataFrame is now populated with the result.
[6]:
df.head()
[6]:
a | b | c | e | |
---|---|---|---|---|
0 | -0.691674 | True | Dan | -6.916743 |
1 | 0.480099 | False | Bob | 4.800994 |
2 | -0.473370 | True | Xavier | -4.733700 |
3 | 0.067479 | True | Alice | 0.674788 |
4 | -0.970850 | False | Sarah | -9.708501 |
Note that, while we’re operating on the Series df['e']
, the kernel executes on the DeviceNDArray “underneath” the Series. If you ever need to access the underlying DeviceNDArray of a Series, you can do so with Series.data.mem
. We’ll use this during an example in the Null Handling section of this guide.
DataFrame UDFs¶
We could apply a UDF on a DataFrame like we did above with forall
. We’d need to write a kernel that expects multiple inputs, and pass multiple Series as arguments when we execute our kernel. Because this is fairly common and can be difficult to manage, cuDF provides two APIs to streamline this: apply_rows
and apply_chunks
. Below, we walk through an example of using apply_rows
. apply_chunks
works in a similar way, but also offers more control over low-level kernel behavior.
Now that we have two numeric columns in our DataFrame, let’s write a kernel that uses both of them.
[7]:
def conditional_add(x, y, out):
for i, (a, e) in enumerate(zip(x, y)):
if a > 0:
out[i] = a + e
else:
out[i] = a
Notice that we need to enumerate
through our zipped
function arguments (which either match or are mapped to our input column names). We can pass this kernel to apply_rows
. We’ll need to specify a few arguments: - incols - A list of names of input columns that match the function arguments. Or, a dictionary mapping input column names to their corresponding function arguments such as {'col1': 'arg1'}
. - outcols - A dictionary defining our output column names and their data types.
These names must match our function arguments. - kwargs (optional) - We can optionally pass keyword arguments as a dictionary. Since we don’t need any, we pass an empty one.
While it looks like our function is looping sequentially through our columns, it actually executes in parallel in multiple threads on the GPU. This parallelism is the heart of GPU-accelerated computing. With that background, we’re ready to use our UDF.
[8]:
df = df.apply_rows(conditional_add,
incols={'a':'x', 'e':'y'},
outcols={'out': np.float64},
kwargs={}
)
df.head()
[8]:
a | b | c | e | out | |
---|---|---|---|---|---|
0 | -0.691674 | True | Dan | -6.916743 | -0.691674 |
1 | 0.480099 | False | Bob | 4.800994 | 5.281093 |
2 | -0.473370 | True | Xavier | -4.733700 | -0.473370 |
3 | 0.067479 | True | Alice | 0.674788 | 0.742267 |
4 | -0.970850 | False | Sarah | -9.708501 | -0.970850 |
As expected, we see our conditional addition worked. At this point, we’ve successfully executed UDFs on the core data structures of cuDF.
Rolling Window UDFs¶
For time-series data, we may need to operate on a small “window” of our column at a time, processing each portion independently. We could slide (“roll”) this window over the entire column to answer questions like “What is the 3-day moving average of a stock price over the past year?”
We can apply more complex functions to rolling windows to rolling
Series and DataFrames using apply
. This example is adapted from cuDF’s API documentation. First, we’ll create an example Series and then create a rolling
object from the Series.
[9]:
ser = cudf.Series([16, 25, 36, 49, 64, 81], dtype='float64')
ser
[9]:
0 16.0
1 25.0
2 36.0
3 49.0
4 64.0
5 81.0
dtype: float64
[10]:
rolling = ser.rolling(window=3, min_periods=3, center=False)
rolling
[10]:
Rolling [window=3,min_periods=3,center=False]
Next, we’ll define a function to use on our rolling windows. We created this one to highlight how you can include things like loops, mathematical functions, and conditionals. Rolling window UDFs do not yet support null values.
[11]:
import math
def example_func(window):
b = 0
for a in window:
b = max(b, math.sqrt(a))
if b == 8:
return 100
return b
We can execute the function by passing it to apply
. With window=3
, min_periods=3
, and center=False
, our first two values are null
.
[12]:
rolling.apply(example_func)
[12]:
0 null
1 null
2 6.0
3 7.0
4 100.0
5 9.0
dtype: float64
We can apply this function to every column in a DataFrame, too.
[13]:
df2 = cudf.DataFrame()
df2['a'] = np.arange(55, 65, dtype='float64')
df2['b'] = np.arange(55, 65, dtype='float64')
df2.head()
[13]:
a | b | |
---|---|---|
0 | 55.0 | 55.0 |
1 | 56.0 | 56.0 |
2 | 57.0 | 57.0 |
3 | 58.0 | 58.0 |
4 | 59.0 | 59.0 |
[14]:
rolling = df2.rolling(window=3, min_periods=3, center=False)
rolling.apply(example_func)
[14]:
a | b | |
---|---|---|
0 | null | null |
1 | null | null |
2 | 7.549834435 | 7.549834435 |
3 | 7.615773106 | 7.615773106 |
4 | 7.681145748 | 7.681145748 |
5 | 7.745966692 | 7.745966692 |
6 | 7.810249676 | 7.810249676 |
7 | 7.874007874 | 7.874007874 |
8 | 7.937253933 | 7.937253933 |
9 | 100.0 | 100.0 |
GroupBy DataFrame UDFs¶
We can also apply UDFs to grouped DataFrames using apply_grouped
. This example is also drawn and adapted from the RAPIDS API documentation.
First, we’ll group our DataFrame based on column b
, which is either True or False. Note that we currently need to pass method="cudf"
to use UDFs with GroupBy objects.
[15]:
df.head()
[15]:
a | b | c | e | out | |
---|---|---|---|---|---|
0 | -0.691674 | True | Dan | -6.916743 | -0.691674 |
1 | 0.480099 | False | Bob | 4.800994 | 5.281093 |
2 | -0.473370 | True | Xavier | -4.733700 | -0.473370 |
3 | 0.067479 | True | Alice | 0.674788 | 0.742267 |
4 | -0.970850 | False | Sarah | -9.708501 | -0.970850 |
[16]:
grouped = df.groupby(['b'], method="cudf")
/opt/conda/envs/rapids/lib/python3.7/site-packages/cudf/core/dataframe.py:2559: UserWarning: as_index==True not supported due to the lack of multi-index with legacy groupby function. Use hash method for multi-index
"as_index==True not supported due to the lack of "
Next we’ll define a function to apply to each group independently. In this case, we’ll take the rolling average of column e
, and call that new column rolling_avg_e
.
[17]:
def rolling_avg(e, rolling_avg_e):
win_size = 3
for i in range(cuda.threadIdx.x, len(e), cuda.blockDim.x):
if i < win_size - 1:
# If there is not enough data to fill the window,
# take the average to be NaN
rolling_avg_e[i] = np.nan
else:
total = 0
for j in range(i - win_size + 1, i + 1):
total += e[j]
rolling_avg_e[i] = total / win_size
We can execute this with a very similar API to apply_rows
. This time, though, it’s going to execute independently for each group.
[18]:
results = grouped.apply_grouped(rolling_avg,
incols=['e'],
outcols=dict(rolling_avg_e=np.float64))
results
[18]:
a | b | c | e | out | rolling_avg_e | |
---|---|---|---|---|---|---|
0 | 0.480099 | False | Bob | 4.800994 | 5.281093 | NaN |
1 | -0.970850 | False | Sarah | -9.708501 | -0.970850 | NaN |
2 | 0.801430 | False | Sarah | 8.014297 | 8.815727 | 1.035597 |
3 | -0.933157 | False | Quinn | -9.331571 | -0.933157 | -3.675258 |
4 | -0.691674 | True | Dan | -6.916743 | -0.691674 | NaN |
5 | -0.473370 | True | Xavier | -4.733700 | -0.473370 | NaN |
6 | 0.067479 | True | Alice | 0.674788 | 0.742267 | -3.658552 |
7 | 0.837494 | True | Wendy | 8.374940 | 9.212434 | 1.438676 |
8 | 0.913899 | True | Ursula | 9.138987 | 10.052885 | 6.062905 |
9 | -0.725581 | True | George | -7.255814 | -0.725581 | 3.419371 |
Notice how, with a window size of three in the kernel, the first two values in each group for our output column are null.
Numba Kernels on CuPy Arrays¶
We can also execute Numba kernels on CuPy NDArrays, again thanks to the __cuda_array_interface__
. We can even run the same UDF on the Series and the CuPy array. First, we define a Series and then create a CuPy array from that Series.
[19]:
import cupy as cp
s = cudf.Series([1.0, 2, 3, 4, 10])
arr = cp.asarray(s)
arr
[19]:
array([ 1., 2., 3., 4., 10.])
Next, we define a UDF and execute it on our Series. We need to allocate a Series of the same size for our output, which we’ll call out
.
[20]:
from cudf.utils import cudautils
@cuda.jit
def multiply_by_5(x, out):
i = cuda.grid(1)
if i < x.size:
out[i] = x[i] * 5
out = cudf.Series(cudautils.zeros(len(s), dtype='int32'))
multiply_by_5.forall(s.shape[0])(s, out)
out
[20]:
0 5
1 10
2 15
3 20
4 50
dtype: int32
Finally, we execute the same function on our array. We allocate an empty array out
to store our results.
[21]:
out = cp.empty_like(arr)
multiply_by_5.forall(arr.size)(arr, out)
out
[21]:
array([ 5., 10., 15., 20., 50.])
Null Handling in UDFs¶
Above, we covered most basic usage of UDFs with cuDF.
The remainder of the guide focuses on considerations for executing UDFs on DataFrames containing null values. If your UDFs will read or write any column containing nulls, you should read this section carefully.
Writing UDFs that can handle null values is complicated by the fact that a separate bitmask is used to identify when a value is valid and when it’s null. By default, DataFrame methods for applying UDFs like apply_rows
will handle nulls pessimistically (all rows with a null value will be removed from the output if they are used in the kernel). Exploring how not handling not pessimistically can lead to undefined behavior is outside the scope of this guide. Suffice it to say, pessimistic null
handling is the safe and consistent approach. You can see an example below.
[22]:
def gpu_add(a, b, out):
for i, (x, y) in enumerate(zip(a, b)):
out[i] = x + y
df = randomdata(nrows=5, dtypes={'a':int, 'b':int, 'c':int}, seed=12)
df.loc[2, 'a'] = None
df.loc[3, 'b'] = None
df.loc[1, 'c'] = None
df.head()
[22]:
a | b | c | |
---|---|---|---|
0 | 963 | 1005 | 997 |
1 | 977 | 1026 | null |
2 | null | 1026 | 1019 |
3 | 1078 | null | 985 |
4 | 979 | 982 | 1011 |
In the dataframe above, there are three null values. Each column has a null in a different row. When we use our UDF with apply_rows
, our output should have two nulls due to pessimistic null handling (because we’re not using column c
, the null value there does not matter to us).
[23]:
df = df.apply_rows(gpu_add,
incols=['a', 'b'],
outcols={'out':np.float64},
kwargs={})
df.head()
[23]:
a | b | c | out | |
---|---|---|---|---|
0 | 963 | 1005 | 997 | 1968.0 |
1 | 977 | 1026 | null | 2003.0 |
2 | null | 1026 | 1019 | null |
3 | 1078 | null | 985 | null |
4 | 979 | 982 | 1011 | 1961.0 |
As expected, we end up with two nulls in our output. The null values from the columns we used propogated to our output, but the null from the column we ignored did not.
Operating on Null Values¶
If you don’t need to conditionally handle null values in your UDFs, feel free to skip these final two sections.
As a developer or data scientist, you may sometimes need to write UDFs that operate on null values. This means you need to think about the null bitmask array when writing your UDF. As a note, cuDF allows you to turn off pessimistic null handling in apply_rows
. Instead of doing this, if you need to operate on null values we recommend writing standard Numba.cuda
kernels. To help you interact with null bitmasks from Python, cuDF provides the mask_get
utility function. The following
example illustrates how you can use mask_get
in Numba kernels like we used earlier in this guide.
Standard Numba Kernels¶
First, we import mask_get
and create a DataFrame with some null values.
[24]:
from cudf.utils.cudautils import mask_get
df = randomdata(nrows=10, dtypes={'a':float, 'b':bool}, seed=12)
df.loc[[2,4], 'a'] = None
df.head()
[24]:
a | b | |
---|---|---|
0 | -0.691674315 | True |
1 | 0.480099393 | False |
2 | null | True |
3 | 0.067478787 | True |
4 | null | False |
Next, we’ll define a simple kernel like before, with a couple of differences. This kernel needs access to the null bitmask, so we include a validity_mask
argument. We also wrap our logic in a conditional based on the results of mask_get
: - If the result of mask_get
for that index is valid (there is a value), do the multiplication - If the result of mask_get
for that index is not valid (it’s null), set the output -999999
[25]:
@cuda.jit
def gpu_kernel_masked(in_col, validity_mask, out_col, multiplier):
i = cuda.grid(1)
if i < in_col.size:
valid = mask_get(validity_mask, i)
if valid:
out_col[i] = in_col[i] * multiplier
else:
out_col[i] = -999999
We now grab the underlying DeviceArrays and execute our kernel like we did previously, except that this time we also pass in the DeviceArray of our column’s null mask. Because Numba doesn’t yet handle masked GPU arrays, we can’t directly pass our Series
here.
[26]:
import rmm # RAPIDS Memory Manager
a_dary = df.a._column.data.mem
a_mask = df.a.nullmask.mem
output_dary = rmm.device_array_like(a_dary)
gpu_kernel_masked.forall(output_dary.size)(a_dary, a_mask, output_dary, 10)
df['result'] = output_dary
df.head()
[26]:
a | b | result | |
---|---|---|---|
0 | -0.691674315 | True | -6.916743 |
1 | 0.480099393 | False | 4.800994 |
2 | null | True | -999999.000000 |
3 | 0.067478787 | True | 0.674788 |
4 | null | False | -999999.000000 |
Summary¶
This guide has covered a lot of content. At this point, you should hopefully feel comfortable writing UDFs (with or without null values) that operate on
Series
DataFrame
Rolling Windows
GroupBy DataFrames
CuPy NDArrays
Numba DeviceNDArrays
For more information please see the cuDF, Numba.cuda, and CuPy documentation.