Alea CUDA DeviceMapModule гораздо медленнее, чем написанное от руки ядро

Question

Я попытался использовать DeviceMapModule и DeviceMap2Module для простых операций сопоставления. К моему удивлению, это примерно в 2 раза медленнее, чем написание ядра вручную. Для написанного от руки ядра я ничего особенного не делал. Я просто скопировал ядро ​​Getting Started Square и изменил его, чтобы сделать сигмоид.

Ниже приведены 2 версии кода:

Сначала медленная версия:

type SigmoidModule(size) =
    inherit Module(size, size) // Note Module is my own abstract class that defines the forward/backward interface.

    [<ReflectedDefinition; AOTCompile>]
    let sigmoid' s y = (1.0f-s) * s * y

    [<ReflectedDefinition; AOTCompile>]
    let sigmoid x = 1.0f / (1.0f + __nv_expf (-x))

    let forwardMap = new DeviceMapModule<float32, float32>(GPUModuleTarget.DefaultWorker, <@ sigmoid @>)
    let backwardMap = new DeviceMap2Module<float32, float32, float32>(
                        GPUModuleTarget.DefaultWorker, <@ sigmoid' @>)

    let mutable activation = Matrix.ZeroCreate 1 size

    override m.Forward (fwd:Matrix) = 
        assert(fwd.Cols = size)
        if activation.Rows <> fwd.Rows then
            activation.Data.Dispose()
            activation <- Matrix.ZeroCreate fwd.Rows fwd.Cols

        forwardMap.Map(activation.Data.Ptr, fwd.Data.Ptr, fwd.Cols * fwd.Rows)
        activation

    override m.Backward (dy:Matrix) = 
        assert(dy.Cols = size)
        assert(activation.Rows = dy.Rows)

        backwardMap.Map(activation.Data.Ptr, activation.Data.Ptr, dy.Data.Ptr, dy.Cols * dy.Rows)
        activation

    interface IDisposable with
        member m.Dispose() = 
            forwardMap.Dispose()
            backwardMap.Dispose()
            activation.Data.Dispose()

Теперь быстрая версия:

type SigmoidModuleFast(size) =
inherit Module(size, size)
let sigmoid' s y = (1.0f-s) * s * y
let worker = Worker.Default

[<ReflectedDefinition; AOTCompile>]
static let sigmoidKernel (outputs:deviceptr<float32>) (inputs:deviceptr<float32>) n =
    let start = blockIdx.x * blockDim.x + threadIdx.x
    let stride = gridDim.x * blockDim.x
    let sigmoid x = 1.0f / (1.0f + __nv_expf (-x))

    let mutable i = start 
    while i < n do
        outputs.[i] <- sigmoid(inputs.[i])
        i <- i + stride

[<ReflectedDefinition; AOTCompile>]
static let sigmoidPrimeKernel (outputs:deviceptr<float32>) (input:deviceptr<float32>) (dy:deviceptr<float32>) n =
    let start = blockIdx.x * blockDim.x + threadIdx.x
    let stride = gridDim.x * blockDim.x

    let mutable i = start 
    while i < n do
        let s = input.[i]
        let y = dy.[i]
        outputs.[i] <- (1.0f-s) * s * y
        i <- i + stride

let mutable activation = Matrix.ZeroCreate 1 size
let mutable lp = LaunchParam(1, size)

override m.Forward (fwd:Matrix) = 
    assert(fwd.Cols = size)
    if activation.Rows <> fwd.Rows then
        activation.Data.Dispose()
        activation <- Matrix.ZeroCreate fwd.Rows fwd.Cols
        let threads = fwd.Rows * fwd.Cols
        if threads < 1024 then
            lp <- LaunchParam(1, threads)
        else
            let blockSize = 256
            let numSm = worker.Device.Attributes.MULTIPROCESSOR_COUNT
            let gridSize = Math.Min(16 * numSm, divup threads blockSize)
            lp <- new LaunchParam(gridSize, blockSize)

    worker.Launch <@ sigmoidKernel @> lp activation.Data.Ptr fwd.Data.Ptr (fwd.Cols*fwd.Rows)
    activation

override m.Backward (dy:Matrix) = 
    assert(dy.Cols = size)
    assert(activation.Rows = dy.Rows)

    worker.Launch <@ sigmoidPrimeKernel @> lp activation.Data.Ptr activation.Data.Ptr dy.Data.Ptr (dy.Cols*dy.Rows)
    activation

interface IDisposable with
    member m.Dispose() = 
        activation.Data.Dispose()

Вот некоторая информация о моей системе:

• Название Intel(R) Core(TM) i5-4590 CPU @ 3,30 ГГц
• Архитектура x64
• Частота 3300 МГц
• Количество ядер 4
• Размер страницы 4,096
• Общая физическая память 8 050,00 МБ
• Доступная физическая память 902.00 МБ
• Гибридная графика включена False
• Название версии Windows 10 Pro
• Номер версии 10.0.10586
• Nsight версия 5.1.0.16062

• Visual Studio версии 14.0

• Графический процессор: NVidia GTX 980 Ti

• .NET CLR:.NET 4.6.

Answer 1

Я не специалист по программированию на GPU, но имею базовое понимание. Я видел этот вопрос во время оценки Alea GPU.

NVidia GTX 980 Ti имеет 2816 ядер. С blockSize 256, что дает размер сетки 2816 / 256 = 11.

Библиотека Alea Community/Basic Developer поддерживает до 2688 ядер. Если библиотека Alea устанавливает blockSize 256 (это предположение), это дает gridSize 10,5.

Может ли быть так, что работа делится на две части? Первый прогон насыщает предел ядра 2688, а затем остальные работают на 2816 - 2688 = 128 ядер?