Как правильно передавать аргументы как структуры в NVRTC?

Question

let prog =
    """//Kernel code:
extern "C" {
    #pragma pack(1)
    typedef struct {
        int length;
        float *pointer;
    } global_array_float;
    __global__ void kernel_main(global_array_float x){
        printf("(on device) x.length=%d\n",x.length); // prints: (on device) x.length=10
        printf("(on device) x.pointer=%lld\n",x.pointer); // prints: (on device) x.pointer=0
        printf("sizeof(global_array_float)=%d", sizeof(global_array_float)); // 12 bytes just as expected
    }
;}"""

printfn "%s" prog
let cuda_kernel = compile_kernel prog "kernel_main"

let test_launcher(str: CudaStream, kernel: CudaKernel, x: CudaGlobalArray<float32>, o: CudaGlobalArray<float32>) =
    let block_size = 1 

    kernel.GridDimensions <- dim3(1)
    kernel.BlockDimensions <- dim3(block_size)
    printfn "(on host) x.length=%i"  x.length // prints: (on host) x.length=10
    printfn "(on host) x.pointer=%i" x.pointer // prints: (on host) x.pointer=21535919104
    let args: obj [] = [|x.length;x.pointer|]
    kernel.RunAsync(str.Stream, args)

let cols, rows = 10, 1
let a = d2M.create((rows,cols)) 
        |> fun x -> fillRandomUniformMatrix ctx.Str x 1.0f 0.0f; x 
let a' = d2MtoCudaArray a

//printfn "%A" (getd2M a)

let o = d2M.create((rows,cols)) // o does nothing here as this is a minimalist example.
let o' = d2MtoCudaArray o

test_launcher(ctx.Str,cuda_kernel,a',o')
cuda_context.Synchronize()

//printfn "%A" (getd2M o)

Вот отрывок из основного репо, над которым я сейчас работаю. Я очень близок к тому, чтобы иметь рабочие цитаты F# для компилятора Cuda C, но я не могу понять, как правильно передать аргументы в функцию со стороны хоста.

Несмотря на прагму пакета, компилятор NVRTC 7.5 Cuda выполняет другую оптимизацию, и я понятия не имею, что это такое.

Поскольку я работаю с цитатами F#, мне нужно передать аргументы в виде единой структуры, чтобы это работало. Если я изменю функцию с kernel_main(global_array_float x) что-то вроде kernel_main(int x_length, float *x_pointer) тогда это работает, но я не являюсь той формой, которую система котировок дает мне заранее, и я хотел бы избежать дополнительной работы, чтобы сделать F# более похожим на C.

Есть идеи, что я могу попробовать?

Answer 1

Я сделал два ошибочных предположения.

Первая ошибка предполагает, что let args: obj [] = [|x.length;x.pointer|] будет аккуратно помещен в стек рядом друг с другом. На самом деле это два разных аргумента, а второй теряется где-то, как показано выше.

Это можно исправить, создав собственный тип структуры и переписав выражение следующим образом: let args: obj [] = [|CudaLocalArray(x.length,x.pointer)|],

Другое ошибочное предположение, которое я нашел, когда переписал его, как указано выше, заключается в том, что использование [<StructLayout(LayoutKind.Sequential>] не означает, что поля будут упакованы вместе. Вместо этого, как и для C, pack является аргументом, поэтому его нужно использовать так: [<StructLayout(LayoutKind.Sequential,Pack=1)>],