1

Alea でさまざまな ML アルゴリズムを実装するいくつかの作業を行ったので、Alea でいくつかの単純だが不可欠なルーチンのベンチマークを試みました。Alea' は、同じことを行う sgeam への同等の cuBLAS 呼び出しよりも約 3 倍の時間がかかることを知って驚きました。共有メモリを操作しなければならない行列乗算のようなもっと複雑なことをしていたら、これは理解できたでしょうが、以下は単純な配列変換です。

let dmat = createRandomUniformMatrix 100 1000 1.0f 0.0f
let dmat2 = createRandomUniformMatrix 100 1000 1.0f 0.0f
let rmat = createEmptyMatrixLike dmat

let m = new DeviceUnaryTransformModule<float32> <@ fun x -> x*2.0f @>

#time
//4.85s/100k
for i=1 to 100000 do
    m.Apply(dmat, rmat) |> ignore
#time

#time
//1.8s/100k
for i=1 to 100000 do
    sgeam2 nT nT 2.0f dmat 0.0f dmat2 rmat  |> ignore
#time

DeviceUnaryTransformModule 変換モジュールのカーネルは、基本的な変換の例と同じです。唯一の違いは、後でホストにデータを収集する代わりに、デバイス上にデータを保持することです。

また、Unbound の reduce は、私にとっては非常にうまく機能しません。行列を合計するために sgeamv を 2 回使用するよりも約 20 倍遅くなります。

let makeReduce (op:Expr<'T -> 'T -> 'T>)  =
    let compileReductionKernel (op:Expr<'T -> 'T -> 'T>) =
        worker.LoadProgram(
                        DeviceReduceImpl.DeviceReduce(op, worker.Device.Arch, PlatformUtil.Instance.ProcessBitness).Template
                        )

    let prog = compileReductionKernel op

    let runReduceProgram (sumProg : Program<DeviceReduceImpl.IDeviceReduceFactory<'A>>) (x: DeviceMemory<'A>) = 
        sumProg.Entry.Create(blob, x.Length)
               .Reduce(None, x.Ptr, x.Length)

    let reduceProg (x: DeviceMemory<'T>) = runReduceProgram prog x
    reduceProg

let sumReduce: DeviceMemory<float32> -> float32 = makeReduce <@ fun (a:float32) b -> a + b @>

#time
//3.5s/10k
for i=1 to 10000 do
    sumReduce dmat.dArray |> ignore
#time

これを CUDA C++ と比較しようとしたことはありませんが、単純なことでは cuBLAS と同等であると思います。最適化フラグがオフになっているのではないかと思ったのですが、デフォルトでオンになっていることがわかりました。ここで見逃している最適化のヒントはありますか?

4

1 に答える 1

3

テストコードにいくつかの問題があると思います：

  1. マッピング モジュールでは、GPUModule をプリロードする必要があります。GPUModule は、初回起動時に JIT コンパイルされます。したがって、実際のタイミング測定には、GPU コードのコンパイル時間が含まれます。

  2. マッピング モジュールでは、Alea コードと cublas コードの両方で、ワーカーを同期する (CUDA コンテキストを同期する) 必要があります。CUDA プログラミングは非同期スタイルです。そのため、カーネルを起動すると、カーネルが完了するのを待たずにすぐに戻ります。ワーカーを同期しない場合、実際にはカーネルの実行時間ではなく、カーネルの起動時間を測定しています。どの Alea gpu の起動時間は、カーネル引数のマーシャリングを行うため、ネイティブ C コードよりも遅くなります。カーネルの起動時間に関連するその他の問題がいくつかあります。次のサンプル コードで示します。

  3. reduce テストは実際に毎回 reduce モジュールをロードします! つまり、リダクションを行うたびに、GPU のコンパイル時間を含めて時間を測定します! GPU モジュールまたはプログラムのインスタンスは、コンパイルされた GPU コードを表しているため、長寿命にすることをお勧めします。

それで、私はあなたの使用法に従ってテストを行いました。ここでは、最初に完全なテスト コードを示します。

#r @"packages\Alea.CUDA.2.1.2.3274\lib\net40\Alea.CUDA.dll"
#r @"packages\Alea.CUDA.IL.2.1.2.3274\lib\net40\Alea.CUDA.IL.dll"
#r @"packages\Alea.CUDA.Unbound.2.1.2.3274\lib\net40\Alea.CUDA.Unbound.dll"
#r "System.Configuration"
open System.IO
Alea.CUDA.Settings.Instance.Resource.AssemblyPath <- Path.Combine(@"packages\Alea.CUDA.2.1.2.3274", "private")
Alea.CUDA.Settings.Instance.Resource.Path <- Path.GetTempPath()

open Alea.CUDA
open Alea.CUDA.Utilities
open Alea.CUDA.CULib
open Alea.CUDA.Unbound
open Microsoft.FSharp.Quotations

type MapModule(target, op:Expr<float32 -> float32>) =
    inherit GPUModule(target)

    [<Kernel;ReflectedDefinition>]
    member this.Kernel (C:deviceptr<float32>) (A:deviceptr<float32>) (B:deviceptr<float32>) (n:int) =
        let start = blockIdx.x * blockDim.x + threadIdx.x
        let stride = gridDim.x * blockDim.x
        let mutable i = start
        while i < n do
            C.[i] <- __eval(op) A.[i] + __eval(op) B.[i]
            i <- i + stride

    member this.Apply(C:deviceptr<float32>, A:deviceptr<float32>, B:deviceptr<float32>, n:int) =
        let lp = LaunchParam(64, 256)
        this.GPULaunch <@ this.Kernel @> lp C A B n

let inline mapTemplate (op:Expr<'T -> 'T>) = cuda {
    let! kernel = 
        <@ fun (C:deviceptr<'T>) (A:deviceptr<'T>) (B:deviceptr<'T>) (n:int) ->
            let start = blockIdx.x * blockDim.x + threadIdx.x
            let stride = gridDim.x * blockDim.x
            let mutable i = start
            while i < n do
                C.[i] <- (%op) A.[i] + (%op) B.[i]
                i <- i + stride @>
        |> Compiler.DefineKernel

    return Entry(fun program ->
        let worker = program.Worker
        let kernel = program.Apply kernel
        let lp = LaunchParam(64, 256)

        let run C A B n =
            kernel.Launch lp C A B n

        run ) }

let test1 (worker:Worker) m n sync iters =
    let n = m * n
    use m = new MapModule(GPUModuleTarget.Worker(worker), <@ fun x -> x * 2.0f @>)
    let rng = System.Random(42)
    use A = worker.Malloc(Array.init n (fun _ -> rng.NextDouble() |> float32))
    use B = worker.Malloc(Array.init n (fun _ -> rng.NextDouble() |> float32))
    use C = worker.Malloc<float32>(n)
    let timer = System.Diagnostics.Stopwatch.StartNew()
    for i = 1 to iters do
        m.Apply(C.Ptr, A.Ptr, B.Ptr, n)
    if sync then worker.Synchronize()
    timer.Stop()
    printfn "%f ms / %d %s (no pre-load module)" timer.Elapsed.TotalMilliseconds iters (if sync then "sync" else "nosync")

let test2 (worker:Worker) m n sync iters =
    let n = m * n
    use m = new MapModule(GPUModuleTarget.Worker(worker), <@ fun x -> x * 2.0f @>)
    // we pre-load the module, this will JIT compile the GPU code
    m.GPUForceLoad()
    let rng = System.Random(42)
    use A = worker.Malloc(Array.init n (fun _ -> rng.NextDouble() |> float32))
    use B = worker.Malloc(Array.init n (fun _ -> rng.NextDouble() |> float32))
    use C = worker.Malloc<float32>(n)
    let timer = System.Diagnostics.Stopwatch.StartNew()
    for i = 1 to iters do
        m.Apply(C.Ptr, A.Ptr, B.Ptr, n)
    if sync then worker.Synchronize()
    timer.Stop()
    printfn "%f ms / %d %s (pre-loaded module)" timer.Elapsed.TotalMilliseconds iters (if sync then "sync" else "nosync")

let test3 (worker:Worker) m n sync iters =
    let n = m * n
    use m = new MapModule(GPUModuleTarget.Worker(worker), <@ fun x -> x * 2.0f @>)
    // we pre-load the module, this will JIT compile the GPU code
    m.GPUForceLoad()
    let rng = System.Random(42)
    use A = worker.Malloc(Array.init n (fun _ -> rng.NextDouble() |> float32))
    use B = worker.Malloc(Array.init n (fun _ -> rng.NextDouble() |> float32))
    use C = worker.Malloc<float32>(n)
    // since the worker is running in a background thread
    // each cuda api will switch to that thread
    // use eval() to avoid the many thread switching
    worker.Eval <| fun _ ->
        let timer = System.Diagnostics.Stopwatch.StartNew()
        for i = 1 to iters do
            m.Apply(C.Ptr, A.Ptr, B.Ptr, n)
        if sync then worker.Synchronize()
        timer.Stop()
        printfn "%f ms / %d %s (pre-loaded module + worker.eval)" timer.Elapsed.TotalMilliseconds iters (if sync then "sync" else "nosync")

let test4 (worker:Worker) m n sync iters =
    use program = worker.LoadProgram(mapTemplate <@ fun x -> x * 2.0f @>)
    let n = m * n
    let rng = System.Random(42)
    use A = worker.Malloc(Array.init n (fun _ -> rng.NextDouble() |> float32))
    use B = worker.Malloc(Array.init n (fun _ -> rng.NextDouble() |> float32))
    use C = worker.Malloc<float32>(n)
    let timer = System.Diagnostics.Stopwatch.StartNew()
    for i = 1 to iters do
        program.Run C.Ptr A.Ptr B.Ptr n
    if sync then worker.Synchronize()
    timer.Stop()
    printfn "%f ms / %d %s (template usage)" timer.Elapsed.TotalMilliseconds iters (if sync then "sync" else "nosync")

let test5 (worker:Worker) m n sync iters =
    use program = worker.LoadProgram(mapTemplate <@ fun x -> x * 2.0f @>)
    let n = m * n
    let rng = System.Random(42)
    use A = worker.Malloc(Array.init n (fun _ -> rng.NextDouble() |> float32))
    use B = worker.Malloc(Array.init n (fun _ -> rng.NextDouble() |> float32))
    use C = worker.Malloc<float32>(n)
    worker.Eval <| fun _ ->
        let timer = System.Diagnostics.Stopwatch.StartNew()
        for i = 1 to iters do
            program.Run C.Ptr A.Ptr B.Ptr n
        if sync then worker.Synchronize()
        timer.Stop()
        printfn "%f ms / %d %s (template usage + worker.Eval)" timer.Elapsed.TotalMilliseconds iters (if sync then "sync" else "nosync")

let test6 (worker:Worker) m n sync iters =
    use cublas = new CUBLAS(worker)
    let rng = System.Random(42)
    use dmat1 = worker.Malloc(Array.init (m * n) (fun _ -> rng.NextDouble() |> float32))
    use dmat2 = worker.Malloc(Array.init (m * n) (fun _ -> rng.NextDouble() |> float32))
    use dmatr = worker.Malloc<float32>(m * n)
    let timer = System.Diagnostics.Stopwatch.StartNew()
    for i = 1 to iters do
        cublas.Sgeam(cublasOperation_t.CUBLAS_OP_N, cublasOperation_t.CUBLAS_OP_N, m, n, 2.0f, dmat1.Ptr, m, 2.0f, dmat2.Ptr, m, dmatr.Ptr, m)
    if sync then worker.Synchronize()
    timer.Stop()
    printfn "%f ms / %d %s (cublas)" timer.Elapsed.TotalMilliseconds iters (if sync then "sync" else "nosync")

let test7 (worker:Worker) m n sync iters =
    use cublas = new CUBLAS(worker)
    let rng = System.Random(42)
    use dmat1 = worker.Malloc(Array.init (m * n) (fun _ -> rng.NextDouble() |> float32))
    use dmat2 = worker.Malloc(Array.init (m * n) (fun _ -> rng.NextDouble() |> float32))
    use dmatr = worker.Malloc<float32>(m * n)
    worker.Eval <| fun _ ->
        let timer = System.Diagnostics.Stopwatch.StartNew()
        for i = 1 to iters do
            cublas.Sgeam(cublasOperation_t.CUBLAS_OP_N, cublasOperation_t.CUBLAS_OP_N, m, n, 2.0f, dmat1.Ptr, m, 2.0f, dmat2.Ptr, m, dmatr.Ptr, m)
        if sync then worker.Synchronize()
        timer.Stop()
        printfn "%f ms / %d %s (cublas + worker.eval)" timer.Elapsed.TotalMilliseconds iters (if sync then "sync" else "nosync")

let test worker m n sync iters =
    test6 worker m n sync iters
    test7 worker m n sync iters
    test1 worker m n sync iters
    test2 worker m n sync iters
    test3 worker m n sync iters
    test4 worker m n sync iters
    test5 worker m n sync iters

let testReduce1 (worker:Worker) n iters =
    let rng = System.Random(42)
    use input = worker.Malloc(Array.init n (fun _ -> rng.NextDouble() |> float32))
    use reduceModule = new DeviceReduceModule<float32>(GPUModuleTarget.Worker(worker), <@ (+) @>)
    // JIT compile and load GPU code for this module
    reduceModule.GPUForceLoad()
    // create a reducer which will allocate temp memory for maxNum=n
    let reduce = reduceModule.Create(n)
    let timer = System.Diagnostics.Stopwatch.StartNew()
    for i = 1 to 10000 do
        reduce.Reduce(input.Ptr, n) |> ignore
    timer.Stop()
    printfn "%f ms / %d (pre-load gpu code)" timer.Elapsed.TotalMilliseconds iters

let testReduce2 (worker:Worker) n iters =
    let rng = System.Random(42)
    use input = worker.Malloc(Array.init n (fun _ -> rng.NextDouble() |> float32))
    use reduceModule = new DeviceReduceModule<float32>(GPUModuleTarget.Worker(worker), <@ (+) @>)
    // JIT compile and load GPU code for this module
    reduceModule.GPUForceLoad()
    // create a reducer which will allocate temp memory for maxNum=n
    let reduce = reduceModule.Create(n)
    worker.Eval <| fun _ ->
        let timer = System.Diagnostics.Stopwatch.StartNew()
        for i = 1 to 10000 do
            reduce.Reduce(input.Ptr, n) |> ignore
        timer.Stop()
        printfn "%f ms / %d (pre-load gpu code and avoid thread switching)" timer.Elapsed.TotalMilliseconds iters

let testReduce worker n iters =
    testReduce1 worker n iters
    testReduce2 worker n iters

let workerDefault = Worker.Default
let workerNoThread = Worker.CreateOnCurrentThread(Device.Default)

Alea GPU ではワーカーが CUDA コンテキストを表し、現在、1 つの GPU が 1 つの専用スレッドを使用し、そのスレッドに CUDA コンテキストがアタッチされるというパターンを使用しています。これを「専用スレッドを持つワーカー」と呼びます。つまり、カーネルの起動など、CUDA API を呼び出すたびに、ワーカー スレッドに切り替える必要があります。カーネルの起動を頻繁に行う場合は、Worker.Eval関数を使用してワーカー スレッド内でコードを実行し、スレッドの切り替えを回避することをお勧めします。現在のスレッドでワーカーを作成する実験的な機能もあります。これにより、スレッドの切り替えが回避されますが、この使用法はまだ最適化されています。詳しくはこちらをご参照ください

ここで、最初にデフォルトのワーカーを使用して、ワーカーを同期せずにテストを行います (つまり、カーネルの起動時間のみを比較していることになります)。デフォルトのワーカーは専用スレッドを持つワーカーなので、 を使用するとパフォーマンスが向上することがわかりますWorker.Eval。ただし、全体として、.net からのカーネルの起動は、ネイティブ C カーネルの起動よりも遅くなります。

> test workerDefault 10000 10000 false 100;;
4.487300 ms / 100 nosync (cublas)
0.560600 ms / 100 nosync (cublas + worker.eval)
304.427900 ms / 100 nosync (no pre-load module)
18.517000 ms / 100 nosync (pre-loaded module)
12.579100 ms / 100 nosync (pre-loaded module + worker.eval)
27.023800 ms / 100 nosync (template usage)
16.007500 ms / 100 nosync (template usage + worker.Eval)
val it : unit = ()
> test workerDefault 10000 10000 false 100;;
3.288600 ms / 100 nosync (cublas)
0.647300 ms / 100 nosync (cublas + worker.eval)
29.129100 ms / 100 nosync (no pre-load module)
18.874700 ms / 100 nosync (pre-loaded module)
12.285000 ms / 100 nosync (pre-loaded module + worker.eval)
20.452300 ms / 100 nosync (template usage)
14.903500 ms / 100 nosync (template usage + worker.Eval)
val it : unit = ()

また、お気付きかもしれませんが、このテストを 2 回実行しました。1 回目は、プリロード モジュールを使用しないテストで 304 ミリ秒を使用しますが、2 回目は、プリロード モジュールを使用しないテストで 29 ミリ秒しか使用しません。その理由は、LLVM P/Invoke を使用してカーネルをコンパイルするためです。また、これらの P/Invoke 関数は遅延関数であるため、最初に使用するときにいくつかの初期化が行われ、その後は高速になります。

ここで、実際のカーネル実行時間を実際に測定したワーカーを同期します。ここで作成したカーネルは非常に単純ですが、行列 A と B の両方で動作します。

> test workerDefault 10000 10000 true 100;;
843.695000 ms / 100 sync (cublas)
841.452400 ms / 100 sync (cublas + worker.eval)
919.244900 ms / 100 sync (no pre-load module)
912.348000 ms / 100 sync (pre-loaded module)
908.909000 ms / 100 sync (pre-loaded module + worker.eval)
914.834100 ms / 100 sync (template usage)
914.170100 ms / 100 sync (template usage + worker.Eval)

スレッドレスワーカーでテストすると、スレッド切り替えがないため、少し高速になります。

> test workerNoThread 10000 10000 true 100;;
842.132100 ms / 100 sync (cublas)
841.627200 ms / 100 sync (cublas + worker.eval)
918.007800 ms / 100 sync (no pre-load module)
908.575900 ms / 100 sync (pre-loaded module)
908.770100 ms / 100 sync (pre-loaded module + worker.eval)
913.405300 ms / 100 sync (template usage)
913.942600 ms / 100 sync (template usage + worker.Eval)

ここで、削減に対するテストを示します。

> testReduce workerDefault 10000000 100;;
7691.335300 ms / 100 (pre-load gpu code)
6448.782500 ms / 100 (pre-load gpu code and avoid thread switching)
val it : unit = ()
> testReduce workerNoThread 10000000 100;;
6467.105300 ms / 100 (pre-load gpu code)
6426.296900 ms / 100 (pre-load gpu code and avoid thread switching)
val it : unit = ()

この縮小テストでは、デバイスからホストへの結果を取得するために、縮小ごとに 1 つのメモリ収集 (memcpyDtoH) があることに注意してください。このメモリ コピー API の呼び出しによって、ワーカーが自動的に同期されます。これは、カーネルが完了していない場合、値が無意味になるためです。したがって、パフォーマンスを C コードと比較したい場合は、結果のスカラーもデバイスからホストにコピーする必要があります。これは 1 つの CUDA API 呼び出しに過ぎませんが、多くの反復 (この例では 100 回) で行ったように、そこで何らかのタイミングが蓄積されます。

これがあなたの質問に答えることを願っています。

于 2015-08-24T06:07:06.823 に答える