private void Prepare()
{
if (prepared)
{
return;
}
Program irprog = this.BuildIR();
Console.WriteLine("Kernels: " +
string.Join("; ", irprog.Kernels.Select(krn => krn.Name)));
Console.WriteLine("Subprograms before inline: " +
string.Join("; ", irprog.GetSubprograms().Except(irprog.Kernels).Select(sp => sp.Name)));
Program irprogInlined = irprog.InlineIR();
Console.WriteLine("Subprograms after inline: " +
string.Join("; ", irprogInlined.GetSubprograms().Except(irprogInlined.Kernels).Select(sp => sp.Name)));
programInlined = irprogInlined.ToGPUClProgram(device, context);
clSafeCall(Cl.BuildProgram(programInlined, 1, new[] { device }, string.Empty, null, IntPtr.Zero));
Assert.AreEqual(Cl.BuildStatus.Success, Cl.GetProgramBuildInfo(programInlined, device, Cl.ProgramBuildInfo.Status, out error).
CastTo <Cl.BuildStatus>());
prepared = true;
}
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 static void Main(string[] args)
{
int dimX = GetParameter("--dimX", 322, args);
int dimY = GetParameter("--dimY", 242, args);
int N = GetParameter("--N", 97760, args);
bool lmem = GetFlag("--lmem", args);
Console.WriteLine("Poisson equation solver: dimX = {0} dimY = {1} N = {2}, LMem: {3}", dimX, dimY, N, lmem);
string options = string.Join(" ", args.Where(arg => arg.IndexOf("--") == -1));
Console.WriteLine("OpenCL program build options: " + options);
Cl.ErrorCode error;
Cl.Device device = (from platformid in Cl.GetPlatformIDs(out error)
from deviceid in Cl.GetDeviceIDs(platformid, Cl.DeviceType.Gpu, out error)
select deviceid).First();
clSafeCall(error);
Cl.Context context = Cl.CreateContext(null, 1, new[] { device }, null, IntPtr.Zero, out error);
clSafeCall(error);
// create program from C# kernel
IRBuildOptions.AutoInline = GetFlag("--inline", args);
IRBuildOptions.WasteRegisters = GetFlag("--rwaste", args);
Console.WriteLine("IR code build options: AutoInline: {0}; WasteRegisters: {1}",
IRBuildOptions.AutoInline, IRBuildOptions.WasteRegisters);
Cl.Program pcsharp = typeof(MainClass).BuildIR().ToGPUClProgram(device, context);
// create program from OpenCL kernel
Cl.Program popencl = Cl.CreateProgramWithSource(context, 1, new[] { PoissonRBSORCl }, null, out error);
clSafeCall(error);
// perform bandwidth comparison
float x0 = (float)(-0.5 * Math.PI);
float y0 = (float)(-0.5 * Math.PI);
float x1 = -x0;
float y1 = -y0;
float omega = 0.8f;
Console.WriteLine("C# benchmark:");
long tcsharp = PoissonRBSOR(device, context, pcsharp, lmem,
x0, y0, x1, y1, dimX, dimY, N, omega, "unigpu.bin", options);
pcsharp.Dispose();
Console.WriteLine("OpenCL benchmark:");
long topencl = PoissonRBSOR(device, context, popencl, lmem,
x0, y0, x1, y1, dimX, dimY, N, omega, "opencl.bin", options);
popencl.Dispose();
Console.WriteLine("OpenCL advantage: {0}", (double)tcsharp / (double)topencl);
context.Dispose();
}
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);
}
private void Prepare(Program irprog)
{
Console.WriteLine("Kernels: " +
string.Join("; ", irprog.Kernels.Select(krn => krn.Name)));
Console.WriteLine("Subprograms: " +
string.Join("; ", irprog.GetSubprograms().Except(irprog.Kernels).Select(sp => sp.Name)));
program = irprog.ToGPUClProgram(device, context);
clSafeCall(Cl.BuildProgram(program, 1, new[] { device }, string.Empty, null, IntPtr.Zero));
Assert.AreEqual(Cl.BuildStatus.Success, Cl.GetProgramBuildInfo(program, device, Cl.ProgramBuildInfo.Status, out error).
CastTo <Cl.BuildStatus>());
}
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 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 void BlockMutation()
{
AssemblyName assemblyName = new AssemblyName("UniGPUTestFixture");
AssemblyBuilder assemblyBuilder = AppDomain.CurrentDomain.DefineDynamicAssembly(assemblyName, AssemblyBuilderAccess.RunAndSave);
ModuleBuilder moduleBuilder = assemblyBuilder.DefineDynamicModule(assemblyName.Name, assemblyName.Name + ".dll");
TypeBuilder typeBuilder = moduleBuilder.DefineType("CILBBTypeMutation", TypeAttributes.Public);
MethodBuilder methodBuilder = typeBuilder.DefineMethod("TestCase", MethodAttributes.Public | MethodAttributes.Static,
typeof(void), new Type[] { typeof(int), typeof(int[]) });
methodBuilder.DefineParameter(1, ParameterAttributes.None, "arg");
methodBuilder.DefineParameter(2, ParameterAttributes.None, "addr");
ILGenerator il = methodBuilder.GetILGenerator();
LocalBuilder lb = il.DeclareLocal(typeof(float));
Label ZERO = il.DefineLabel();
Label LOOP = il.DefineLabel();
Label LOOP_FLT_MUTATOR = il.DefineLabel();
Label LOOP_INT_MUTATOR = il.DefineLabel();
il.Emit(OpCodes.Ldarg_1);
il.Emit(OpCodes.Ldc_I4_0);
il.Emit(OpCodes.Ldarg_0);
il.Emit(OpCodes.Ldarg_0);
il.Emit(OpCodes.Brfalse, ZERO);
il.MarkLabel(LOOP);
il.Emit(OpCodes.Conv_I2);
il.Emit(OpCodes.Starg, 0);
il.Emit(OpCodes.Ldarga, 0);
il.Emit(OpCodes.Dup);
il.Emit(OpCodes.Ldind_I4);
il.Emit(OpCodes.Dup);
il.Emit(OpCodes.Ldc_I4_2);
il.Emit(OpCodes.Rem);
il.Emit(OpCodes.Not);
il.Emit(OpCodes.Ldc_I4_1);
il.Emit(OpCodes.And);
il.Emit(OpCodes.Brtrue, LOOP_FLT_MUTATOR);
il.MarkLabel(LOOP_INT_MUTATOR);
il.Emit(OpCodes.Conv_I4);
il.Emit(OpCodes.Starg, 0);
il.Emit(OpCodes.Ldind_I4);
il.Emit(OpCodes.Ldarg_0);
il.Emit(OpCodes.Add);
il.Emit(OpCodes.Ldc_I4_2);
il.Emit(OpCodes.Div);
il.Emit(OpCodes.Ldc_I4_M1);
il.Emit(OpCodes.Neg);
il.Emit(OpCodes.Sub);
il.Emit(OpCodes.Dup);
il.Emit(OpCodes.Ldc_I4_1);
il.Emit(OpCodes.Bge, LOOP);
il.Emit(OpCodes.Br, ZERO);
il.MarkLabel(LOOP_FLT_MUTATOR);
il.Emit(OpCodes.Conv_R4);
il.Emit(OpCodes.Stloc_0);
il.Emit(OpCodes.Pop);
il.Emit(OpCodes.Ldloc_0);
il.Emit(OpCodes.Ldc_R4, 1.0f);
il.Emit(OpCodes.Sub);
il.Emit(OpCodes.Dup);
il.Emit(OpCodes.Ldc_R4, 1.0f);
il.Emit(OpCodes.Bge, LOOP);
il.Emit(OpCodes.Conv_I4);
il.MarkLabel(ZERO);
il.Emit(OpCodes.Ldc_I4_1);
il.Emit(OpCodes.Add);
il.Emit(OpCodes.Stelem_I4);
il.Emit(OpCodes.Ret);
MethodInfo method = typeBuilder.CreateType().GetMethod(methodBuilder.Name);
int[] res = { 0 };
method.Invoke(null, new object[] { 8, res });
//Assert.AreEqual(1, res[0]);
Cl.Program program = method.BuildIR().ToGPUClProgram(device, context);
clSafeCall(Cl.BuildProgram(program, 1, new[] { device }, string.Empty, null, IntPtr.Zero));
Assert.AreEqual(Cl.BuildStatus.Success, Cl.GetProgramBuildInfo(program, device, Cl.ProgramBuildInfo.Status, out error).
CastTo <Cl.BuildStatus>());
Cl.Kernel kernel = Cl.CreateKernel(program, "TestCase", out error);
clSafeCall(error);
Cl.Mem cl_res = Cl.CreateBuffer(context, Cl.MemFlags.WriteOnly, (IntPtr)sizeof(int), IntPtr.Zero, out error);
clSafeCall(error);
Cl.CommandQueue cmdQueue = Cl.CreateCommandQueue(context, device, (Cl.CommandQueueProperties) 0, out error);
clSafeCall(error);
clSafeCall(Cl.SetKernelArg(kernel, 0, 8));
clSafeCall(Cl.SetKernelArg(kernel, 1, cl_res));
clSafeCall(Cl.EnqueueNDRangeKernel(cmdQueue, kernel, 1, null, new[] { (IntPtr)1 }, null, 0, null, out clevent));
clSafeCall(Cl.EnqueueReadBuffer(cmdQueue, cl_res, Cl.Bool.True, IntPtr.Zero, (IntPtr)sizeof(int), res, 0, null, out clevent));
clSafeCall(Cl.Finish(cmdQueue));
clSafeCall(Cl.ReleaseMemObject(cl_res));
program.Dispose();
Assert.AreEqual(1, res[0]);
}
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]);
}
private static long PoissonRBSOR(Cl.Device device, Cl.Context context, Cl.Program program, bool lmem,
float x0, float y0, float x1, float y1, int dimX, int dimY, int N, float omega,
string fileName = null, string options = "")
{
Cl.ErrorCode error;
Cl.Event clevent;
// build program
clSafeCall(Cl.BuildProgram(program, 1, new[] { device }, options, 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());
}
// save binary
if (fileName != null)
{
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));
byte[] binary = binaries[0].CastToArray <byte>(binarySizes.CastTo <IntPtr>(0).ToInt32());
File.WriteAllBytes(fileName, binary);
}
// create kernel
Cl.Kernel kernel = Cl.CreateKernel(program, "PoissonRBSOR" + (lmem ? "_LMem" : ""), out error);
clSafeCall(error);
// create command queue
Cl.CommandQueue cmdQueue = Cl.CreateCommandQueue(context, device, Cl.CommandQueueProperties.None, out error);
clSafeCall(error);
float hx = (x1 - x0) / dimX;
float hy = (y1 - y0) / dimY;
// boundary values
float[] hgrid = new float[dimX * dimY];
int gstride = dimX;
for (int i = 1; i < dimY - 1; i++)
{
int y_idx = i * gstride;
float y_val = y0 + i * hy;
hgrid[y_idx] = u(x0, y_val);
hgrid[y_idx + dimX - 1] = u(x0 + (dimX - 1) * hx, y_val);
}
for (int j = 1; j < dimX - 1; j++)
{
float x_val = x0 + j * hx;
hgrid[j] = u(x_val, y0);
hgrid[j + (dimY - 1) * gstride] = u(x_val, y0 + (dimY - 1) * hy);
}
// laplacian values
float[] hlaplacian = new float[(dimX - 2) * (dimY - 2)];
int lstride = dimX - 2;
for (int i = 1; i < dimY - 1; i++)
{
for (int j = 1; j < dimX - 1; j++)
{
hlaplacian[j - 1 + (i - 1) * lstride] = J(x0 + j * hx, y0 + i * hy);
}
}
// allocate device vectors
Cl.Mem dgrid = Cl.CreateBuffer(context, Cl.MemFlags.CopyHostPtr | Cl.MemFlags.ReadWrite,
(IntPtr)(sizeof(float) * hgrid.Length), hgrid, out error);
clSafeCall(error);
Cl.Mem dlaplacian = Cl.CreateBuffer(context, Cl.MemFlags.CopyHostPtr | Cl.MemFlags.ReadOnly,
(IntPtr)(sizeof(float) * hlaplacian.Length), hlaplacian, out error);
clSafeCall(error);
// setup kernel arguments
clSafeCall(Cl.SetKernelArg(kernel, 0, dgrid));
clSafeCall(Cl.SetKernelArg(kernel, 1, dlaplacian));
clSafeCall(Cl.SetKernelArg(kernel, 2, dimX));
clSafeCall(Cl.SetKernelArg(kernel, 3, dimY));
clSafeCall(Cl.SetKernelArg(kernel, 4, gstride));
clSafeCall(Cl.SetKernelArg(kernel, 5, lstride));
clSafeCall(Cl.SetKernelArg(kernel, 6, hx));
clSafeCall(Cl.SetKernelArg(kernel, 7, hy));
clSafeCall(Cl.SetKernelArg(kernel, 8, omega));
if (lmem)
{
clSafeCall(Cl.SetKernelArg(kernel, 10, (AREA_SIZE_Y + 2) * (AREA_SIZE_X + 2) * sizeof(float), null));
}
IntPtr[] lo = { (IntPtr)TILE_SIZE_X, (IntPtr)TILE_SIZE_Y };
IntPtr[] gl =
{
(IntPtr)((dimX - 2 + (lmem ? AREA_SIZE_X : TILE_SIZE_X) - 1) /
(lmem ? AREA_SIZE_X : TILE_SIZE_X) * TILE_SIZE_X),
(IntPtr)((dimY - 2 + (lmem ? AREA_SIZE_Y : TILE_SIZE_Y) - 1) /
(lmem ? AREA_SIZE_Y : TILE_SIZE_Y) * TILE_SIZE_Y)
};
// execute RED kernel
clSafeCall(Cl.SetKernelArg(kernel, 9, 1));
clSafeCall(Cl.EnqueueNDRangeKernel(cmdQueue, kernel, 2, null, gl, lo, 0, null, out clevent));
// execute BLACK kernel
clSafeCall(Cl.SetKernelArg(kernel, 9, 0));
clSafeCall(Cl.EnqueueNDRangeKernel(cmdQueue, kernel, 2, null, gl, lo, 0, null, out clevent));
clSafeCall(Cl.Finish(cmdQueue));
Stopwatch stopwatch = new Stopwatch();
stopwatch.Start();
for (int idx = 1; idx < N; idx++)
{
// execute RED kernel
clSafeCall(Cl.SetKernelArg(kernel, 9, 1));
clSafeCall(Cl.EnqueueNDRangeKernel(cmdQueue, kernel, 2, null, gl, lo, 0, null, out clevent));
// execute BLACK kernel
clSafeCall(Cl.SetKernelArg(kernel, 9, 0));
clSafeCall(Cl.EnqueueNDRangeKernel(cmdQueue, kernel, 2, null, gl, lo, 0, null, out clevent));
}
clSafeCall(Cl.Finish(cmdQueue));
stopwatch.Stop();
// copy results from device back to host
clSafeCall(Cl.EnqueueReadBuffer(cmdQueue, dgrid, Cl.Bool.True, IntPtr.Zero,
(IntPtr)(sizeof(float) * hgrid.Length), hgrid, 0, null, out clevent));
clSafeCall(Cl.Finish(cmdQueue));
cmdQueue.Dispose();
kernel.Dispose();
dgrid.Dispose();
float avgerr = 0, maxerr = 0;
for (int i = 1; i < dimY - 1; i++)
{
for (int j = 1; j < dimX - 1; j++)
{
float theory = u(x0 + j * hx, y0 + i * hy);
float err = Math.Abs(theory - hgrid[j + i * gstride]) / Math.Abs(theory);
avgerr += err;
maxerr = Math.Max(maxerr, err);
}
}
avgerr /= dimX * dimY;
long elapsedTime = stopwatch.ElapsedMilliseconds;
Console.WriteLine("average error = {0}%\nmaximal error = {1}%\nelapsed time: {2}ms\niterations per second: {3}",
avgerr * 100, maxerr * 100, elapsedTime, (double)N / (double)elapsedTime * 1000.0d);
return(elapsedTime);
}
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());
}
