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