Интернет магазин китайских планшетных компьютеров



Компьютеры - OpenCL - Особенности языка

14 июня 2011


Оглавление:
1. OpenCL
2. История
3. Особенности языка



Ключевыми отличиями используемого языка от C99 являются:

  • Отсутствие поддержки указателей на функции, рекурсии, битовых полей, массивов переменной длины, стандартных заголовочных файлов
  • Расширения языка для параллелизма: векторные типы, синхронизация, функции для Work-items/Work-Groups
  • Квалификаторы типов памяти: __global, __local, __constant, __private
  • Иной набор встроенных функций


Примеры

Пример вычисления БПФ:

  // создание вычислительного контекста для GPU
  context = clCreateContextFromType;
 
  // создание очереди команд
  queue = clCreateCommandQueue;
 
  // выделение памяти в виде буферов
  memobjs = clCreateBuffer*2*num_entries, srcA, NULL);
  memobjs = clCreateBuffer*2*num_entries, NULL, NULL);
 
  // создание программы из исходных текстов
  program = clCreateProgramWithSource;
 
  // компиляция программы
  clBuildProgram;
 
  // создание объекта kernel из скомпилированной прогаммы
  kernel = clCreateKernel;
 
  // подготовка аргументов
  clSetKernelArg, &memobjs);
  clSetKernelArg, &memobjs);
  clSetKernelArg**16, NULL);
  clSetKernelArg**16, NULL);
 
  // задание N-D диапазона с размерностями work-item и отправка в очередь исполнения
  global_work_size = num_entries;
  local_work_size = 64;
  clEnqueueNDRangeKernel;

Непосредственные вычисления

  // Данный код вычисляет FFT длины 1024, путем разбиения на 16, 16 и 4
 
  __kernel void fft1D_1024 (__global float2 *in, __global float2 *out,
                          __local float *sMemx, __local float *sMemy) {
    int tid = get_local_id;
    int blockIdx = get_group_id * 1024 + tid;
    float2 data;
 
    // адрес начала обрабатываемых данных в глобальной памяти
    in = in + blockIdx;  out = out + blockIdx;
 
    globalLoads; // coalesced global reads
    fftRadix16Pass;      // in-place radix-16 pass
    twiddleFactorMul;
 
    // локальная перестановка с использованием локальной памяти
    localShuffle * 65) + ));
    fftRadix16Pass;               // in-place radix-16 pass
    twiddleFactorMul; // twiddle factor multiplication
 
    localShuffle * 64) + ));
 
    // 4 вызова БПФ порядка 4
    fftRadix4Pass;      // radix-4 function number 1
    fftRadix4Pass;  // radix-4 function number 2
    fftRadix4Pass;  // radix-4 function number 3
    fftRadix4Pass; // radix-4 function number 4
 
    // coalesced global writes
    globalStores;
  }

Полноценная реализация БПФ на OpenCL доступна на сайте Apple



Просмотров: 1860


<<< DirectCompute
Echelon (процессор) >>>