Как сделать Alea быстрее?
Проделав некоторую работу по реализации различных алгоритмов ML в Alea, я попытался сравнить некоторые простые, но важные процедуры в Alea. Меня удивило, что Alea'занимает примерно в 3 раза больше времени, чем эквивалентный вызов cuBLAS для sgeam. Если бы я делал что-то более сложное, например, умножение матриц, где мне пришлось бы манипулировать разделяемой памятью, это было бы понятно, но ниже приведены лишь простые преобразования массива.
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 работает для меня очень плохо, настолько плохо, что в том, как я его использую, должна быть ошибка. Это примерно в 20 раз медленнее, чем использование sgeamv дважды для суммирования матрицы.
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. Я думал, что флаг оптимизации мог быть выключен, но потом обнаружил, что он включен по умолчанию. Какие-нибудь советы по оптимизации, которые мне здесь не хватает?
1 ответ
Я думаю, что в вашем тестовом коде есть некоторые проблемы:
В вашем картографическом модуле вы должны предварительно загрузить модуль GPUModule. GPUModule JIT-компилируется при первом запуске. Так что на самом деле ваши измерения времени включают в себя время компиляции кода GPU;
В вашем модуле отображения, как в коде Alea, так и в коде cublas, вы должны синхронизировать работника (синхронизировать контекст CUDA). Программирование CUDA - асинхронный стиль. Поэтому, когда вы запускаете ядро, оно сразу же возвращается, не дожидаясь его завершения. Если вы не синхронизируете работника, на самом деле вы измеряете время запуска ядра, а не время выполнения ядра. Какое время запуска Alea GPU будет медленнее, чем собственный код C, так как он будет выполнять маршалинг аргументов ядра. Есть и другие проблемы, связанные со временем запуска ядра, которые я покажу вам в следующем примере кода.
Ваш тест на уменьшение фактически загружает модуль уменьшения каждый раз! Это означает, что каждый раз, когда вы делаете сокращение, вы измеряете время, включая время компиляции 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, и в настоящее время мы используем шаблон, в котором один графический процессор использует один выделенный поток, и к этому потоку присоединен контекст 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 = ()
Кроме того, как вы могли заметить, я запускаю этот тест дважды, и в первый раз тест без предварительно загруженного модуля использует 304 мс, но во второй раз тест без предварительно загруженного модуля использует только 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 = ()
Обратите внимание, что в этом тесте сокращения для каждого сокращения используется один сбор памяти (memcpyDtoH) для получения результата от устройства к хосту. и этот вызов API копирования памяти автоматически синхронизирует работника, потому что, если ядро не завершено, значение не имеет смысла. Поэтому, если вы хотите сравнить производительность с кодом C, вам также следует скопировать скаляр результата с устройства на хост. хотя это всего лишь один вызов API CUDA, но, как вы делали это для многих итераций (в этом примере 100), он будет накапливать некоторое время там.
Надеюсь, что это отвечает на ваш вопрос.