Wie übergeben Sie Argumente als Strukturen an 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) 

Hier ist ein Auszug aus dem main repo, an dem ich gerade arbeite. Ich bin sehr nah daran, eine funktionierende F # -Angaben zu Cuda C-Compiler, aber ich kann nicht herausfinden, wie die Argumente in die Funktion ordnungsgemäß von der Host-Seite übergeben.

Trotz der Pack Pragma, der NVRTC 7.5 Cuda Compiler macht einige andere Optimierung und ich habe keine Ahnung, was es ist.

Da ich von F # -Angebote arbeite, muss ich die Argumente als eine einzige Struktur übergeben, damit dies funktioniert. Wenn ich die Funktion von kernel_main(global_array_float x) zu so etwas wie kernel_main(int x_length, float *x_pointer) dann funktioniert es ändern, aber ich, dass ist nicht die Form, die die Notierungen System gibt mir im Voraus, und ich möchte dabei zusätzliche Arbeit vermeiden F # mehr wie C.

Beliebig zu machen Idee was ich ausprobieren könnte?

Answer 1

Ich habe zwei falsche Annahmen gemacht.

Der erste Fehler geht davon aus, dass let args: obj [] = [|x.length;x.pointer|] sauber nebeneinander auf den Stapel gelegt werden. In Wirklichkeit sind das zwei verschiedene Argumente und der zweite geht irgendwo verloren, wenn er wie oben weitergereicht wird.

Es kann behoben werden, indem Sie einen benutzerdefinierten Strukturtyp erstellen und den Ausdruck wie folgt umschreiben: let args: obj [] = [|CudaLocalArray(x.length,x.pointer)|]. Die andere falsche Annahme, die ich fand, als ich es wie oben beschrieben umformte, ist, dass die Verwendung von [<StructLayout(LayoutKind.Sequential>] nicht bedeutet, dass die Felder zusammen gepackt werden. Stattdessen ist pack wie für C ein Argument, also muss es wie folgt verwendet werden: [<StructLayout(LayoutKind.Sequential,Pack=1)>].