KernelTuner / kernel_tuner

Kernel Tuner
https://kerneltuner.github.io/kernel_tuner/
Apache License 2.0
285 stars 49 forks source link

Populating pointers of struct from cupy #211

Closed neworderofjamie closed 11 months ago

neworderofjamie commented 1 year ago

Not really an issue but, I'm trying to populate a struct in constant memory containing (device) pointers in order to run my kernel through kernel tuner. In C the struct looks something like:

struct MergedNeuronUpdateGroup0
 {
    unsigned int* spkCnt;
    unsigned int* spk;
};
__device__ __constant__ MergedNeuronUpdateGroup0 d_mergedNeuronUpdateGroup0[8];

I can see how I would pack pointers into this with struct.pack("PP0l", spk_cnt, spk) and then convert to numpy as described in your example before passing via the cmem_args kwarg to tune_kernel. But, I've no idea how I would create (and initialise) these pointers from Python! Any pointers?

Thanks in advance for your help!

benvanwerkhoven commented 1 year ago

I'm not sure how to do this from cupy, but with pycuda it should be easy. If you have PyCuda device allocation you can just cast it to an int to get the GPU memory address. From a cupy array you may be able to extract the pointer somehow, but I just had a look at cupy and DLpack, but I'm not sure how to extract the device pointer from those data types. If using PyCuda is an option for you, I think you could make it work.

neworderofjamie commented 1 year ago

I haven't used either PyCuda or cupy so happy to use whichever is easier - thanks for the suggestion