public CommandQueue(ComputeProvider provider, Cl.Device device, bool outOfOrderExecution = false)
{
Cl.ErrorCode error;
_queue = Cl.CreateCommandQueue(provider.Context, device, outOfOrderExecution ? Cl.CommandQueueProperties.OutOfOrderExecModeEnable : Cl.CommandQueueProperties.None, out error);
if (error != Cl.ErrorCode.Success)
throw new CLException(error);
}
public override Brahma.CommandQueue Finish()
{
ClNet.ErrorCode error = Cl.Finish(_queue);
if (error != ClNet.ErrorCode.Success)
{
throw new Cl.Exception(error);
}
return(this);
}
public CommandQueue(ComputeProvider provider, ClNet.Device device, bool outOfOrderExecution = false)
{
ClNet.ErrorCode error;
_queue = Cl.CreateCommandQueue
(provider.Context
, device
, outOfOrderExecution
? ClNet.CommandQueueProperties.OutOfOrderExecModeEnable
: ClNet.CommandQueueProperties.None, out error);
if (error != ClNet.ErrorCode.Success)
{
throw new Cl.Exception(error);
}
}
public Buffer(ComputeProvider provider, Operations operations, Memory memory, IntPtr data, int length) // Create and copy/use data from host
{
ClNet.ErrorCode error;
_length = length;
_mem = Cl.CreateBuffer(provider.Context, (ClNet.MemFlags)operations | (memory == Memory.Host ? ClNet.MemFlags.UseHostPtr : (ClNet.MemFlags)memory | (data != IntPtr.Zero ? ClNet.MemFlags.CopyHostPtr : 0)),
(IntPtr)(_elementSize * _length), data, out error);
if (error != ClNet.ErrorCode.Success)
{
throw new Cl.Exception(error);
}
Operations = operations;
Memory = memory;
}
public Buffer(ComputeProvider provider, Operations operations, bool hostAccessible, int length) // Create, no data
{
ClNet.ErrorCode error;
_length = length;
var size = (IntPtr)(_length * _elementSize);
_mem = Cl.CreateBuffer(
provider.Context
, (ClNet.MemFlags)operations | (hostAccessible ? ClNet.MemFlags.AllocHostPtr : 0)
, size, null, out error);
if (error != ClNet.ErrorCode.Success)
{
throw new Cl.Exception(error);
}
Operations = operations;
Memory = Memory.Device;
}
public void Retain()
{
Cl.RetainCommandQueue(this);
}
internal static void AddEvent(string name, Cl.Event ev)
{
_namedEvents.Add(name, ev);
}
public void ExternalLoopBody(Cl.Program program)
{
// create kernel
Cl.Kernel kernel = Cl.CreateKernel(program, "ExternalLoopBody", out error);
clSafeCall(error);
// create command queue
Cl.CommandQueue cmdQueue = Cl.CreateCommandQueue(context, device, Cl.CommandQueueProperties.None, out error);
clSafeCall(error);
// allocate host vectors
int[] hres = { 0, 1, 2, 3, 4, 5 };
// allocate device vectors
Cl.Mem dres = Cl.CreateBuffer(context, Cl.MemFlags.ReadWrite | Cl.MemFlags.CopyHostPtr,
(IntPtr)(sizeof(int) * hres.Length), hres, out error);
clSafeCall(error);
// setup kernel arguments
clSafeCall(Cl.SetKernelArg(kernel, 0, dres));
clSafeCall(Cl.SetKernelArg(kernel, 1, hres.Length));
// execute kernel
clSafeCall(Cl.EnqueueNDRangeKernel(cmdQueue, kernel, 1, null, new[] { (IntPtr)1 }, null, 0, null, out clevent));
// copy results from device back to host
clSafeCall(Cl.EnqueueReadBuffer(cmdQueue, dres, Cl.Bool.True, IntPtr.Zero,
(IntPtr)(sizeof(int) * hres.Length), hres, 0, null, out clevent));
clSafeCall(Cl.Finish(cmdQueue));
Assert.AreEqual(new[] { 1, 4, 3, 6, 5, 8 }, hres);
}
public void ArrayCompare(Cl.Program program)
{
// create kernel
Cl.Kernel kernel = Cl.CreateKernel(program, "ArrayCompare", out error);
clSafeCall(error);
// create command queue
Cl.CommandQueue cmdQueue = Cl.CreateCommandQueue(context, device, Cl.CommandQueueProperties.None, out error);
clSafeCall(error);
// allocate host vectors
bool[] res = { true, false, true, false };
// allocate device vectors
Cl.Mem dp1 = Cl.CreateBuffer(context, Cl.MemFlags.WriteOnly, (IntPtr)(sizeof(int)), IntPtr.Zero, out error);
clSafeCall(error);
Cl.Mem dp2 = Cl.CreateBuffer(context, Cl.MemFlags.WriteOnly, (IntPtr)(sizeof(int)), IntPtr.Zero, out error);
clSafeCall(error);
Cl.Mem dp3 = Cl.CreateBuffer(context, Cl.MemFlags.WriteOnly, (IntPtr)(sizeof(bool) * res.Length), IntPtr.Zero, out error);
clSafeCall(error);
// setup kernel arguments
clSafeCall(Cl.SetKernelArg(kernel, 0, dp1));
clSafeCall(Cl.SetKernelArg(kernel, 1, dp2));
clSafeCall(Cl.SetKernelArg(kernel, 2, dp3));
// execute kernel
clSafeCall(Cl.EnqueueNDRangeKernel(cmdQueue, kernel, 1, null, new[] { (IntPtr)1 }, null, 0, null, out clevent));
// copy results from device back to host
clSafeCall(Cl.EnqueueReadBuffer(cmdQueue, dp3, Cl.Bool.True, IntPtr.Zero,
(IntPtr)(sizeof(bool) * res.Length), res, 0, null, out clevent));
clSafeCall(Cl.Finish(cmdQueue));
Assert.AreEqual(new[] { false, true, false, true }, res);
// setup kernel arguments
clSafeCall(Cl.SetKernelArg(kernel, 0, dummy));
clSafeCall(Cl.SetKernelArg(kernel, 1, dummy));
// execute kernel
clSafeCall(Cl.EnqueueNDRangeKernel(cmdQueue, kernel, 1, null, new[] { (IntPtr)1 }, null, 0, null, out clevent));
// copy results from device back to host
clSafeCall(Cl.EnqueueReadBuffer(cmdQueue, dp3, Cl.Bool.True, IntPtr.Zero,
(IntPtr)(sizeof(bool) * res.Length), res, 0, null, out clevent));
clSafeCall(Cl.Finish(cmdQueue));
Assert.AreEqual(new[] { true, false, true, false }, res);
}
public void Release()
{
Cl.ReleaseContext(this);
}
public void Release()
{
Cl.ReleaseMemObject(this);
}
public void Release()
{
Cl.ReleaseProgram(this);
}
private static byte[] ToAMDClBinary(this Program irprog, Cl.Device device)
{
// create OpenCL stub before registers naming
string clstub = string.Join("\n", irprog.Kernels.Select(kernel => string.Format(
"__kernel void {0}({1}) {{ {2} }}",
kernel.Name,
string.Join(", ", kernel.FormalParameters.Select(
(fp, idx) => {
switch (fp.StateSpace)
{
case StateSpaces.GLOBAL:
return string.Format("__global {0}* p{1}", fp.UnderlyingType.FormatC(), idx);
case StateSpaces.SHARED:
return string.Format("__local {0}* p{1}", fp.UnderlyingType.FormatC(), idx);
case StateSpaces.CONSTANT:
return string.Format("__constant {0}* p{1}", fp.UnderlyingType.FormatC(), idx);
case StateSpaces.REG:
default:
return string.Format("{0} p{1}", fp.UnderlyingType.FormatC(), idx);
}})),
string.Join(" ", kernel.FormalParameters.Where(
fp => fp.StateSpace != StateSpaces.REG && fp.StateSpace != StateSpaces.CONSTANT).Select(
fp => string.Format("*p{0} = 0;", kernel.FormalParameters.IndexOf(fp)))))));
// create template binary from OpenCL stub
Cl.ErrorCode error;
Cl.Context context = Cl.CreateContext(null, 1, new[] { device }, null, IntPtr.Zero, out error);
clSafeCall(error);
Cl.Program program = Cl.CreateProgramWithSource(context, 1, new[] { clstub }, null, out error);
clSafeCall(error);
clSafeCall(Cl.BuildProgram(program, 1, new[] { device },
"-fno-bin-source -fno-bin-llvmir -fno-bin-exe -fbin-amdil", null, IntPtr.Zero));
Cl.BuildStatus status = Cl.GetProgramBuildInfo(program, device, Cl.ProgramBuildInfo.Status, out error).CastTo<Cl.BuildStatus>();
if (status != Cl.BuildStatus.Success)
throw new Exception(status.ToString());
Cl.InfoBuffer binarySizes = Cl.GetProgramInfo(program, Cl.ProgramInfo.BinarySizes, out error);
clSafeCall(error);
Cl.InfoBufferArray binaries = new Cl.InfoBufferArray(
binarySizes.CastToEnumerable<IntPtr>(Enumerable.Range(0, 1)).Select(sz => new Cl.InfoBuffer(sz)).ToArray());
IntPtr szRet;
clSafeCall(Cl.GetProgramInfo(program, Cl.ProgramInfo.Binaries, binaries.Size, binaries, out szRet));
program.Dispose();
context.Dispose();
// inject generated code into the elf binary
LinkingView elf = new LinkingView(binaries[0].CastToArray<byte>(binarySizes.CastTo<IntPtr>(0).ToInt32()));
SymTabSection symtab = (SymTabSection)elf[".symtab"];
Section amdil = elf[".amdil"];
Section rodata = elf[".rodata"];
MemoryStream amdilcode = new MemoryStream();
foreach (Kernel kernel in irprog.Kernels)
{
SymbolWrapper _metadata = symtab["__OpenCL_" + kernel.Name + "_metadata"];
string[] str_metadata = Marshal.PtrToStringAnsi(Marshal.UnsafeAddrOfPinnedArrayElement(
rodata.Data, (int)_metadata.st_value), (int)_metadata.st_size).Split('\n');
int setup_id = (from line in str_metadata let prms = line.Split(':')
where prms[0] == ";uniqueid" select int.Parse(prms[1])).Single();
int raw_uav_id = (from line in str_metadata let prms = line.Split(':')
where prms[0] == ";uavid" select int.Parse(prms[1])).Single();
SymbolWrapper _fmetadata = symtab["__OpenCL_" + kernel.Name + "_fmetadata"];
string[] str_fmetadata = Marshal.PtrToStringAnsi(Marshal.UnsafeAddrOfPinnedArrayElement(
rodata.Data, (int)_fmetadata.st_value), (int)_fmetadata.st_size).Split('\n');
int func_id = (from line in str_fmetadata let prms = line.Split(':')
where prms[0] == ";uniqueid" select int.Parse(prms[1])).Single();
// ugly, i know!!!
raw_uav_id = Math.Max(raw_uav_id, 11);
int arena_uav_id = raw_uav_id;
byte[] code = Encoding.Convert(Encoding.Unicode, Encoding.ASCII, Encoding.Unicode.GetBytes(
irprog.ToAMDIL(kernel.Name, setup_id, func_id, raw_uav_id, arena_uav_id)));
SymbolWrapper _amdil = symtab["__OpenCL_" + kernel.Name + "_amdil"];
_amdil.st_value = (uint)amdilcode.Position;
_amdil.st_size = (uint)code.Length;
foreach (byte b in code)
amdilcode.WriteByte(b);
}
amdil.Data = amdilcode.ToArray();
return elf.BuildBinary();
}
private static void clSafeCall(Cl.ErrorCode error)
{
if (Cl.ErrorCode.Success != error)
throw new Exception(error.ToString());
}
public static Cl.Program ToGPUClProgram(this Program irprog, Cl.Device device, Cl.Context context)
{
byte[] code = irprog.ToGPUClBinary(device);
Cl.ErrorCode error;
Cl.ErrorCode[] binariesStatus = { Cl.ErrorCode.InvalidBinary };
Cl.Program program = Cl.CreateProgramWithBinary(
context,
1,
new[] { device },
new IntPtr[] { (IntPtr)code.Length },
new Cl.InfoBufferArray(new Cl.InfoBuffer(code)),
binariesStatus,
out error
);
clSafeCall(error);
clSafeCall(binariesStatus[0]);
return program;
}
public static byte[] ToGPUClBinary(this Program irprog, Cl.Device device)
{
Cl.ErrorCode error;
Cl.Platform platform = Cl.GetDeviceInfo(device, Cl.DeviceInfo.Platform, out error).CastTo<Cl.Platform>();
clSafeCall(error);
string platformName = Cl.GetPlatformInfo(platform, Cl.PlatformInfo.Name, out error).ToString();
clSafeCall(error);
switch (platformName)
{
case "NVIDIA CUDA":
return irprog.ToNVIDIAClBinary(device);
case "AMD Accelerated Parallel Processing":
return irprog.ToAMDClBinary(device);
default:
throw new NotSupportedException(platformName);
}
}
private static byte[] ToNVIDIAClBinary(this Program irprog, Cl.Device device)
{
Cl.ErrorCode error;
int cchi = Cl.GetDeviceInfo(device, Cl.DeviceInfo.ComputeCapabilityMajorNV, out error).CastTo<int>();
clSafeCall(error);
int cclo = Cl.GetDeviceInfo(device, Cl.DeviceInfo.ComputeCapabilityMinorNV, out error).CastTo<int>();
clSafeCall(error);
return Encoding.Convert(
Encoding.Unicode, Encoding.ASCII,
Encoding.Unicode.GetBytes(irprog.ToPTX(string.Format("sm_{0}{1}", cchi, cclo))));
}
public CLException(Cl.ErrorCode error, Exception inner) : base(error.ToString(), inner)
{
}
public void Release()
{
Cl.ReleaseCommandQueue(this);
}
public void Retain()
{
Cl.RetainProgram(this);
}
public void Retain()
{
Cl.RetainSampler(this);
}
public void Retain()
{
Cl.RetainMemObject(this);
}
public void Release()
{
Cl.ReleaseSampler(this);
}
public void Retain()
{
Cl.RetainContext(this);
}
public void Retain()
{
Cl.RetainEvent(this);
}
public void ArrayRefOut(Cl.Program program)
{
// create kernel
Cl.Kernel kernel = Cl.CreateKernel(program, "ArrayRefOut", out error);
clSafeCall(error);
// create command queue
Cl.CommandQueue cmdQueue = Cl.CreateCommandQueue(context, device, Cl.CommandQueueProperties.None, out error);
clSafeCall(error);
// allocate host vectors
int[] hp1 = { 1 };
int[] hp2 = { 2 };
// allocate device vectors
Cl.Mem dp1 = Cl.CreateBuffer(context, Cl.MemFlags.CopyHostPtr | Cl.MemFlags.ReadWrite,
(IntPtr)(sizeof(int) * hp1.Length), hp1, out error);
clSafeCall(error);
Cl.Mem dp2 = Cl.CreateBuffer(context, Cl.MemFlags.CopyHostPtr | Cl.MemFlags.ReadWrite,
(IntPtr)(sizeof(int) * hp2.Length), hp2, out error);
clSafeCall(error);
// setup kernel arguments
clSafeCall(Cl.SetKernelArg(kernel, 0, dp1));
clSafeCall(Cl.SetKernelArg(kernel, 1, dp2));
clSafeCall(Cl.SetKernelArg(kernel, 2, dummy));
// execute kernel
clSafeCall(Cl.EnqueueNDRangeKernel(cmdQueue, kernel, 1, null, new[] { (IntPtr)1 }, null, 0, null, out clevent));
// copy results from device back to host
clSafeCall(Cl.EnqueueReadBuffer(cmdQueue, dp1, Cl.Bool.True, IntPtr.Zero,
(IntPtr)(sizeof(int) * hp1.Length), hp1, 0, null, out clevent));
clSafeCall(Cl.EnqueueReadBuffer(cmdQueue, dp2, Cl.Bool.True, IntPtr.Zero,
(IntPtr)(sizeof(int) * hp1.Length), hp2, 0, null, out clevent));
clSafeCall(Cl.Finish(cmdQueue));
Assert.AreEqual(5, hp1[0]);
Assert.AreEqual(4, hp2[0]);
}
public CLException(Cl.ErrorCode error) : base(error.ToString())
{
}
public void SmallTypes(Cl.Program program)
{
// create kernel
Cl.Kernel kernel = Cl.CreateKernel(program, "SmallTypes", out error);
clSafeCall(error);
// create command queue
Cl.CommandQueue cmdQueue = Cl.CreateCommandQueue(context, device, Cl.CommandQueueProperties.None, out error);
clSafeCall(error);
// allocate host vectors
short[] hres1 = { 0 };
short[] hres2 = { 0 };
// allocate device vectors
Cl.Mem dres1 = Cl.CreateBuffer(context, Cl.MemFlags.WriteOnly,
(IntPtr)(sizeof(short) * hres1.Length), IntPtr.Zero, out error);
clSafeCall(error);
Cl.Mem dres2 = Cl.CreateBuffer(context, Cl.MemFlags.WriteOnly,
(IntPtr)(sizeof(short) * hres2.Length), IntPtr.Zero, out error);
clSafeCall(error);
// setup kernel arguments
clSafeCall(Cl.SetKernelArg(kernel, 0, dres1));
clSafeCall(Cl.SetKernelArg(kernel, 1, dres2));
clSafeCall(Cl.SetKernelArg(kernel, 2, (byte)1));
clSafeCall(Cl.SetKernelArg(kernel, 3, (sbyte)-20));
clSafeCall(Cl.SetKernelArg(kernel, 4, (ushort)30));
clSafeCall(Cl.SetKernelArg(kernel, 5, (short)-4));
clSafeCall(Cl.SetKernelArg(kernel, 6, true));
// execute kernel
clSafeCall(Cl.EnqueueNDRangeKernel(cmdQueue, kernel, 1, null, new[] { (IntPtr)1 }, null, 0, null, out clevent));
// copy results from device back to host
clSafeCall(Cl.EnqueueReadBuffer(cmdQueue, dres1, Cl.Bool.True, IntPtr.Zero,
(IntPtr)(sizeof(short) * hres1.Length), hres1, 0, null, out clevent));
clSafeCall(Cl.EnqueueReadBuffer(cmdQueue, dres2, Cl.Bool.True, IntPtr.Zero,
(IntPtr)(sizeof(short) * hres1.Length), hres2, 0, null, out clevent));
clSafeCall(Cl.Finish(cmdQueue));
Assert.AreEqual(7, hres1[0]);
Assert.AreEqual(-7, hres2[0]);
// setup kernel arguments
clSafeCall(Cl.SetKernelArg(kernel, 6, false));
// execute kernel
clSafeCall(Cl.EnqueueNDRangeKernel(cmdQueue, kernel, 1, null, new[] { (IntPtr)1 }, null, 0, null, out clevent));
// copy results from device back to host
clSafeCall(Cl.EnqueueReadBuffer(cmdQueue, dres1, Cl.Bool.True, IntPtr.Zero,
(IntPtr)(sizeof(short) * hres1.Length), hres1, 0, null, out clevent));
clSafeCall(Cl.EnqueueReadBuffer(cmdQueue, dres2, Cl.Bool.True, IntPtr.Zero,
(IntPtr)(sizeof(short) * hres1.Length), hres2, 0, null, out clevent));
clSafeCall(Cl.Finish(cmdQueue));
Assert.AreEqual(-7, hres1[0]);
Assert.AreEqual(7, hres2[0]);
}
public void Release()
{
Cl.ReleaseKernel(this);
}
public void Release()
{
Cl.ReleaseEvent(this);
}
public static ComputeProvider Create(string platformName = "*", Cl.DeviceType deviceType = Cl.DeviceType.Default)
{
var platformNameRegex = new Regex(WildcardToRegex(platformName), RegexOptions.IgnoreCase);
Cl.Platform? currentPlatform = null;
Cl.ErrorCode error;
foreach (Cl.Platform platform in Cl.GetPlatformIDs(out error))
if (platformNameRegex.Match(Cl.GetPlatformInfo(platform, Cl.PlatformInfo.Name, out error).ToString()).Success)
{
currentPlatform = platform;
break;
}
if (currentPlatform == null)
throw new PlatformNotSupportedException(string.Format("Could not find a platform that matches {0}", platformName));
var compatibleDevices = from device in Cl.GetDeviceIDs(currentPlatform.Value, deviceType, out error)
select device;
if (compatibleDevices.Count() == 0)
throw new PlatformNotSupportedException(string.Format("Could not find a device with type {0} on platform {1}",
deviceType, Cl.GetPlatformInfo(currentPlatform.Value, Cl.PlatformInfo.Name, out error)));
return new ComputeProvider(compatibleDevices.ToArray().First());
}
private void clSafeCall(Cl.ErrorCode error)
{
Assert.AreEqual(Cl.ErrorCode.Success, error, error.ToString());
}
public void Retain()
{
Cl.RetainKernel(this);
}
