如何获取通用类型的默认值（用于 warp shuffle）？

Question

        let t = if warpid = 0 then mean.[i / num_rows] else (Unchecked.defaultof<'T>)
        __syncthreads()
        let v = __shfl t 0 32

我想获取 'T 的默认值，但由于 Unchecked.defaultof<'T>，上面的代码片段给出了编译错误。在 Alea 中进行 warp shuffle 的首选方法是什么？

现在我遇到了一个问题，许多线程从同一个位置读取一次，我正在尝试测试是否只让第一个线程从那个位置读取然后将值洗牌到经线中的其他人。（编辑：一点也不。缓存工作得很好。）

Answer 1

支持Unchecked.defaultof是个好主意，我会检查一下，谢谢。

目前有两种方法可以获取'T类型的默认值：

使用Alea.CUDA.Intrinsic.__default_value<'T>()（参见here）。 Intrinsic是一个自动打开的模块，所以如果你打开了Alea.CUDA命名空间，你可以直接在你的代码中使用__default_value()。
第二种方法是打开命名空间Alea.CUDA.Utilities，并使用自动打开的NumericLiteralG模块（参见here），然后在你的内联泛型函数中你可以直接写0G, 1G, 等等..

对于你的第二个问题，我贴了一些helper warp shuffle类型的源代码，其中包括广播的使用。这些辅助静态方法在模块 Alea.CUDA.Intrinsic:

中可用

///A helper static class providing shuffle instructions.
[<AbstractClass;Sealed>]
type WarpShuffle private () =

    [<ReflectedDefinition>]
    static member Broadcast(input:'T, srcLane:int, width:int) =
        __shfl input srcLane width

    [<ReflectedDefinition>]
    static member Broadcast(input:'T, srcLane:int) =
        let width = __warp_size()
        __shfl input srcLane width

    [<ReflectedDefinition>]
    static member Up(input:'T, delta:int, width:int) =
        __shfl_up input delta width

    [<ReflectedDefinition>]
    static member Up(input:'T, delta:int) =
        let width = __warp_size()
        __shfl_up input delta width

    [<ReflectedDefinition>]
    static member Down(input:'T, delta:int, width:int) =
        __shfl_down input delta width

    [<ReflectedDefinition>]
    static member Down(input:'T, delta:int) =
        let width = __warp_size()
        __shfl_down input delta width

    [<ReflectedDefinition>]
    static member Xor(input:'T, laneMask:int, width:int) =
        __shfl_xor input laneMask width

    [<ReflectedDefinition>]
    static member Xor(input:'T, laneMask:int) =
        let width = __warp_size()
        __shfl_xor input laneMask width

///[omit]
[<AbstractClass;Sealed>]
type FullWarpShuffle private () =

    [<ReflectedDefinition>]
    static member Broadcast(input:'T, srcLane:int, logicWarpThreads:int) =
        let shflC = logicWarpThreads - 1
        __shfl_raw input srcLane shflC

    [<ReflectedDefinition>]
    static member Broadcast(input:'T, srcLane:int) =
        let shflC = __warp_size() - 1
        __shfl_raw input srcLane shflC

    [<ReflectedDefinition>]
    static member Up(input:'T, srcOffset:int) =
        let shflC = 0
        __shfl_up_raw input srcOffset shflC

    [<ReflectedDefinition>]
    static member Down(input:'T, srcOffset:int) =
        let shflC = __warp_size() - 1
        __shfl_down_raw input srcOffset shflC

    [<ReflectedDefinition>]
    static member Down(input:'T, srcOffset:int, warpThreads:int) =
        let shflC = warpThreads - 1
        __shfl_down_raw input srcOffset shflC

在上面的代码中，它使用了 __shf_raw，在线文档已过时。这是 ptx 代码 shfl.idx 的原始版本，其中 shflC 包含两个打包值，指定用于逻辑上将 warp 拆分为子段的掩码和用于限制源通道索引的上限。在 here.

阅读更多内容

如何获取通用类型的默认值（用于 warp shuffle）？

How to get the default value of a generic type (for the warp shuffle)?

aleagpu