search

rapidsai / rapids-examples

33 stars 24 forks source link

Explore pycuda as an option to write custom cuda code with cudf #10

Closed VibhuJawa closed 3 years ago

VibhuJawa commented 3 years ago

Explore pycuda as an option to write custom cuda code with cudf

With a recent PR pycuda got __cuda_array_interface__ allowing interoperability with numba,cupy etc. Pycuda in theory allows us access to the whole cuda driver API. It might be worth exploring how we can use it with cudf.

Having examples that do operations like https://github.com/rapidsai/rapids-examples/tree/main/shareable-dataframes but with Pycuda might make access to CUDA kernels easier without too much glue code.

An example to create will be something on the lines @jdye64 did but with pycuda

Helpful links:

  1. https://documen.tician.de/pycuda/tutorial.html#executing-a-kernel
  2. https://documen.tician.de/pycuda/tutorial.html?highlight=numba#interoperability-with-other-libraries-using-the-cuda-array-interface

CC: @beckernick , @randerzander , @jdye64 .

randerzander commented 3 years ago

Would the CUDA code strings be JIT compiled with pycuda?

If so, is there a cache so that you only pay the compilation tax the first time you run?

VibhuJawa commented 3 years ago

Would the CUDA code strings be JIT compiled with pycuda?

Yup, they would be JIT compiled, cached as well as loaded to the device when we call pycuda.compiler.SourceModule('''Cuda code striing''') so it should be a one time cost.