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()); }