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 TearDown()
{
program.Dispose();
context.Dispose();
}
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());
}
