Таким образом, если задача может быть разбита на небольшие блоки, параллельно обрабатывающие небольшой фрагмент блока данных, такая задача может эффективно быть решена на GPU.
Для запуска OpenCL-программы из Python, необходимо поставить библиотеку pyopencl, скачать ее дистрибутив можно со страницы https://wiki.tiker.net/PyOpenCL/Installation/Windows. Для установки библиотеки под Windows необходимо скачать файл и ввести команду pip install pyopencl-2018.1.1+cl12-cp27-cp27m-win_amd64.whl для 64-разрядной версии Python, или pip install pyopencl‑2018.1.1+cl12‑cp27‑cp27m‑win32.whl для 32-разрядной.
Для тестирования pyopencl можно запустить программу, выводящую информацию о системе:
import pyopencl as cl
for plat in cl.get_platforms():
print 'Platform: {}'.format(plat.name)
print 'Version: ' + plat.version
devices = plat.get_devices(cl.device_type.ALL)
print 'Devices:'
for dev in devices:
print('\t{} ({})'.format(dev.name, dev.vendor))
flags = [('Version', dev.version),
('Type', cl.device_type.to_string(dev.type)),
('Memory (global), MB', str(dev.global_mem_size/(1024*1024))),
('Memory (local), KB', str(dev.local_mem_size/1024)),
('Max work item dims', str(dev.max_work_item_dimensions)),
('Max work group size', str(dev.max_work_group_size)),
('Max compute units', str(dev.max_compute_units)),
('Driver version', dev.driver_version),
('Device available', str(bool(dev.available))),
('Compiler available', str(bool(dev.compiler_available)))]
for name, flag in flags:
print '\t\t{0:<25}{1:<10}'.format(name + ':', flag)
Если pyopencl и драйвер видеокарты установлен корректно, мы увидим на экране примерно такой вывод (программа запускалась на Macbook Pro):
Platform: Apple
Version: OpenCL 1.2 (Sep 12 2017 16:28:17)
Devices:
Iris Pro (Intel)
Version: OpenCL 1.2
Type: GPU
Memory (global), MB: 1536
Memory (local), KB: 64
Max work item dims: 3
Max work group size: 512
Max compute units: 40
Driver version: 1.2 (Oct 4 2017 01:28:36)
Device available: True
Compiler available: True
Программа удобна своей кросс-платформенностью, один и тот же код может работать и на OSX и на Windows без каких-либо изменений.
Рассмотрим пример: сформировать массив простых чисел от 1 до N. Для решения такой задачи можно использовать алгоритм “решето Эратосфена”, но мы в учебных целях будем решать задачу напрямую, проверяя каждое число отдельно.
Код программы:
import pyopencl as cl
import pyopencl.array as cl_array
import numpy, time
type = cl.device_type.GPU # cl.device_type.CPU
platform = cl.get_platforms()[0]
devices = platform.get_devices(device_type=type)
ctx = cl.Context(devices)
queue = cl.CommandQueue(ctx)
prg = cl.Program(ctx, """
__kernel void primes(__global int *results)
{
int gid = get_global_id(0);
if (gid < 2) return;
int num = gid + 1;
for(int p=2; p<=num/2 + 1; p++) {
if ((num % p) == 0) {
return;
}
}
int val = atomic_add(&results[0], 1);
results[val+1] = num;
}
""").build()
N = 500000
results = numpy.zeros(N, dtype=numpy.int32)
dest_dev = cl_array.to_device(queue, results)
size = (results.shape[0]-1,)
prg.primes(queue, size, None, dest_dev.data)
res = dest_dev.get()
count = res[0]
primes = numpy.sort(res[1:count+1])
print "Found prime numbers:", count
# for p in primes:
# print p
Разберем текст программы подробнее.
1. Мы создаем контекст устройства , передав ему в качестве параметра тип устройства device_type.GPU. Переменная типа CommandQueue хранит очередь команд, которые будут посланы на устройство.
2. Класс cl.Program получает в качестве параметра программу ядра (kernel) и компилирует ее с помощью вызова функции build. Функция primes написана на языке Си, и будет выполняться параллельно на всех ядрах видеокарты.
3. С помощью функции numpy.zeros мы создаем заполненный нулями массив размера N, затем копируем его на видеокарту с помощью вызова dest_dev = cl_array.to_device. Важно понимать, что существуют две разных копии массива - один в памяти компьютера, второй на видеокарте, с которым и будут выполняться вычисления.
4. Вызовом функции prg.primes выполняется параллельный расчет на видеокарте. При этом в качестве параметра функция получает массив, общий для всех экземпляров функции. Библиотека OpenCL сама распараллеливает вызовы функций, общее количество вызовов равно числу элементов массива, а для получения конкретного идентификатора мы используем функцию get_global_id. В нашем случае идентификатор соответствует числу, которое мы хотим проверить.