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