F# и пример расчета на GPU.

В статье рассматривается проект обработки изображения на массиве процессоров. В качестве такого массива используется видеокарта с поддержкой технологии CUDA от компании NVIDIA. В качестве технологии, обеспечивающей доступ к видеокарте, применяется библиотека Alea.cuBase от компании Quantalea. В качестве языка программирования - F Sharp. Ну, а в качесте инструмента разработки Microsoft Visual Studio 2012.
Известно, что графические процессоры могут содержать тысячи ядер, способных работать параллельно. Известно так-же, что графические процессоры могут выполнять не только обработку графических данных. Однако начнем именно с графики. Возьмем простую и конкретную задачу - увеличение контрастности изображения, имеющего, к напримеру, 10000 х 5000 пикселей. Вся задача выглядит следующим образом.Считаем с диска Jpg файл, увелим его контрастность на определенный коэффициeнт, затем модифицированный файл запишем на диск под другим именем, по ходу дела подсчитаем время, затрачиваемое на обработку изображения, и полное время работы программы. Поступим так дважды. Сначала выполним все на CPU, а затем на GPU выполним модификацию изображения, а все остальное на CPU. Ну и, конечно, проанализируем, что у нас получилоcь. Проект можно свободно загрузить.
1. Что требуется для запуска проекта.
Перечислим состав аппаратных и программных средств, необходимых для запуска проекта:

1. На компьютере должна быть установлена видеокарта NVIDIA, поддерживающая технологию CUDA.
2. Должен быть установлен драйвер этой видиокарты.
3. Необходимо загрузить библиотеку Alea.cuBase от компании Quantalea.
4. Должена быть установлена среда разработки. Автор использовал Visual Studio 2012 под Windows 7.

2. Обновление изображения на CPU.
Прежде чем приступить к обработке изображения, необходимо выполнить несколько подготовительных операций. Все они помещены в функцию readImageBytes(), перечислим их (смотрите текст программы):

1. Создаем новый объект типв Bitmap со ссылкой на указанный файл изображения.
2. Создаем новый объект типа Rectangle с размерами файла изображения.
3. Заблокируем созданный объект Bitmap.
4. Вычислим размер изображения в байтах.
6. Создадим байтовый массив imageBytes.
7. Скопируем данные из bitmapData в массив imageBytes.

Отметим, что эта функция, с таким же множеством шагов выполняется и для варианта расчета на GPU. Далее запускается функция updateImageCPU(), в которой собственно и происходит процесс обновления изображения. После обновления выполняются функции по сохранению изображения в другом файле:

1. Копируем данные обратно из imageBytes в bitmapData
2. Разблокируем изображение.
3. Записываем данные в выходной файл.
4. Остановливаем подсчет времени.

Отметим, что совершенно аналогичные функции выполняются и при окончании работы GPU.
Code
// 25.05.2014 Программа увеличения контрастности в JPG файле. Как пример
// расчета на массиве процессоров.
// Применен язык F#, технология CUDA от NVIDIA и Alea.cuBase от Quantalea.

open System
open System.IO
open System.Drawing
open System.Drawing.Imaging
open System.Runtime.InteropServices
open Microsoft.FSharp.Quotations
open Alea.CUDA
open Alea.CUDA.Utilities

// Имя входного файла
let fileNameImage       = @"Test.jpg"
// Имя выходного файла
let outFileNameImage    = @"Test_out.jpg"
// Среднее значение контрастности
let original    = 128
// Коэффициент увеличения контрастности
let amount      = 3

// Функция ядра GPU
let pfunct = cuda {
    let! kernel =
        <@ fun n (x:deviceptr<byte>)
                (z:deviceptr<byte>) ->
            let start = blockIdx.x * blockDim.x + threadIdx.x
            let stride = gridDim.x * blockDim.x
            let mutable i = start
            while i < n do
                // Здесь original = 128, amount = 3
                z.[i]  <- if x.[i] <= 85uy then 0uy elif x.[i] > 170uy then 255uy else  128uy + (x.[i] - 128uy) * 3uy    
                i <- i + stride @>
        |> Compiler.DefineKernel

    let divUp num den = (num + den - 1) / den

// Функция подготовоки и запуска расчета на GPU
    return Entry(fun program ->
        let worker = program.Worker
        let kernel = program.Apply kernel   
    
         // Определим функцию run
        let run (x:byte[])  =      
            let n = x.Length
            use x = worker.Malloc(x)
            use output = worker.Malloc(n)
                     
            // Определим параметры запуска
            let blockSize = 128
            let numSm = worker.Device.Attributes.MULTIPROCESSOR_COUNT
            let gridSize = min (numSm * 16) (divup n blockSize)
           //  let lp = LaunchParam(gridSize, blockSize)
            let lp = LaunchParam(1024, 512)

            // Создаддим два события для фиксации пуска и останова расчета GPU
            use start = worker.CreateEvent()
            use stop = worker.CreateEvent()
            worker.Synchronize()
            start.Record()

            // Запустим GPU
            kernel.Launch lp n x.Ptr output.Ptr

            // Остановим таймер и считаем время работы
            stop.Record()
            stop.Synchronize()
            let msec = Event.ElapsedMilliseconds(start, stop)
            // Считаем выходные данные
            let output = output.Gather()
            output, msec
        run) }

// Функция чтения из файла в байтовый массив
let readImageBytes() =
    // Создадим Bitmap
    let image = new Bitmap(fileNameImage)
    // Создадим прямоугольник с размерами изображения
    let rect  =   new Rectangle(0,0, image.Width,image.Height)
    // Заблокируем прямоугольник
    let bitmapData  =   image.LockBits(rect, ImageLockMode.ReadWrite, PixelFormat.Format32bppArgb)
    // Вычисленим размер image в байтах
    let all_Bytes = Math.Abs(bitmapData.Stride) * bitmapData.Height
    // Созданим массив байтов imageBytes
    let imgBytes:byte[] = Array.zeroCreate<byte> all_Bytes
    // Скопируем данные в массив imageBytes
    let copy() = Marshal.Copy(bitmapData.Scan0, imgBytes, 0, all_Bytes)
    copy()
    imgBytes, image, bitmapData, all_Bytes

// Обновление jpg файла на GPU
let updateImageGPU() =
    use program = pfunct |> Compiler.load Worker.Default
    let calc =
        // Запустим подсчет полного времени расчета на GPU
        let watchGPU = System.Diagnostics.Stopwatch.StartNew()
        // Считаем jpg файл с диска в байтовый массив
        let imageBytes, image, bitmapData, all_Bytes  = readImageBytes()
        // Запустим обновление
        let gpuResults, gpuTime = program.Run imageBytes
        //Выведем время преобразования
        printfn "GPU %10.6f ms"  gpuTime
        // Скопируем данные обратно, из gpuResults в bitmapData
        Marshal.Copy(gpuResults, 0, bitmapData.Scan0, all_Bytes)
        // Разблокируем изображение
        image.UnlockBits(bitmapData)
        // Запишем данные в выходной файл
        image.Save(outFileNameImage, Imaging.ImageFormat.Jpeg)
        // Выведем полное время работы
        printfn "GPU %10.6f ms"   watchGPU.Elapsed.TotalMilliseconds
    calc

// Обновление jpg файла на CPU
let updateImageCPU() =
    // Запусти подсчет полного времени работы на CPU
    let watchCPU = System.Diagnostics.Stopwatch.StartNew()
    // Считаем jpg файл с диска в байтовый массив
    let imageBytes, image, bitmapData, all_Bytes  = readImageBytes()
    
    // Функция увеличения контрастности
    let setupBrightness() =
        // Запомним время старта преобразования массива на CPU
        let watchImg = System.Diagnostics.Stopwatch.StartNew()
        for y = 0 to image.Height - 1 do
            for x = 0 to image.Width - 1 do
                // Ширина шага
                let i =    y * bitmapData.Stride + x * 4
                // blue
                let tempColor   =  original + (int imageBytes.[i] - original) * amount
                let blue        = if tempColor < 0 then  0  else  if tempColor > 255 then 255 else tempColor
                imageBytes.[i]  <-  byte blue
                //green
                let tempColor   =  original + (int imageBytes.[i + 1] - original) * amount
                let green       = if tempColor < 0 then  0  else  if tempColor > 255 then 255 else tempColor
                imageBytes.[i + 1]  <-  byte green   
                //red
                let tempColor   =  original + (int imageBytes.[i + 2] - original) * amount
                let red         = if tempColor < 0 then  0  else  if tempColor > 255 then 255 else tempColor
                imageBytes.[i + 2]  <-  byte red  
                // alpha
                imageBytes.[i + 3] <- imageBytes.[i + 3]
        // Остановим время
        watchImg.Stop()
        // Выведем время преобразования на CPU
        printfn "(CPU %10.6f ms)" watchImg.Elapsed.TotalMilliseconds   

    // Запустим функция увеличения контрастности      
    setupBrightness()
    // Скопируем данные обратно, из ImageBytes в bitmapData
    Marshal.Copy(imageBytes, 0, bitmapData.Scan0, all_Bytes)
    // Разблокируем изображение
    image.UnlockBits(bitmapData)
    // Запишем данные в выходной файл
    image.Save(outFileNameImage, Imaging.ImageFormat.Jpeg)
    // Остановим подсчет времени
    watchCPU.Stop()
    // Выведем полное время работы на CPU на консоль
    printfn "(CPU %10.6f ms)" watchCPU.Elapsed.TotalMilliseconds   

[<EntryPoint>]
let main argv =
    updateImageCPU()
    updateImageGPU()
    System.Console.ReadKey() |> ignore
    0
3. Обновление изображения на GPU.
В состав программы обновления Image, работающей на GPU входят:

1. Шаблон pfunct.
2. Функция updateImageGPU().
3. Функция readImageBytes(), которая работает, как уже указывалось, на CPU в контексте задачи GPU.

pfunct - это один из возможных и рекомендуемых компанией Quantalea шаблонов для запуска вычислений на GPU. Здесь определяется рабочая функция для вычислений на массиве процессоров. Определяются параметры запуска GPU, в том числе размер и число блоков, число потоков, входные и выходные массивы данных, их типы и размеры, а так-же ряд других параметров.
Результатом работы этого шаблона являются выходной байтовый массив, содержащий модифицированное изображение и время работы GPU в формате float.

Функция updateImageGPU() - это основная функция вычислений на GPU. Она содержит все необходимые вызовы других функций, записывает выходные данные в файл, подсчитывает полное время работы своей работы.
4. Экспресс анализ результатов работы программы.
В конфигурации аппаратных средств, установленных на машине автора, выигрыш в производительности GPU по отношению к CPU составляет около двух порядков, но только для функций обновления байтовых массивов. Здесь выигрыш бесспорен.
Однако полные времена работы программ на CPU и GPU окозались уже вполне сопоставимы. При этом надо помнить, что чтение и запись на диск выполнялись совершенно одинаково на CPU для обеих версий программы, значит, длительности работы на этих этапах были одинаковы. Причина увеличения общего времени расчета на GPU по мнению автора кроется в необходимости загрузки (10000 x 5000 x 4) байтового массива в память GPU, а после модификации выгрузки его из памяти GPU в память CPU.
Вероятно для расчета на GPU подходят задачи с большим числом вычислений, хорошо поддающиеся распараллеливанию и, по возможности, не имеющие большого трафика обмена между памятью устройства и памятью центрального процессора.

Проект приложения можно свободно загрузить.
Евгений Вересов.
25.05.2014 года.