如何将参数作为结构正确地传递给 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=%dn",x.length); // prints: (on device) x.length=10
        printf("(on device) x.pointer=%lldn",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)

这是我目前正在处理的主存储库的摘录。我非常接近向 Cuda C 编译器提供有效的 F# 报价，但我无法弄清楚如何将参数从主机端正确传递到函数中。

尽管有包编译指示，但 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)>] 。