Can this project support pytorch、tensorflow?

[jeremyyx]

Reading through the code, I found that the current project is currently unable to support gpu virtualisation for python as it has to be a cuda binary?

[n-eiling]

pytorch is not supported because pytorch loads the CUDA code as a shared object. I started implementing support for this, but it is not finished, yet.

[jeremyyx]

pytorch is not supported because pytorch loads the CUDA code as a shared object. I started implementing support for this, but it is not finished, yet.

Do you have any idea? I've been experimenting with this recently.

[n-eiling]

Once a shared object with CUDA code is loaded it calls this function (or something similiar):

cricket/cpu/cpu-client.c

Line 215 in c0dec62

void __cudaRegisterFunction(void **fatCubinHandle, const char *hostFun, char *deviceFun,

We cannot simply forward this call to the server because the shared object is not loaded in the server. So, before these calls we need to send the shared object to the server, load it and then forward the cudaRegisterFunction calls.
It's probably not that much work, but I am focusing on something different right now.

Hope this answers your questions.

[jeremyyx]

Once a shared object with CUDA code is loaded it calls this function (or something similiar):

cricket/cpu/cpu-client.c

Line 215 in c0dec62

void __cudaRegisterFunction(void **fatCubinHandle, const char *hostFun, char *deviceFun,

We cannot simply forward this call to the server because the shared object is not loaded in the server. So, before these calls we need to send the shared object to the server, load it and then forward the cudaRegisterFunction calls.
It's probably not that much work, but I am focusing on something different right now.
Hope this answers your questions.

Yes, I do find this API, and I test the GVirtuS repo. this repo can send some version of the shared object. However, not all cubin files are in the same format. In the GVirtuS project, he used the ELF parsing of fatcubin files to obtain the parameter characteristics of the cuda kernel, but older versions of the compiled ELF format could not be parsed properly.

[hudengjunai]

I find the GVirtuS project. Does GViruS support pytorch or tensorflow in shared lib load app?

[xial-thu]

I find the GVirtuS project. Does GViruS support pytorch or tensorflow in shared lib load app?

No, hijacking of cuda internal APIs require huge amount of reverse engineering, and you'll waste you life dealing with fatbinary. It is not that simple as GVirtus shows. I would suggest you hijack pytorch API for remote execution.

[nravic]

It's probably not that much work, but I am focusing on something different right now.

@n-eiling could you elaborate on what needs to be done? I'm interested in using the project for some orchestration of tensorflow/pytorch and would be happy to contribute to getting this done

[n-eiling]

Any contribution is highly welcome! A while back I looked into building support for pytorch. The issue is that in pytorch the cuda kernels are located in shared objects that are loaded at runtime using dlopen/dlsym. For shared objects CUDA inserts some initialization code that registers the new kernels with the CUDA APIs. Without running this initialization code we cannot call the kernels using cuLaunchKernel. If I remember correctly I already tacked down the initialization to __cudaRegisterFunction.

So to make pytorch work with cricket, we would need to detect loading of CUDA shared objects on the client side, then copy the shared object to the server and there load it also and register the kernels.

[nravic]

Makes sense to me, thanks! I'll give this a try and come back with questions.

Will have to reread your paper with more detail :)

[jin-zhengnan]

Once a shared object with CUDA code is loaded it calls this function (or something similiar):

cricket/cpu/cpu-client.c

Line 215 in c0dec62

void __cudaRegisterFunction(void **fatCubinHandle, const char *hostFun, char *deviceFun,

We cannot simply forward this call to the server because the shared object is not loaded in the server. So, before these calls we need to send the shared object to the server, load it and then forward the cudaRegisterFunction calls.

It's probably not that much work, but I am focusing on something different right now.
Hope this answers your questions.

I have been following your project for a long time and have some questions that I would like to inquire about. Please help me answer them when you have time. Thank you very much. Cuda has some hidden functions, for example__ cudaRegisterFatBinary，__ cudaRegisterFatBinaryEnd,__ PushCallConfiguration,__ PopCallConfiguration and so on, why not directly call these hidden functions on the server after client hijacking, but instead convert them through other methods similar to cuModuleLoadData?@n-eiling

[wzhao18]

Hi @jin-zhengnan, I recently experimented forwarding cudaRegisterFatBinary，__ cudaRegisterFatBinaryEnd,__ PushCallConfiguration,__ PopCallConfiguration calls from client applications to the server. And it is doable. You just have to intercept those calls and store the necessary data, such as the cubin data. And you repeat the same procedure in the server.

[n-eiling]

I also experimented with using these functions. These hidden functions are part of the runtime API, while I am using the cuModule APIs from the driver API. This is because I also use the driver API to launch kernels (cuLaunchKernekl) and, when I register the kernels using the hidden APIs I do not easily get the cuModule pointer required for this.

The hidden APIs seem to do basically the same as the cuModule APIs, however they register the host kernel stubs. These are the functions you call when launching a kernel in the runtime api (e.g., kernel<<<1, 1>>>()). We can't really intercept these, so instead Cricket intercepts the cuLaunchKernel call these stubs actually do.

I think working at the driver API level as much as possible is anyway desireable, as it reduces the server side code. So in a virtualized environment most code is executed in the VM or Container. Also the cuModule API is actually documented and hopefully more stable.

A few years back I also tried to implement Cricket entirely on the driver API level, so that the runtime API can also run on the client side. However, there are a lot more hidden functions called between both, which don't even show up in any public header. With more investigations this might be possible though...

[wzhao18]

hi @n-eiling, thanks for responding. Sorry I just found out your repo a few days ago and I haven't checked your implementation. I plan to read your paper soon.

I have a question regarding your answer. What do you mean that "runtime api (e.g., kernel<<<1, 1>>>()). We can't really intercept these"? Because when I check the compilation, it seems all kernel<<<1, 1>>> calls are in fact macro to launch the kernel using cudalaunchkernel. From my experiments, all kernels launches can be intercepted from cudalaunchkernel. In fact, I don't know when cuLaunchKernel are used. If you could provide some insights on the difference between cudalaunchkernel vs culaunchkernel, and what cuModule APIs are for, it will be very helpful for me.

Thanks in advance!

[n-eiling]

Ah okay. So cudaLaunchKernel requires as first parameter the kernel stub. These are not available on the server. We could register something using cudaRegisterFunction but it is not as simple as copying a bunch of memory to the server and calling cudaRegisterFunction on it. There are a lot of data structure being used when launching a kernel. I imagine CUDA is doing something similar as Cricket is doing (i.e. creating a list of available kernels and storing some useful information on them), because for launching a kernel we need to know what architecture the kernel is compile for, what parameters it has at which offsets, if it is compressed, etc.
I am also pretty sure that the Runtime API is actually also using the driver API to implement this .