Программирование на видеокартах GPGPU (1184391), страница 8
Текст из файла (страница 8)
"воображаемые" итераторы (transform_iterator<>,permutation_iterator<>, zip_iterator<> и пр.).Интересно также познакомиться с библиотекой VexCL: (https://github.com/ddemidov/vexcl)VexCL: Vector expression template library for OpenCLhttp://www.codeproject.com/Articles/415058/Это C++-библиотека, генерирующая ядра OpenCL/CUDA на основе векторных выражений.30Стандарт OpenCLOpenCL — это открытый стандарт для параллельного программирования и работы с широкимнабором современных параллельных вычислителей (многоядерных процессоров, GPU, FPGA).Расшифровывается это название как Open Computing Language (что-то вроде: Открытый языкдля вычислений).
Изначально OpenCL был предложен фирмой Apple, но впоследствии получилподдержку многих представителей отрасли, в том числе Intel, AMD, IBM, NVIDIA, ARM, Samsungи др. Ожидается, что каждый изготовитель процессоров реализует доступ к своимвычислительным ресурсам с помощью собственной OpenCL-библиотеки. Есть такая поддержка и врамках CUDA (см. заголовочные файлы в подкаталоге CL/ каталога с заголовочными файлами).Предполагается (в идеале), что предоставляемая библиотека OpenCL может использовать вседоступные ресурсы (GPU, CPU) параллельно, её программная модель основана на C имаксимально абстрагирована от конкретной реализации вычислительных устройств.
Она можетопрашивать и выбирать имеющиеся вычислительные ресурсы, инициализировать их, создаватьтак называемые вычислительные контексты и рабочие очереди. Кроме того, она можеткомпилировать и создавать программы (они тоже называются ядрами), которые затемисполняются на этих ресурсах. Для описания ядер используется C99-подмножество языка C снекоторыми расширениями.Параллельность вычислений обеспечивается независимой работой отдельных вычислителей (илирабочих единиц, в терминологии OpenCL — work-item), которые могут быть объединены врабочие группы (work-group); полное число вычислителей, работающих параллельно, называетсяglobal work size.
Эти вычислители могут взаимодействовать друг с другом, их работа может бытьсинхронизирована в рабочей группе для координации доступа к памяти. Нетрудно заметить, чтопонятие рабочей группы в OpenCL соответствует понятию блока нитей в CUDA, а т.н. глобальныйразмер задаёт аналог сетки блоков из CUDA.Работа программы, использующей OpenCL, протекает примерно так: опрашиваются имеющиесявычислительные ресурсы, выбираются те, что будут далее использоваться, вычислительные ядрасоздаются из исходного кода и распределяются для запуска по вычислительным ресурсам.Таким образом, разработка OpenCL-программы сводится к написанию ядер и host-приложениядля PC, которое распределяет нужные ядра по доступным устройствам.
Такое приложениедолжно использовать пять структур: cl_device_id, cl_kernel, cl_program,cl_command_queue, cl_context. Оно распределяет ядра (cl_kernel), полученные изисходного кода (cl_program) по устройствам (cl_device_id), эти ядра попадают устройствамчерез очередь команд (cl_command_queue); контекст (cl_context) позволяет устройствамполучать ядра и обмениваться данными.Более конкретно, подобное приложение (в качестве примера будем иметь в виду простуюпрограмму вроде hello.c, но не полностью совпадающую с ней) должно получить данные обустройстве, которое будет далее исполнять функцию-ядро, например, как-нибудь так:clGetPlatformIDs(1, &platform, NULL); // первая обнаруженная платформаclGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL); // первое у-во (GPU!)Далее приложение создаёт контекст, например, только с одним обнаруженным устройством:context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);31В последнем вызове и ряде последующих вызовов функциям передаётся адрес переменнойerr, в которую записывается код возможной ошибки при выполнении; реальная программадолжна анализировать этот код, используя его значение для вывода сообщений об ошибках.
Вприводимых здесь строках кода — для простоты — этого не делается.После этого приложение должно получить программу из кода ядра, содержащегося, например, вфайле hello.cl, для чего содержимое этого файла читается в массив, передаваемый функцииclCreateProgramWithSource:program = clCreateProgramWithSource(context, 1,(const char**)&program_buffer, &program_size, &err);clBuildProgram(program, 0, &device, NULL, NULL, NULL);Отсутствующие в последнем вызове параметры могут определять варианты компиляции. Посленеё из заданной функции создаётся ядро (здесь строка "hello" — название функции ядра):kernel = clCreateKernel(program, "hello", &err);Для распределения ядер по устройствам необходимо создавать очереди к устройствам:queue = clCreateCommandQueue(context, device, 0, &err);Поскольку ядру понадобится память для вывода, необходимо также создать буфер памяти:mem = clCreateBuffer(context, CL_MEM_READ_WRITE, MEM_SIZE * sizeof(char), NULL, &ret);Теперь, когда все компоненты окружения (т.е., структуры cl_device_id, cl_kernel,cl_program, cl_command_queue, cl_context) созданы, следует подготовить ядрунеобходимые параметры вызова, например, в случае ядра hello — это один параметр (адресбуфера памяти для вывода):ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&mem);и можно отправлять ядро в очередь на исполнение:global_size = 8;local_size = 4;clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_size,&local_size, 0, NULL, NULL);Последняя функция (как и в случае CUDA) не только обеспечивает запуск ядра на устройстве, но иопределяет, как много рабочих единиц должно быть создано (параметр global_size), а такжесколько рабочих единиц будет в рабочей группе (параметр local_size).Для чтения полученных результатов из буфера памяти mem в массив (в данном случае —символов) string вызывается функция clEnqueueReadBuffer:ret = clEnqueueReadBuffer(queue, mem, CL_TRUE, 0,MEM_SIZE * sizeof(char), string, 0, NULL, NULL);Она возвращает значение кода возможной ошибки (или значение CL_SUCCESS в случаеуспешного завершения).32Синтаксис оформления текста ядер в OpenCL включает в себя спецификаторы: функции ядра(__kernel), памяти (__global, __local) и некоторые другие, встречающиеся более редко.Для того, чтобы каждый исполняемый экземпляр кода ядра мог иметь доступ к «своим» данным ивообще «знать» параметры конфигурации, используются вызовы специальных функций:OpenCLget_local_size(0)// 1, 2get_group_id(0)get_num_groups(0)get_local_id(0)get_global_id(0)Размер рабочей группы(work-group)Номер рабочей группыЧисло рабочих группНомер элемента (work-item)Глобальный номер эл-таget_global_size(0)Глобальный размерCUDAblockDim.x// y, zblockIdx.xgridDim.xthreadIdx.xblockDim.x*blockIdx.x+threadIdx.xblockDim.x*gridDim.xРазмер блока нитей(thread block)Номер блока нитейЧисло блоковНомер нити в блокеГлобальный номер нитиГлобальный размерБолее детальную информацию по спецификации OpenCL можно получить в документе "TheOpenCL Specification" (скажем, файл opencl-1.2.pdf), хотя не факт, что доступная версияOpenCL не окажется более ранней, например, 1.1...
На момент написания пособия уже одобренаверсия 2.0 спецификации.Из приведённого выше краткого описания становится понятно, что самое главное преимуществоOpenCL — переносимость. Ядра OpenCL могут выполняться на различных GPU и CPU самых разныхизготовителей: Intel, AMD, Nvidia, IBM и др., а также на устройствах программируемой логики(FPGA), причём одно и то же приложение может распределять выполнение ядер на многих такихустройствах одновременно. Некоторым недостатком является необходимость освоения новогоподхода со своим набором функций. Избежать этого можно применением какой-нибудь«обёртки» (C++-библиотеки), например, VexCL.В качестве первоначального тестового примера работы с OpenCL предлагаются файлы hello.c,hello.cl.
Первый файл — программа, реализующая почти все описанные выше шаги,необходимые OpenCL, а второй из этих файлов — исходный текст довольно примитивноготестирующего ядра. Их можно поместить в каталог обычного C-проекта (даже не CUDA) и всвойства проекта добавить пути поиска заголовочных файлов $(CUDA_INC_PATH), а также путипоиска библиотек $(CUDA_LIB_PATH) и саму нужную библиотеку OpenCL.lib.Если есть желание попрактиковаться в доведении OpenCL-программ до работоспособногосостояния в системе Windows, то можно взять весьма интересные примеры с сайтаhttp://www.caam.rice.edu/~timwar/HPC12/ (Tim Warburton, профессор факультетаComputational and Applied Mathematics университета Rice).
С одной стороны, эти примерынаписаны на C++, причём довольно изобретательно, с другой стороны, они точноработоспособны, правда, под Windows — далеко не сразу.Тексты с упомянутого сайта написаны, скорее всего, под Mac, поэтому для компиляции их подWindows (конкретно, файла clhelper.hpp с помощью VS2008) придётся кое-чтоподправить. Во-первых, лучше сразу удалить подключение заголовочного файла unistd.h,поскольку под Windows его всё равно нет, а в программе он (возможно, и некоторые другие)не используется.
![](https://s.studizba.com/z.php?f=/templates/si/i/noavatar.png&w=100&h=100&t=1"){kind=avatar}
