Passing arbitrary struct arguments by value to kernels with StrPack?

[ikirill]

If I understand correctly, I can use StrPack to ensure that a Julia value of a composite type can be converted to/from a string of bytes consistent with the binary representation expected in C code. I'm not too sure, but IIUC just passing a pointer to a Julia object is not (?) necessarily safe.

But I can't figure out how to launch a kernel using the binary representation of a Julia composite type. In other words, (but this is clearly wrong, because it can't be distinguished from passing a host pointer to a kernel, and I also get a totally different error below):

@struct type A; x :: Cint; end
iostr = IOBuffer(); pack(iostr, A(1))
...
launch(..., (iostr.data,))

where in C code:

struct A { int x; }
__global__ void kernel_fun(A a);

I looked in execute.jl, and it seems it's not implemented: https://github.com/JuliaGPU/CUDArt.jl/blob/master/src/execute.jl#L2 and I get an error that rawpointer is undefined for Array{UInt8,1}, which is what is in IOBuffer data.

I think whenever a value of a composite type is passed to a kernel, it might make sense to pass a pointer to its binary representation to cuLaunchKernel, so that the argument gets passed by value to the kernel.

[I should add] that calling cuLaunchKernel with ccall and constructing kernel arguments myself seems to work fine so far (although I haven't tested everything very much yet).

[timholy]

Presumably you first have to move the memory over to the GPU. Does it work if you say launch(..., (CudaArray(iostr.data),))?

[ikirill]

No, I don't. If I understand correctly from documentation, cuLaunchKernel copies the arguments itself when those arguments are passed by value. For example, if you pass it a pointer to a device pointer, the device pointer will be copied and used to initialize the function argument, and if you give it a pointer to an int, the int will be copied by cuLaunchKernel, then used to initialize an argument of type int (not int* or something). So all I need is to be able to give it a pointer to the structure in host memory, and it copies the structure itself, initializing function argument from the host memory.

It's just that the interface of launch doesn't let me do that, hence the issue. The way to get around is to ccall((:cuLaunchKernel, ... with kernel arguments set to Ptr{Void}[iostr.data].

Also, I believe arguments are copied to either constant or shared memory (I'm not sure which), not global memory.

[timholy]

With Ptr{Void}, how does it know how many bytes to copy over? Different structures have different sizes.

[ikirill]

It knows the kernel function definition, and the definitions of types of its arguments, like a C compiler would, so that's how it knows how many bytes to copy and how to interpret them. So if I tell StrPack to use the right (for nvcc) alignment (align_native? It seems distinct from align_default for some reason), then the structure is laid out in memory correctly and cuda can read it itself without being told anything else.

[timholy]

You're right that the kernel function itself is a good source of information. If that really works, fantastic.

It's certainly fine to change the interface of launch if it expands the scope of what it can do. Once you get something working that passes the tests, please do submit a PR.