Работает с OpenCL

Среда PascalABC.NET
#1

Недавно особо приспичило, по сразу нескольким причинам понадобилось выполнять код на GPU на 2 компах. CUDA не подошла, ибо на 1 из тех компов стоит видеокарта AMD. И вот мой первый опыт:

  1. Драйверы:

Я опросил всех знакомых (давая им запустить простейшую программку, которая проверяет кол-во девайсов поддерживающих OpenCL), таким образом проверив компов 30 - у всех OpenCL стоял и всё было ок. У некоторых было даже не 1 а 2 поддерживающих девайса. Но мне при этом - так уж повезло, на обоих моих компах OpenCL не работал.

С NVidia-видеокартой всё было предельно просто:
1) Зашёл на сайт чтоб скачать установщик - он сразу на первой странице подобрал мне драйвер по характеристикам видеокарты, так что ничего искать и вводить не пришлось.
2) Сохранил раздел реестра HKEY_CURRENT_USER\SOFTWARE\Microsoft\Windows\Shell\Bags\1\Desktop (положение значков на рабочем столе)
3) Снёс драйверы с помощью DDU и запустил тот установщик.

С первого раз всё получилось, без каких либо осложнений.
Решив что переустановка драйверов это предельно легко - я сразу стал переустанавливать драйверы и на второму компу, ноутбуку с видеокартой AMD.

Но не тут то было, в итоге, переустановил я всё раз 15, и никак оно прямо не встанет. При чём, каждый раз я скачивал установщик драйвера из другого места, всё ещё оставаясь на офф сайте AMD. У них всё ужасно запутано.

Потом ещё прочитал - оказывается у меня снова звёзды совпали: у ноутбуков и у AMD могут быть серьёзные осложнения при установке драйверов, а у меня тут комбо…

В общем если что - лучше сразу мастера звать.

В деле помогали следующие утилиты:

  • DDU - чистелка старых драйверов.
  • GPU Caps и GPU-Z - показывают много подробной инфы о железе, и умеют запускать примеры, позволяющие быстро протестировать встал ли OpenCL.
  1. Использование функционала:

Тыря куски кода и другую инфу из следующих источников:

  • Кишки NOpenCL, C# врапера OpenCL. Он и в целом виде работает, но для моих целей не подходил. Из него я тырил большинство кода.
  • Разные C++ примеры OpenCL. В них перепроверял типы параметров и ещё некоторых моментов.
  • Кривая справка OpenCL. В целом она была бесполезна, но мало ли, может ещё проявит себя…
  • Разные статьи на хабре. Вот та что помогла больше всего. А меньше всего помогли те что говорят использовать всякие враперы вроде cloo, ибо ни 1 из этих враперов стабильно не работает.
  • Ну и конечно, родимые гугл, стак оверфлоу и msdn, без них никак.

Я смог собрать этого монстра на 300+ строк:

uses System;
uses System.Runtime.InteropServices;

{$region ErrorCode}

type
  ErrorCode = (
    Success = 0,
    
    DeviceNotFound = -1,
    DeviceNotAvailable = -2,
    CompilerNotAvailable = -3,
    MemObjectAllocationFailure = -4,
    OutOfResources = -5,
    OutOfHostMemory = -6,
    ProfilingInfoNotAvailable = -7,
    MemCopyOverlap = -8,
    ImageFormatMismatch = -9,
    ImageFormatNotSupported = -10,
    BuildProgramFailure = -11,
    MapFailure = -12,
    MisalignedSubBufferOffset = -13,
    ExecStatusErrorForEventsInWaitList = -14,
    CompileProgramFailure = -15,
    LinkerNotAvailable = -16,
    LinkProgramFailure = -17,
    DevicePartitionFailed = -18,
    KernelArgInfoNotAvailable = -19,
    
    InvalidValue = -30,
    InvalidDeviceType = -31,
    InvalidPlatform = -32,
    InvalidDevice = -33,
    InvalidContext = -34,
    InvalidQueueProperties = -35,
    InvalidCommandQueue = -36,
    InvalidHostPtr = -37,
    InvalidMemObject = -38,
    InvalidImageFormatDescriptor = -39,
    InvalidImageSize = -40,
    InvalidSampler = -41,
    InvalidBinary = -42,
    InvalidBuildOptions = -43,
    InvalidProgram = -44,
    InvalidProgramExecutable = -45,
    InvalidKernelName = -46,
    InvalidKernelDefinition = -47,
    InvalidKernel = -48,
    InvalidArgIndex = -49,
    InvalidArgValue = -50,
    InvalidArgSize = -51,
    InvalidKernelArgs = -52,
    InvalidWorkDimension = -53,
    InvalidWorkGroupSize = -54,
    InvalidWorkItemSize = -55,
    InvalidGlobalOffset = -56,
    InvalidEventWaitList = -57,
    InvalidEvent = -58,
    InvalidOperation = -59,
    InvalidGlObject = -60,
    InvalidBufferSize = -61,
    InvalidMipLevel = -62,
    InvalidGlobalWorkSize = -63,
    InvalidProperty = -64,
    InvalidImageDescriptor = -65,
    InvalidCompilerOptions = -66,
    InvalidLinkerOptions = -67,
    InvalidDevicePartitionCount = -68,
    
    CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR = -1000,
    CL_PLATFORM_NOT_FOUND_KHR = -1001,
    CL_INVALID_D3D10_DEVICE_KHR = -1002,
    CL_INVALID_D3D10_RESOURCE_KHR = -1003,
    CL_D3D10_RESOURCE_ALREADY_ACQUIRED_KHR = -1004,
    CL_D3D10_RESOURCE_NOT_ACQUIRED_KHR = -1005);
  
{$endregion ErrorCode}
  
{$region MemoryFlags}
  
  [Flags]
  MemoryFlags = (
    None = 0,
    ReadWrite = 1 shl 0,
    WriteOnly = 1 shl 1,
    ReadOnly = 1 shl 2,
    UseHostPointer = 1 shl 3,
    AllocateHostPointer = 1 shl 4,
    CopyHostPointer = 1 shl 5,
    //Reserved = 1 shl 6,
    HostWriteOnly = 1 shl 7,
    HostReadOnly = 1 shl 8,
    HostNoAccess = 1 shl 9
  );

{$endregion MemoryFlags}

{$region CommandQueueProperties}

  [Flags]
  CommandQueueProperties = (
    cqNone = 0,
    OutOfOrderExecutionModeEnable = 1 shl 0,
    ProfilingEnable = 1 shl 1
  );

{$endregion CommandQueueProperties}

type
  [UnmanagedFunctionPointer(CallingConvention.StdCall)]
  CreateContextCallback = procedure(
    [MarshalAs(UnmanagedType.LPStr)] errorInfo: string;
    privateInfo: IntPtr;
    cb: UIntPtr;
    userData: IntPtr
  );
  
  [UnmanagedFunctionPointer(CallingConvention.StdCall)]
  BuildProgramCallback = procedure(prog, userData: IntPtr);

function clGetPlatformIDs(
  numEntries: cardinal;
  [MarshalAs(UnmanagedType.LPArray)] platforms: array of IntPtr;
  var numPlatforms: cardinal
): ErrorCode; external 'opencl.dll';

function clGetDeviceIDs(
  platform: IntPtr;
  deviceType: uint64;
  numEntries: cardinal;
  [MarshalAs(UnmanagedType.LPArray)] devices: array of IntPtr;
  var numDevices: cardinal
): ErrorCode; external 'opencl.dll';

function clCreateContext(
  [MarshalAs(UnmanagedType.LPArray)] properties: array of IntPtr;
  numDevices: cardinal;
  [MarshalAs(UnmanagedType.LPArray)] devices: array of IntPtr;
  pfnNotify: CreateContextCallback;
  userData: IntPtr;
  var ec: ErrorCode
): IntPtr; external 'opencl.dll';

function clCreateCommandQueue(
  context: IntPtr;
  device: IntPtr;
  properties: CommandQueueProperties;
  var ec: ErrorCode
): IntPtr; external 'opencl.dll';

function clCreateProgramWithSource(
  context: IntPtr;
  count: cardinal;
  [MarshalAs(UnmanagedType.LPArray, ArraySubType = UnmanagedType.LPStr)] strings: array of string;
  [MarshalAs(UnmanagedType.LPArray)] lengths: array of IntPtr;
  var ec: ErrorCode
): IntPtr; external 'opencl.dll';

function clBuildProgram(
  prog: IntPtr;
  numDevices: cardinal;
  [MarshalAs(UnmanagedType.LPArray)] devices: array of IntPtr;
  [MarshalAs(UnmanagedType.LPStr)] options: string;
  pfnNotify: BuildProgramCallback;
  userData: IntPtr
): ErrorCode; external 'opencl.dll';

function clGetProgramBuildInfo(
  prog: IntPtr;
  device: IntPtr;
  paramName: integer;
  paramValueSize: UIntPtr;
  [MarshalAs(UnmanagedType.LPArray)] paramValue: array of char;
  var paramValueSizeRet: UIntPtr
): ErrorCode; external 'opencl.dll';

function clCreateKernel(
  prog: IntPtr;
  [MarshalAs(UnmanagedType.LPStr)] kernelName: string;
  var ec: ErrorCode
): IntPtr; external 'opencl.dll';

function clCreateBuffer(
  context: IntPtr;
  flags: uint64;
  size: IntPtr;
  hostPointer: IntPtr;
  var ec: ErrorCode
): IntPtr; external 'opencl.dll';

function clSetKernelArg(
  kernel: IntPtr;
  argumentIndex: integer;
  argSize: UIntPtr;
  argValue: IntPtr
): ErrorCode; external 'opencl.dll';

function clEnqueueNDRangeKernel(
  commandQueue: IntPtr;
  kernel: IntPtr;
  workDim: cardinal;
  [MarshalAs(UnmanagedType.LPArray)] globalWorkOffset: array of IntPtr;
  [MarshalAs(UnmanagedType.LPArray)] globalWorkSize: array of IntPtr;
  [MarshalAs(UnmanagedType.LPArray)] localWorkSize: array of IntPtr;
  numEventsInWaitList: cardinal;
  eventWaitList: IntPtr; // [MarshalAs(UnmanagedType.CustomMarshaler, MarshalTypeRef = typeof(SafeHandleArrayMarshaler))] array of IntPtr;
  var ev: IntPtr
): ErrorCode; external 'opencl.dll';

function clEnqueueReadBuffer(
  commandQueue: IntPtr;
  buffer: IntPtr;
  [MarshalAs(UnmanagedType.Bool)] blockingRead: boolean;
  offset: IntPtr;
  size: IntPtr;
  destination: IntPtr;
  numEventsInWaitList: cardinal;
  eventWaitList: IntPtr; // [MarshalAs(UnmanagedType.CustomMarshaler, MarshalTypeRef = typeof(SafeHandleArrayMarshaler))] array of IntPtr;
  var ev: IntPtr
): ErrorCode; external 'opencl.dll';

begin
  
  try
    writeln('-' * 30);
    var c: cardinal;
    var pfs: array of IntPtr;
    
    clGetPlatformIDs(0, pfs, c).ToString.Println;
    c.ToString.Println;
    
    pfs := new IntPtr[c];
    clGetPlatformIDs(c, pfs, c).ToString.Println;
    pfs.Println;
    
    
    
    writeln('-' * 30);
    var dvs: array of IntPtr;
    c := 0;
    
    clGetDeviceIDs(pfs[0], 1, 0, dvs, c).ToString.Println;
    c.ToString.Println;
    
    dvs := new IntPtr[c];
    clGetDeviceIDs(pfs[0], 1, c, dvs, c).ToString.Println;
    dvs.PrintLines;
    
    
    
    writeln('-' * 30);
    var ec: ErrorCode;
    var context := clCreateContext(nil, dvs.Length, dvs, nil, IntPtr.Zero, ec);
    ec.ToString.Println;
    context.ToString.Println;
    
    
    
    writeln('-' * 30);
    var command_queue := clCreateCommandQueue(context, dvs[0], cqNone, ec);
    ec.ToString.Println;
    command_queue.ToString.Println;
    
    
    
    writeln('-' * 30);
    {$resource test.cl}
    var prog_str := System.IO.StreamReader.Create(GetResourceStream('test.cl')).ReadToEnd;
    writeln(prog_str);
    var prog := clCreateProgramWithSource(
      context,
      1,
      new string[](prog_str),
      new IntPtr[](new IntPtr(prog_str.Length)),
      ec);
    ec.ToString.Println;
    prog.ToString.Println;
    
    
    
    writeln('-' * 30);
    ec := clBuildProgram(prog, dvs.Length, dvs, nil, nil, IntPtr.Zero);
    ec.ToString.Println;
    if ec <> ErrorCode.Success then
    begin
      writeln('Getting error log:');
      //CL_PROGRAM_BUILD_LOG = 4483
      var sz: UIntPtr;
      clGetProgramBuildInfo(prog, dvs[0], 4483, UIntPtr.Zero, nil, sz).ToString.Println;
      sz.ToString.Println;
      var chs := new char[sz.ToUInt64];
      clGetProgramBuildInfo(prog, dvs[0], 4483, sz, chs, sz).ToString.Println;
      sz.ToString.Println;
      chs.Select(ch -> word(ch)).Println; // почему то всё нули, в инете не нашёл про это инфу
      //writeln(string.Create(chs));
    end;
    
    
    
    writeln('-' * 30);
    var kernel := clCreateKernel(prog, 'TEST', ec);
    ec.ToString.Println;
    kernel.ToString.Println;
    
    
    
    writeln('-' * 30);
    var mem := Marshal.AllocHGlobal(40);
    Marshal.Copy(ArrFill(10,1),0,mem,10);
    var memobj := clCreateBuffer(context,uint64(MemoryFlags.ReadWrite or MemoryFlags.UseHostPointer),IntPtr(40),mem,ec); // UseHostPointer работает вместо clEnqueueWriteBuffer
    ec.ToString.Println;
    memobj.ToString.Println;
    
    
    
    writeln('-' * 30);
    clSetKernelArg(kernel,0,new UIntPtr(sizeof(pointer)),new IntPtr(@memobj)).ToString.Println;
    
    
    
    writeln('-' * 30);
    var ev: IntPtr;
    clEnqueueNDRangeKernel(command_queue, kernel, 1, nil, new IntPtr[](IntPtr(10)),nil, 0,IntPtr.Zero,ev).ToString.Println;
    
    
    
    writeln('-' * 30);
    clEnqueueReadBuffer(command_queue,memobj,true,IntPtr(0),IntPtr(40),mem,0,IntPtr.Zero,ev).ToString.Println;
    
    
    writeln('-' * 30);
    var res := new integer[10];
    Marshal.Copy(mem,res,0,10);
    res.Println;
    
  except
    on e: Exception do writeln(e);
  end;
  
  if not System.Console.IsInputRedirected then readln; // срабатывает только в Shift+F9 режиме (ну и запуске самого .exe)
end.

И ядро:

__kernel void TEST(__global int* message)
{
	int gid = get_global_id(0);

	message[gid] += gid;
}

Архив с исходниками и .exe (18,5 КБ)

Предполагается - что эту тему будут использовать для обсуждения прочих подводных камней связанных с OpenCL.