引数を構造体として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) 

私が現在取り組んでいるmain repoの抜粋です。私はCuda Cコンパイラに作業中のF＃を引用することに非常に近いですが、ホスト側から引数を関数に正しく渡す方法を理解することはできません。

パック・プラグマにもかかわらず、NVRTC 7.5 Cudaコンパイラは他の最適化を行っていますが、それは何か分かりません。

私はF＃の引用を取り組んでいるので、これを機能させるには、引数を単一の構造体として渡す必要があります。関数をkernel_main(global_array_float x)からkernel_main(int x_length, float *x_pointer)に変更すると、それは動作しますが、引用システムが私に前払いする形式ではないので、F＃をCのようにするために余分な作業をしないでください。

Any私は何を試すことができるアイデア？

Answer 1

私は2つの誤った仮定をしました。

最初のエラーは、let args: obj [] = [|x.length;x.pointer|]がお互いに隣に積み重ねられていると仮定しています。実際にはこれらは2つの異なる議論であり、2つ目は上記のように渡されるとどこかで失われます。

これは、カスタム構造体型を作成し、次のように式を書き換えることで修正できます。let args: obj [] = [|CudaLocalArray(x.length,x.pointer)|]。

私が上記のように書き直したときに見つけたもう1つの誤った仮定は、[<StructLayout(LayoutKind.Sequential>]を使用しても、フィールドが一緒にパックされるわけではありません。代わりに、Cの場合と同様に、packは引数なので、それを使用する必要があります：[<StructLayout(LayoutKind.Sequential,Pack=1)>]。