public void Run(ComputeContext context, TextWriter log)
{
try
{
ComputeProgram program = new ComputeProgram(context, kernelSources);
program.Build(null, null, null, IntPtr.Zero);
log.WriteLine("Program successfully built.");
ICollection<ComputeKernel> kernels = program.CreateAllKernels();
log.WriteLine("Kernels successfully created.");
// cleanup kernels
foreach (ComputeKernel kernel in kernels)
{
kernel.Dispose();
}
kernels.Clear();
// cleanup program
program.Dispose();
}
catch (Exception e)
{
log.WriteLine(e.ToString());
}
}
public KernelManager(GraphicsInterop interop, InputManager input, string source)
{
_input = input;
var localSizeSingle = (long)Math.Sqrt(interop.Device.MaxWorkGroupSize);
_localSize = new[] { localSizeSingle, localSizeSingle };
//_localSize = new[] { interop.Device.MaxWorkGroupSize, 1 };
_program = new ComputeProgram(interop.Context, source);
try
{
_program.Build(new[] { interop.Device }, "", null, IntPtr.Zero);
}
catch (InvalidBinaryComputeException)
{
Console.WriteLine(_program.GetBuildLog(interop.Device));
return;
}
catch (BuildProgramFailureComputeException)
{
Console.WriteLine(_program.GetBuildLog(interop.Device));
return;
}
Console.WriteLine(_program.GetBuildLog(interop.Device));
_kernels = _program.CreateAllKernels().ToArray();
}
protected override void RunInternal()
{
ComputeProgram program = new ComputeProgram(context, new string[] { kernelSource });
program.Build(null, null, null, IntPtr.Zero);
byte[] bytes = program.Binaries[0];
Console.WriteLine("Compiled program head:");
Console.WriteLine(BitConverter.ToString(bytes, 0, 16) + "...");
}
public static void Run(TextWriter log, ComputeContext context)
{
StartTest(log, "Vector addition test");
try
{
int count = 10;
float[] arrA = new float[count];
float[] arrB = new float[count];
float[] arrC = new float[count];
Random rand = new Random();
for (int i = 0; i < count; i++)
{
arrA[i] = (float)(rand.NextDouble() * 100);
arrB[i] = (float)(rand.NextDouble() * 100);
}
ComputeBuffer<float> a = new ComputeBuffer<float>(context, ComputeMemoryFlags.ReadOnly | ComputeMemoryFlags.CopyHostPointer, arrA);
ComputeBuffer<float> b = new ComputeBuffer<float>(context, ComputeMemoryFlags.ReadOnly | ComputeMemoryFlags.CopyHostPointer, arrB);
ComputeBuffer<float> c = new ComputeBuffer<float>(context, ComputeMemoryFlags.WriteOnly, arrC.Length);
ComputeProgram program = new ComputeProgram(context, kernelSource);
program.Build(null, null, null, IntPtr.Zero);
ComputeKernel kernel = program.CreateKernel("VectorAdd");
kernel.SetMemoryArgument(0, a);
kernel.SetMemoryArgument(1, b);
kernel.SetMemoryArgument(2, c);
ComputeCommandQueue commands = new ComputeCommandQueue(context, context.Devices[0], ComputeCommandQueueFlags.None);
ICollection<ComputeEventBase> events = new Collection<ComputeEventBase>();
// BUG: ATI Stream v2.2 crash if event list not null.
commands.Execute(kernel, null, new long[] { count }, null, events);
//commands.Execute(kernel, null, new long[] { count }, null, null);
arrC = new float[count];
GCHandle arrCHandle = GCHandle.Alloc(arrC, GCHandleType.Pinned);
commands.Read(c, true, 0, count, arrCHandle.AddrOfPinnedObject(), events);
arrCHandle.Free();
for (int i = 0; i < count; i++)
log.WriteLine("{0} + {1} = {2}", arrA[i], arrB[i], arrC[i]);
}
catch (Exception e)
{
log.WriteLine(e.ToString());
}
EndTest(log, "Vector addition test");
}
public void Run(ComputeContext context, TextWriter log)
{
this.log = log;
try
{
program = new ComputeProgram(context, clSource);
program.Build(null, null, notify, IntPtr.Zero);
}
catch (Exception e)
{
log.WriteLine(e.ToString());
}
}
public void Run(ComputeContext context, TextWriter log)
{
try
{
ComputeProgram program = new ComputeProgram(context, kernelSources);
program.Build(null, null, null, IntPtr.Zero);
log.WriteLine("Program successfully built.");
program.CreateAllKernels();
log.WriteLine("Kernels successfully created.");
}
catch (Exception e)
{
log.WriteLine(e.ToString());
}
}
public OpenCLPasswordMatcher ()
{
if (ComputePlatform.Platforms.Count == 0) {
Console.WriteLine ("Cound not find any OpenCL platforms");
Environment.Exit (1);
}
var platform = ComputePlatform.Platforms [0];
logger.Info ("Found {0} computing devices:", platform.Devices.Count);
foreach (var d in platform.Devices) {
logger.Info ("* {0}", d.Name);
}
Context = new ComputeContext (ComputeDeviceTypes.All,
new ComputeContextPropertyList (platform), null, IntPtr.Zero);
Device = Context.Devices [0];
logger.Info ("Using first device.");
// load opencl source
StreamReader streamReader = new StreamReader (MD5_OPENCL_FILE);
string clSource = streamReader.ReadToEnd ();
streamReader.Close ();
// create program with opencl source
ComputeProgram program = new ComputeProgram (Context, clSource);
// compile opencl source
try {
program.Build (null, null, null, IntPtr.Zero);
} catch (Exception e) {
logger.Error ("Build log: " + program.GetBuildLog(Device));
throw e;
}
// load chosen kernel from program
Kernel = program.CreateKernel ("crackMD5");
}
public static void Run(TextWriter log, ComputeContext context)
{
StartTest(log, "Program test");
ProgramTest.log = log;
try
{
ComputeProgram program = new ComputeProgram(context, clSource);
program.Build(null, null, notify, IntPtr.Zero);
byte[] bytes = program.Binaries[0];
log.WriteLine("Compiled program head:");
log.WriteLine(BitConverter.ToString(bytes, 0, 16) + "...");
}
catch (Exception e)
{
log.WriteLine(e.ToString());
}
EndTest(log, "Program test");
}
public void Run(ComputeContext context, TextWriter log)
{
this.log = log;
try
{
DateTime ExecutionStartTime; //Var will hold Execution Starting Time
DateTime ExecutionStopTime;//Var will hold Execution Stopped Time
TimeSpan ExecutionTime;//Var will count Total Execution Time-Our Main Hero
ExecutionStartTime = DateTime.Now; //Gets the system Current date time expressed as local time
program = new ComputeProgram(context, clSource);
program.Build(null, null, notify, IntPtr.Zero);
//program.Build(null, null, null , IntPtr.Zero);
ExecutionStopTime = DateTime.Now;
ExecutionTime = ExecutionStopTime - ExecutionStartTime;
log.WriteLine("Use {0} ms", ExecutionTime.TotalMilliseconds.ToString());
}
catch (Exception e)
{
log.WriteLine(e.ToString());
}
}
/// <summary>
/// OpenCL関係の準備をする
/// </summary>
static void InitializeOpenCL(Real[] result, Real[] left, Real[] right)
{
// プラットフォームを取得
var platform = ComputePlatform.Platforms[0];
Console.WriteLine("プラットフォーム：{0} ({1})", platform.Name, platform.Version);
// コンテキストを作成
var context = new ComputeContext(Cloo.ComputeDeviceTypes.Gpu, new ComputeContextPropertyList(platform), null, IntPtr.Zero);
// 利用可能なデバイス群を取得
var devices = context.Devices;
Console.WriteLine("デバイス数：{0}", devices.Count);
// 1デバイスで使う要素数を計算
countPerDevice = (int)Math.Ceiling((double)COUNT / devices.Count);
// キューの配列を作成
queues = new ComputeCommandQueue[devices.Count];
// 利用可能なデバイスすべてに対して
for(int i = 0; i < devices.Count; i++)
{
var device = devices[i];
// キューを作成
queues[i] = new ComputeCommandQueue(context, device, ComputeCommandQueueFlags.None);
// デバイス情報を表示
Console.WriteLine("* {0} ({1})", device.Name, device.Vendor);
}
// プログラムを作成
var program = new ComputeProgram(context, Properties.Resources.MultiGpu);
// ビルドしてみて
try
{
string realString = ((typeof(Real) == typeof(Double)) ? "double" : "float");
program.Build(devices,
string.Format(" -D REAL={0} -D REALV={0}{1} -D VLOADN=vload{1} -D VSTOREN=vstore{1} -D COUNT_PER_WORKITEM={2} -Werror", realString, VECTOR_COUNT, COUNT_PER_WORKITEM),
null, IntPtr.Zero);
}
// 失敗したら
catch(BuildProgramFailureComputeException ex)
{
// ログを表示して例外を投げる
throw new ApplicationException(string.Format("{0}\n{1}", ex.Message, program.GetBuildLog(devices[0])), ex);
}
// カーネルを作成
addOneElement = new ComputeKernel[devices.Count];
for(int i = 0; i < devices.Count; i++)
{
addOneElement[i] = program.CreateKernel("AddOneElement");
}
// バッファーを作成
bufferLeft = new ComputeBuffer<Real>(context, ComputeMemoryFlags.ReadOnly | ComputeMemoryFlags.UseHostPointer, left);
bufferRight = new ComputeBuffer<Real>(context, ComputeMemoryFlags.ReadOnly | ComputeMemoryFlags.UseHostPointer, right);
bufferResult = new ComputeBuffer<Real>(context, ComputeMemoryFlags.WriteOnly | ComputeMemoryFlags.UseHostPointer, result);
buffersLeft = new ComputeSubBuffer<Real>[devices.Count];
buffersRight = new ComputeSubBuffer<Real>[devices.Count];
buffersResult = new ComputeSubBuffer<Real>[devices.Count];
for(int i = 0; i < devices.Count; i++)
{
buffersLeft[i] = new ComputeSubBuffer<Real>(bufferLeft, ComputeMemoryFlags.ReadOnly, countPerDevice * i, countPerDevice);
buffersRight[i] = new ComputeSubBuffer<Real>(bufferRight, ComputeMemoryFlags.ReadOnly, countPerDevice * i, countPerDevice);
buffersResult[i] = new ComputeSubBuffer<Real>(bufferResult, ComputeMemoryFlags.WriteOnly, countPerDevice * i, countPerDevice);
}
}
public void Run(ComputeContext context, TextWriter log)
{
try
{
// Create the arrays and fill them with random data.
int count = 640*480; //
float[] arrA = new float[count];
float[] arrB = new float[count];
float[] arrC = new float[count];
Random rand = new Random();
for (int i = 0; i < count; i++)
{
arrA[i] = (float)(rand.NextDouble() * 100);
arrB[i] = (float)(rand.NextDouble() * 100);
}
// Create the input buffers and fill them with data from the arrays.
// Access modifiers should match those in a kernel.
// CopyHostPointer means the buffer should be filled with the data provided in the last argument.
program = new ComputeProgram(context, clProgramSource);
program.Build(null, null, null, IntPtr.Zero);
ComputeBuffer<float> a = new ComputeBuffer<float>(context, ComputeMemoryFlags.ReadOnly | ComputeMemoryFlags.CopyHostPointer, arrA);
//ComputeBuffer<float> b = new ComputeBuffer<float>(context, ComputeMemoryFlags.ReadOnly | ComputeMemoryFlags.CopyHostPointer, arrB);
// The output buffer doesn't need any data from the host. Only its size is specified (arrC.Length).
ComputeBuffer<float> c = new ComputeBuffer<float>(context, ComputeMemoryFlags.WriteOnly, arrC.Length);
// Create and build the opencl program.
// Create the kernel function and set its arguments.
ComputeKernel kernel = program.CreateKernel("CompareGPUCPU");
DateTime ExecutionStartTime; //Var will hold Execution Starting Time
DateTime ExecutionStopTime;//Var will hold Execution Stopped Time
TimeSpan ExecutionTime;//Var will count Total Execution Time-Our Main Hero
ComputeCommandQueue commands = new ComputeCommandQueue(context, context.Devices[0], ComputeCommandQueueFlags.None);
ExecutionStartTime = DateTime.Now; //Gets the system Current date time expressed as local time
int repeatTimes = 100;
for (int repeatCounter = 0; repeatCounter < repeatTimes; repeatCounter++)
{
kernel.SetMemoryArgument(0, a);
//kernel.SetMemoryArgument(1, b);
//kernel.SetMemoryArgument(2, c);
kernel.SetMemoryArgument(1, c);
// Create the event wait list. An event list is not really needed for this example but it is important to see how it works.
// Note that events (like everything else) consume OpenCL resources and creating a lot of them may slow down execution.
// For this reason their use should be avoided if possible.
//ComputeEventList eventList = new ComputeEventList();
// Create the command queue. This is used to control kernel execution and manage read/write/copy operations.
// Execute the kernel "count" times. After this call returns, "eventList" will contain an event associated with this command.
// If eventList == null or typeof(eventList) == ReadOnlyCollection<ComputeEventBase>, a new event will not be created.
//commands.Execute(kernel, null, new long[] { count }, null, eventList);
commands.Execute(kernel, null, new long[] { count }, null, null);
// Read back the results. If the command-queue has out-of-order execution enabled (default is off), ReadFromBuffer
// will not execute until any previous events in eventList (in our case only eventList[0]) are marked as complete
// by OpenCL. By default the command-queue will execute the commands in the same order as they are issued from the host.
// eventList will contain two events after this method returns.
//commands.ReadFromBuffer(c, ref arrC, false, eventList);
commands.ReadFromBuffer(c, ref arrC, false, null);
// A blocking "ReadFromBuffer" (if 3rd argument is true) will wait for itself and any previous commands
// in the command queue or eventList to finish execution. Otherwise an explicit wait for all the opencl commands
// to finish has to be issued before "arrC" can be used.
// This explicit synchronization can be achieved in two ways:
// 1) Wait for the events in the list to finish,
//eventList.Wait();
// 2) Or simply use
commands.Finish();
}
ExecutionStopTime = DateTime.Now;
ExecutionTime = ExecutionStopTime - ExecutionStartTime;
double perTaskTime = ExecutionTime.TotalMilliseconds / repeatTimes;
log.WriteLine("Use {0} ms using GPU", perTaskTime);
// Do that using CPU
/*
ExecutionStartTime = DateTime.Now; //Gets the system Current date time expressed as local time
for (int repeatCounter = 0; repeatCounter < repeatTimes; repeatCounter++)
{
for (int i = 0; i < count; i++)
{
//arrC[i] = arrA[i] + arrB[i];
int j;
for (j = 0; j < 330 * 10; j++)
arrC[i] = arrA[i] + j;
}
}
ExecutionStopTime = DateTime.Now;
ExecutionTime = ExecutionStopTime - ExecutionStartTime;
perTaskTime = ExecutionTime.TotalMilliseconds / repeatTimes;
log.WriteLine("Use {0} ms using CPU", ExecutionTime.TotalMilliseconds.ToString());
*/
log.WriteLine("arrA[0]:{0}, arrC[0]:{1}", arrA[0], arrC[0]);
}
catch (Exception e)
{
log.WriteLine(e.ToString());
}
}
static void Main(string[] args)
{
#region
const string programName = "Prime Number";
Stopwatch stopWatch = new Stopwatch();
string clProgramSource = KernelProgram();
Console.WriteLine("Environment OS:");
Console.WriteLine("-----------------------------------------");
Console.WriteLine(Environment.OSVersion);
#endregion
if (ComputePlatform.Platforms.Count == 0)
{
Console.WriteLine("No OpenCL Platforms are availble!");
}
else
{
#region 1
// step 1 choose the first available platform
ComputePlatform platform = ComputePlatform.Platforms[0];
// output the basic info
BasicInfo(platform);
Console.WriteLine("Program: " + programName);
Console.WriteLine("-----------------------------------------");
#endregion
//Cpu 10 seconds Gpu 28 seconds
int count = 64;
int[] output_Z = new int[count * count * count];
int[] input_X = new int[count * count * count];
for (int x = 0; x < count * count * count; x++)
{
input_X[x] = x;
}
#region 2
// step 2 create context for that platform and all devices
ComputeContextPropertyList properties = new ComputeContextPropertyList(platform);
ComputeContext context = new ComputeContext(platform.Devices, properties, null, IntPtr.Zero);
// step 3 create and build program
ComputeProgram program = new ComputeProgram(context, clProgramSource);
program.Build(platform.Devices, null, null, IntPtr.Zero);
#endregion
// step 4 create memory objects
ComputeBuffer<int> a = new ComputeBuffer<int>(context, ComputeMemoryFlags.ReadOnly | ComputeMemoryFlags.CopyHostPointer, input_X);
ComputeBuffer<int> z = new ComputeBuffer<int>(context, ComputeMemoryFlags.WriteOnly, output_Z.Length);
// step 5 create kernel object with same kernel programe name VectorAdd
ComputeKernel kernel = program.CreateKernel("PrimeNumber");
// step 6 set kernel arguments
//kernel.SetMemoryArgument(0, a);
kernel.SetMemoryArgument(0, a);
kernel.SetMemoryArgument(1, z);
ComputeEventList eventList = new ComputeEventList();
//for (int j = 0; j < context.Devices.Count; j++)
// query available devices n,...,1,0. cpu first then gpu
for (int j = context.Devices.Count-1; j > -1; j--)
{
#region 3
stopWatch.Start();
// step 7 create command queue on that context on that device
ComputeCommandQueue commands = new ComputeCommandQueue(context, context.Devices[j], ComputeCommandQueueFlags.None);
// step 8 run the kernel program
commands.Execute(kernel, null, new long[] { count, count, count }, null, eventList);
//Application.DoEvents();
#endregion
// step 9 read results
commands.ReadFromBuffer(z, ref output_Z, false, eventList);
#region 4
commands.Finish();
string fileName = "C:\\primenumber\\PrimeNumberGPU.txt";
StreamWriter file = new StreamWriter(fileName, true);
FileInfo info = new FileInfo(fileName);
long fs = info.Length;
// 1 MegaByte = 1.049e+6 Byte
int index = 1;
if (fs == 1.049e+6)
{
fileName = "C:\\primenumber\\PrimeNumberGPU" + index.ToString() + ".txt";
file = new System.IO.StreamWriter(fileName, true);
index++;
}
#endregion
for (uint xx = 0; xx < count * count * count; xx++)
{
if (output_Z[xx] != 0 && output_Z[xx] != 1)
{
Console.WriteLine(output_Z[xx]);
file.Write(output_Z[xx]);
file.Write("x");
}
}
#region 5
file.Close();
stopWatch.Stop();
ComputeCommandProfilingInfo start = ComputeCommandProfilingInfo.Started;
ComputeCommandProfilingInfo end = ComputeCommandProfilingInfo.Ended;
double time = 10e-9 * (end - start);
//Console.WriteLine("Nanosecond: " + time);
TimeSpan ts = stopWatch.Elapsed;
string elapsedTime = String.Format("{0:00}:{1:00}:{2:00}.{3:00}", ts.Hours, ts.Minutes, ts.Seconds, ts.Milliseconds);
Console.WriteLine(context.Devices[j].Name.Trim() + " Elapsed Time " + elapsedTime);
Console.WriteLine("-----------------------------------------");
#endregion
}
Console.ReadLine();
}
}
bool useGPU = true; // GPU code enabled (from commandline)
#endregion Fields
#region Methods
// initialize renderer: takes in command line parameters passed by template code
public void Init( int rt, bool gpu, int platformIdx )
{
// pass command line parameters
runningTime = rt;
useGPU = gpu;
gpuPlatform = platformIdx;
// initialize accumulator
accumulator = new Vector3[screen.width * screen.height];
ClearAccumulator();
// setup scene
scene = new Scene();
// setup camera
camera = new Camera( screen.width, screen.height );
// Generate randoms
Console.Write("Generating randoms....\t");
randoms = new float[1000];
Random r = RTTools.GetRNG();
for (int i = 0; i < 1000; i++)
randoms[i] = (float)r.NextDouble();
int variable = r.Next();
Console.WriteLine("Done!");
// initialize required opencl things if gpu is used
if (useGPU)
{
StreamReader streamReader = new StreamReader("../../kernel.cl");
string clSource = streamReader.ReadToEnd();
streamReader.Close();
platform = ComputePlatform.Platforms[0];
context = new ComputeContext(ComputeDeviceTypes.Gpu, new ComputeContextPropertyList(platform), null, IntPtr.Zero);
queue = new ComputeCommandQueue(context, context.Devices[0], ComputeCommandQueueFlags.None);
program = new ComputeProgram(context, clSource);
try
{
program.Build(null, null, null, IntPtr.Zero);
kernel = program.CreateKernel("Main");
sceneBuffer = new ComputeBuffer<Vector4>(context, ComputeMemoryFlags.ReadOnly | ComputeMemoryFlags.UseHostPointer, scene.toCL());
rndBuffer = new ComputeBuffer<float>(context, ComputeMemoryFlags.ReadOnly | ComputeMemoryFlags.UseHostPointer, randoms);
cameraBuffer = new ComputeBuffer<Vector3>(context, ComputeMemoryFlags.ReadOnly | ComputeMemoryFlags.UseHostPointer, camera.toCL());
outputBuffer = new ComputeBuffer<int>(context, ComputeMemoryFlags.WriteOnly | ComputeMemoryFlags.UseHostPointer, screen.pixels);
skydome = new ComputeBuffer<float>(context, ComputeMemoryFlags.ReadOnly | ComputeMemoryFlags.UseHostPointer, scene.Skydome);
kernel.SetMemoryArgument(0, outputBuffer);
kernel.SetValueArgument(1, screen.width);
kernel.SetValueArgument(2, screen.height);
kernel.SetMemoryArgument(3, sceneBuffer);
kernel.SetValueArgument(4, scene.toCL().Length);
kernel.SetMemoryArgument(5, skydome);
kernel.SetMemoryArgument(6, cameraBuffer);
kernel.SetMemoryArgument(7, rndBuffer);
}
catch (ComputeException e) {
Console.WriteLine("Error in kernel code: {0}", program.GetBuildLog(context.Devices[0]));
Console.ReadLine();
useGPU = false;
}
}
else {
return;
}
}
protected virtual void buildOpenCLProgram()
{
if (CLSourcePaths == null)
{
System.Diagnostics.Trace.Write("No CL source defined.\n");
return;
}
String[] sourceArray = new String[CLSourcePaths.Length];
try
{
for (int i = 0; i < CLSourcePaths.Length; i++)
{
StreamReader sourceReader = new StreamReader(CLSourcePaths[i]);
sourceArray[i] = sourceReader.ReadToEnd();
}
}
catch (FileNotFoundException e)
{
System.Diagnostics.Trace.Write("Can't find: " + e.FileName + "\n");
Environment.Exit(-1);
}
// Build and compile the OpenCL program
_renderKernel = null;
_renderProgram = new ComputeProgram(_commandQueue.Context, sourceArray);
try
{
// build the program
_renderProgram.Build(null, "-cl-nv-verbose", null, IntPtr.Zero);
// create a reference a kernel function
_renderKernel = _renderProgram.CreateKernel("render");
}
catch (BuildProgramFailureComputeException)
{
printBuildLog();
Environment.Exit(-1);
}
catch (InvalidBuildOptionsComputeException)
{
printBuildLog();
Environment.Exit(-1);
}
catch (InvalidBinaryComputeException)
{
printBuildLog();
Environment.Exit(-1);
}
}
private static ComputeKernel Compile(ComputeContext context, string[] sourcecodes, Dictionary<string, string> defines)
{
var program = new ComputeProgram(context, sourcecodes);
var device = context.Devices.Single();
try
{
foreach (var define in defines.Where(define => define.Key.Any(char.IsWhiteSpace) || define.Value.Any(char.IsWhiteSpace)))
{
MessageBox.Show("Invalid define \"" + define.Key + "=" + define.Value + "\": define contained whitespace", "Error");
return null;
}
var options = string.Join(" ", defines.Where(kvp => !string.IsNullOrEmpty(kvp.Value)).Select(kvp => "-D " + kvp.Key + "=" + kvp.Value));
program.Build(new[] { device }, options + " " + StaticSettings.Fetch.OpenClOptions, null, IntPtr.Zero);
var str = program.GetBuildLog(device).Trim();
if (string.IsNullOrEmpty(str) == false)
MessageBox.Show(str, "Build log");
return program.CreateKernel("Main");
}
catch (InvalidBinaryComputeException)
{
MessageBox.Show(program.GetBuildLog(device), "Build error (invalid binary)");
return null;
}
catch (BuildProgramFailureComputeException)
{
MessageBox.Show(program.GetBuildLog(device), "Build error (build program failure)");
return null;
}
}
// Use this for initialization
void Awake()
{
var platform = ComputePlatform.Platforms[0];
_context = new ComputeContext(ComputeDeviceTypes.Cpu,
new ComputeContextPropertyList(platform), null, System.IntPtr.Zero);
_queue = new ComputeCommandQueue(_context, _context.Devices[0], ComputeCommandQueueFlags.None);
string clSource = System.IO.File.ReadAllText(clProgramPath);
_program = new ComputeProgram(_context, clSource);
try {
_program.Build(null, null, null, System.IntPtr.Zero);
} catch(BuildProgramFailureComputeException) {
Debug.Log(_program.GetBuildLog(_context.Devices[0]));
throw;
}
_events = new ComputeEventList();
_updateGridKernel = _program.CreateKernel(clUpdateGridKernelName);
_updateBoidsKernel = _program.CreateKernel(clUpdateBoidsKernelName);
_boundaryKernel = _program.CreateKernel(clBoundaryKernelName);
_pointCounters = new int[nGridPartitions * nGridPartitions * nGridPartitions];
_pointIndices = new int[_pointCounters.Length * maxIndices];
_pointCountersBuffer = new Cloo.ComputeBuffer<int>(
_context, ComputeMemoryFlags.WriteOnly, _pointCounters.Length);
_pointIndicesBuffer = new Cloo.ComputeBuffer<int>(
_context, ComputeMemoryFlags.WriteOnly, _pointIndices.Length);
_gridInfo = new GridInfo() {
worldOrigin = gridbounds.min,
worldSize = gridbounds.size,
cellSize = gridbounds.size * (1f / nGridPartitions),
nGridPartitions = nGridPartitions,
maxIndices = maxIndices
};
_boundaryKernel.SetValueArgument(1, _gridInfo);
_updateGridKernel.SetMemoryArgument(1, _pointCountersBuffer);
_updateGridKernel.SetMemoryArgument(2, _pointIndicesBuffer);
_updateGridKernel.SetValueArgument(3, _gridInfo);
_updateBoidsKernel.SetMemoryArgument(2, _pointCountersBuffer);
_updateBoidsKernel.SetMemoryArgument(3, _pointIndicesBuffer);
_updateBoidsKernel.SetValueArgument(4, _gridInfo);
}
/// <summary>Compiles the program. Returns the build logs for each device.</summary>
/// <param name="SourceCode">Source code array to compile</param>
/// <param name="BuildLogs">Build logs for each device</param>
public static void Compile(string[] SourceCode, out List<string> BuildLogs)
{
//CLProgram Prog = OpenCLDriver.clCreateProgramWithSource(ContextoGPUs, 1, new string[] { sProgramSource }, null, ref Err);
Prog = new ComputeProgram(Context, SourceCode);
//Verifica se compilou em algum device
bool funcionou = false;
for (int i = 0; i < CLCalc.CLDevices.Count; i++)
{
try
{
Prog.Build(new List<ComputeDevice>() { CLCalc.CLDevices[i] }, "", null, IntPtr.Zero);
funcionou = true;
}
catch
{
}
}
//Build Information
BuildLogs = new List<string>();
for (int i = 0; i < CLDevices.Count; i++)
{
string LogInfo = "";
try
{
LogInfo = Prog.GetBuildLog(CLCalc.CLDevices[i]);
}
catch
{
LogInfo = "Error retrieving build info";
}
//if (!CLCalc.CLDevices[i].CLDeviceAvailable) LogInfo = "Possible compilation failure for device " + i.ToString() + "\n" + LogInfo;
BuildLogs.Add(LogInfo);
}
//Nao compilou em nenhum, joga exception
if (!funcionou)
{
throw new Exception("Could not compile program");
}
}
/// <summary>
/// Attempts to initialize OpenCL for the selected GPU.
/// </summary>
internal void InitializeOpenCL()
{
// only initialize once
if (clKernel != null)
return;
// unused memory so Cloo doesn't break with a null ptr
var userDataPtr = Marshal.AllocCoTaskMem(512);
try
{
clDevice = Gpu.CLDevice;
// context we'll be working underneath
clContext = new ComputeContext(
new[] { clDevice },
new ComputeContextPropertyList(clDevice.Platform),
(p1, p2, p3, p4) => { },
userDataPtr);
// queue to control device
clQueue = new ComputeCommandQueue(clContext, clDevice, ComputeCommandQueueFlags.None);
// buffers to store kernel output
clBuffer0 = new ComputeBuffer<uint>(clContext, ComputeMemoryFlags.ReadOnly, 16);
clBuffer1 = new ComputeBuffer<uint>(clContext, ComputeMemoryFlags.ReadOnly, 16);
// obtain the program
clProgram = new ComputeProgram(clContext, Gpu.GetSource());
var b = new StringBuilder();
if (Gpu.WorkSize > 0)
b.Append(" -D WORKSIZE=").Append(Gpu.WorkSize);
if (Gpu.HasBitAlign)
b.Append(" -D BITALIGN");
if (Gpu.HasBfiInt)
b.Append(" -D BFIINT");
try
{
// build kernel for device
clProgram.Build(new[] { clDevice }, b.ToString(), (p1, p2) => { }, userDataPtr);
}
catch (ComputeException)
{
throw new Exception(clProgram.GetBuildLog(clDevice));
}
clKernel = clProgram.CreateKernel("search");
}
finally
{
Marshal.FreeCoTaskMem(userDataPtr);
}
}
public ICalculator GenFractalCalc(List<ProcessLayer> LayerData, FractalType fractaltype, string code, ProcessLayer deflayer)
{
string macros = @"
#pragma OPENCL EXTENSION cl_amd_printf : enable
inline float ABS(float a) {
return a>0?a:-a;
}
inline float ARGC(float2 a) {
return atan2(a.y,a.x);
}
inline float NORM(float2 a) {
return a.x*a.x+a.y*a.y;
}
inline float ABSC(float2 a) {
return sqrt(NORM(a));
}
inline float2 MULC(float2 a, float2 b) {
return (float2)( a.x*b.x-a.y*b.y, a.y*b.x+a.x*b.y );
}
inline float2 DIVC(float2 a, float2 b) {
return (float2)( (a.x*b.x+a.y*b.y)/(b.x*b.x+b.y*b.y), (a.y*b.x-a.x*b.y)/(b.x*b.x+b.y*b.y) );
}
inline float2 lnc(float2 c) {
float r = ABSC(c);
float a = ARGC(c);
return (float2)(log(r),a);
}
inline float2 arctanc(float2 c) {
float2 io = (float2)(0.0f,1.0f);
float2 two = (float2)(2.0f,0.0f);
float2 one = (float2)(1.0f,0.0f);
return (float2)(MULC(DIVC(io,two),lnc(one - MULC(io,c))-lnc(one + MULC(io,c))));
}
inline float2 powc(float2 c, float p) {
if (NORM(c)==0) {
return (float2)(0.0f,0.0f);
} else {
float r = pow(ABSC(c),p);
float a = ARGC(c)*p;
return (float2)(r*cos(a),r*sin(a));
}
}
struct ProcessLayer {
float2 c_old2x;
float2 c_oldx;
float2 c_x;
float2 c_resx;
float c_calc;
float c_cmean;
float c_cvarsx;
float c_cvariance;
int c_active;
int c_isin;
int c_n;
int c_resn;
};
kernel void FractalCalc (
global read_only float2* in_x,
global read_only float2* in_c,
";
StringBuilder kernel = new StringBuilder(macros);
for (int i=0; i< LayerData.Count; i++) {
kernel.Append(" global write_only struct ProcessLayer* out_p" + i);
kernel.Append(i+1==LayerData.Count ? "\n){" : ",\n");
}
bool hastriangle = false;
bool fractdiv = true;
SeqType modesused = 0;
foreach (var it in LayerData)
{
if (it.c_checktype.HasFlag(SeqCheck.MPL_CHECK_TRIANGLE)) hastriangle = true;
if (it.c_checktype.HasFlag(SeqCheck.MPL_CHECK_TRIANGLE_SMOOTH)) hastriangle = true;
modesused |= it.c_seqtype;
}
if (modesused.HasFlag(SeqType.MPL_SEQ_STDDEV)) modesused |= SeqType.MPL_SEQ_VARIANCE;
if (modesused.HasFlag(SeqType.MPL_SEQ_VARIANCE)) modesused |= SeqType.MPL_SEQ_VARSX;
if (modesused.HasFlag(SeqType.MPL_SEQ_VARSX)) modesused |= SeqType.MPL_SEQ_MEAN;
kernel.Append("float2 sumx = (float2)(0.0f,0.0f);");
kernel.Append("float2 meanx = (float2)(0.0f,0.0f);");
kernel.Append("float2 varsx = (float2)(0.0f,0.0f);");
kernel.Append("float2 variacex = (float2)(0.0f,0.0f);");
kernel.Append("float2 sdx = (float2)(0.0f,0.0f);");
kernel.Append("float2 minx = (float2)(0.0f,0.0f);");
kernel.Append("float2 maxx = (float2)(0.0f,0.0f);");
kernel.Append("float2 deltax = (float2)(0.0f,0.0f);");
kernel.Append("float2 deltac = (float2)(0.0f,0.0f);");
kernel.Append("float delta = 0.0f;");
kernel.Append("float newxnorm = 0.0f;");
kernel.Append("float lowbound = 0.0f;");
kernel.Append("float newd = 0.0f;");
kernel.Append("int end = 0;");
kernel.Append("int n = 0;");
kernel.Append("float2 newx = (float2)(0.0f,0.0f);");
kernel.Append("int index = get_global_id(0);");
kernel.Append("float2 x = in_x[index];");
kernel.Append("float2 c = in_c[index];");
for (int i = 0; i < LayerData.Count; i++)
{
kernel.Append("struct ProcessLayer p"+i+";");
kernel.Append("p"+i+".c_active = 1;");
kernel.Append("p"+i+".c_isin = 0;");
kernel.Append("p"+i+".c_x = x;");
kernel.Append("p"+i+".c_oldx = x;");
kernel.Append("p"+i+".c_old2x = x;");
kernel.Append("p"+i+".c_calc = 0;");
kernel.Append("p"+i+".c_cmean = 0;");
kernel.Append("p"+i+".c_cvarsx = 0;");
kernel.Append("p"+i+".c_cvariance = 0;");
}
kernel.Append("struct ProcessLayer* p = 0;");
if (hastriangle)
{
if (fractaltype == FractalType.FRACTAL_TYPE_MANDEL)
{
kernel.Append("float trinorm = ABSC(c);");
// trinorm = c.Magnitude;
}
else
{
kernel.Append("float trinorm = NORM(c);");
// trinorm = c.Norm;
}
}
kernel.Append("while (!end) {");
// while (!end)
kernel.Append("n++;");
// n++;
switch (fractaltype)
{
case FractalType.FRACTAL_TYPE_MANDEL:
kernel.Append("newx = (float2)(x.x*x.x - x.y*x.y,2*x.x*x.y) + c;");
//kernel.Append(@"printf(""%f %f - "",newx.x,newx.y);");
//double sx = x.Real;
//double sy = x.Imaginary;
//return new Complex(sx * sx - sy * sy + c.Real, 2 * sx * sy + c.Imaginary);
break;
case FractalType.FRACTAL_TYPE_MANDEL_N:
kernel.Append("newx = powc(x,pr) + c;");
// return Complex.Pow(x, param) + c;
break;
case FractalType.FRACTAL_TYPE_BURNINGSHIP:
kernel.Append("newx = (float2)(x.x*x.x-x.y*x.y,2*ABS(x.x*x.y)) + c;");
// double sx = x.Real;
// double sy = x.Imaginary;
// return new Complex(sx * sx - sy * sy + c.Real, 2 * absval(sx * sy) + c.Imaginary);
break;
case FractalType.FRACTAL_TYPE_BURNINGSHIP_N:
kernel.Append("newx = powc((ABS(x.x),ABS(x.y)),pr) + c;");
// return Complex.Pow(new Complex(absval(x.Real), absval(x.Imaginary)), n) + c;
break;
case FractalType.FRACTAL_TYPE_DIVERGENT:
kernel.Append("newx = " + code + ";");
// newx = code.eval(x, c, n, param);
break;
case FractalType.FRACTAL_TYPE_CONVERGENT:
kernel.Append("newx = " + code + ";");
fractdiv = false;
break;
default:
throw new NotSupportedException("Unknown FractalType");
}
if (modesused.HasFlag(SeqType.MPL_SEQ_SUM))
{
kernel.Append("sumx += newx;");
//sumx+=newx;
}
if (modesused.HasFlag(SeqType.MPL_SEQ_MEAN))
{
kernel.Append("deltax = newx-meanx;");
kernel.Append("meanx += deltax/(float)n;");
/*Complex delta = newx-meanx;
meanx = meanx+delta/(double)n;*/
if (modesused.HasFlag(SeqType.MPL_SEQ_VARSX))
{
kernel.Append("varsx += MULC(deltax,(newx-meanx));");
//varsx = varsx + delta*(newx-meanx);
if (modesused.HasFlag(SeqType.MPL_SEQ_VARIANCE))
{
kernel.Append("if (n!=1) {");
// if (n!=1) {
kernel.Append("variacex = varsx / (float)((float)n-(float)1.0f);");
//variacex = varsx/((double)n-(double)1);
if (modesused.HasFlag(SeqType.MPL_SEQ_STDDEV))
{
kernel.Append("sdx = powc(variacex,0.5f);");
//sdx = Complex.Sqrt(variacex);
}
kernel.Append("}");
}
}
}
if (modesused.HasFlag(SeqType.MPL_SEQ_MIN))
{
kernel.Append("if (n==1) minx = newx; else {");
kernel.Append("if (NORM(newx)<NORM(minx)) { minx = newx; } }");
//if (n==1) minx=newx; else if (Complex.Abs(newx)<Complex.Abs(minx)) minx=newx;
}
if (modesused.HasFlag(SeqType.MPL_SEQ_MAX))
{
kernel.Append("if (n==1) maxx = newx; else {");
kernel.Append("if (NORM(newx)>NORM(maxx)) { maxx = newx; } }");
//if (n==1) maxx=newx; else if (Complex.Abs(newx)>Complex.Abs(maxx)) maxx=newx;
}
if (modesused.HasFlag(SeqType.MPL_SEQ_DELTA))
{
kernel.Append("deltax = newx - x");
//deltax = newx-x;
}
for (int i=0; i< LayerData.Count; i++)
{
var p = LayerData[i];
kernel.Append("p = &p"+i+";");
kernel.Append("if (p->c_active) {");
//if (p.c_active) {
kernel.Append("p->c_n = n;");
//p.c_n = n;
kernel.Append("p->c_old2x = p->c_oldx;");
kernel.Append("p->c_oldx = p->c_x;");
//p.c_old2x = p.c_oldx;
//p.c_oldx = p.c_x;
switch (p.c_seqtype)
{
case SeqType.MPL_SEQ_NORMAL: kernel.Append("p->c_x = newx;"); break; // p.c_x = newx; break;
case SeqType.MPL_SEQ_SUM: kernel.Append("p->c_x = sumx;"); break; // p.c_x = sumx; break;
case SeqType.MPL_SEQ_MEAN: kernel.Append("p->c_x = meanx;"); break;// p.c_x = meanx; break;
case SeqType.MPL_SEQ_VARSX: kernel.Append("p->c_x = varsx;"); break;
case SeqType.MPL_SEQ_VARIANCE: kernel.Append("p->c_x = variacex;"); break; // p.c_x = variacex; break;
case SeqType.MPL_SEQ_STDDEV: kernel.Append("p->c_x = sdx;"); break; // p.c_x = sdx; break;
case SeqType.MPL_SEQ_MIN: kernel.Append("p->c_x = minx;"); break; // p.c_x = minx; break;
case SeqType.MPL_SEQ_MAX: kernel.Append("p->c_x = maxx;"); break; // p.c_x = maxx; break;
case SeqType.MPL_SEQ_DELTA: kernel.Append("p->c_x = deltax;"); break; // p.c_x = deltax; break;
default: kernel.Append("p->c_x = newx;"); break; // p.c_x = newx; break;
}
kernel.Append("newd = 0;");
//double newd = 0;
switch (p.c_checktype)
{
case SeqCheck.MPL_CHECK_SMOOTH:
if (fractdiv)
{
kernel.Append("newd = exp(-ABSC(p->c_x));");
//newd = Math.Exp(-Complex.Abs(p.c_x));
}
else
{
kernel.Append("newd = exp(-ABSC(p->c_x-p->c_oldx));");
//newd = Math.Exp(-Complex.Abs(p.c_x-p.c_oldx));
}
break;
case SeqCheck.MPL_CHECK_REAL:
kernel.Append("newd = p->c_x.x;");
//newd = p.c_x.Real;
break;
case SeqCheck.MPL_CHECK_IMAG:
kernel.Append("newd = p->c_x.y;");
//newd = p.c_x.Imaginary;
break;
case SeqCheck.MPL_CHECK_ARG:
kernel.Append("newd = atan2(p->c_x.y,p->c_x.x);");
//newd = p.c_x.Phase;
break;
case SeqCheck.MPL_CHECK_ABS:
kernel.Append("newd = ABSC(p->c_x);");
//newd = p.c_x.Magnitude;
break;
case SeqCheck.MPL_CHECK_CURVATURE:
kernel.Append("if (isnotequal(p.c_oldx,p.c_old2x)) { newd = ABSC(atanc(DIVC(p->c_x-p->c_oldx,p->c_oldx-p->c_old2x))); } else newd = 0;");
//if ((p.c_oldx!=p.c_old2x)) newd=Complex.Abs(Complex.Atan((p.c_x-p.c_oldx) / (p.c_oldx-p.c_old2x))); else newd=0; }
break;
case SeqCheck.MPL_CHECK_TRIANGLE:
if (fractaltype == FractalType.FRACTAL_TYPE_MANDEL)
{
kernel.Append("newxnorm = NORM(p->c_oldx);");
//double newxnorm = p.c_oldx.Norm();
kernel.Append("lowbound = ABS(newxnorm-trinorm);");
//double lowbound = absval(newxnorm-trinorm);
kernel.Append("if ((newxnorm+trinorm-lowbound)==0) newd = 0; else newd = (ABSC(p->c_x)-lowbound)/(newxnorm+trinorm-lowbound);");
//if ((newxnorm+trinorm-lowbound)==0) newd=0; else
// newd = (p.c_x.Magnitude-lowbound)/(newxnorm+trinorm-lowbound);
}
else
{
kernel.Append("newxnorm = ABSC(p->c_x);");
//double newxnorm = p.c_x.Magnitude;
kernel.Append("lowbound = ABS(newxnorm-trinorm);");
//double lowbound = absval(newxnorm-trinorm);
kernel.Append("if ((newxnorm+trinorm-lowbound)==0) newd = 0; else newd = (ABSC(p->c_x-c)-lowbound)/(newxnorm+trinorm-lowbound);");
//if ((newxnorm+trinorm-lowbound)==0) newd=0; else
// newd = ((Complex.Abs(p.c_x-c)-lowbound)/(newxnorm+trinorm-lowbound));
}
break;
case SeqCheck.MPL_CHECK_TRIANGLE_SMOOTH:
if (fractaltype == FractalType.FRACTAL_TYPE_MANDEL)
{
kernel.Append("newxnorm = NORM(p->c_oldx);");
//double newxnorm = p.c_oldx.Norm();
kernel.Append("lowbound = ABS(newxnorm-trinorm);");
//double lowbound = absval(newxnorm-trinorm);
kernel.Append("if ((newxnorm+trinorm-lowbound)==0) newd = 0; else newd = (ABSC(p->c_x)-lowbound)/(newxnorm+trinorm-lowbound);");
//if ((newxnorm+trinorm-lowbound)==0) newd=0; else
// newd = (p.c_x.Magnitude-lowbound)/(newxnorm+trinorm-lowbound);
}
else
{
kernel.Append("newxnorm = ABSC(p->c_x);");
//double newxnorm = p.c_x.Magnitude;
kernel.Append("lowbound = ABS(newxnorm-trinorm);");
//double lowbound = absval(newxnorm-trinorm);
kernel.Append("if ((newxnorm+trinorm-lowbound)==0) newd = 0; else newd = (ABSC(p->c_x-c)-lowbound)/(newxnorm+trinorm-lowbound);");
//if ((newxnorm+trinorm-lowbound)==0) newd=0; else
// newd = ((Complex.Abs(p.c_x-c)-lowbound)/(newxnorm+trinorm-lowbound));
}
break;
case SeqCheck.MPL_CHECK_ORBIT_TRAP:
switch (p.c_orbittraptype)
{
case OrbitTrap.MPL_ORBIT_TRAP_POINT:
kernel.Append("newd = ABSC(p->c_x - p->c_pointA);");
//newd = Complex.Abs(p.c_x - p.c_pointA);
break;
case OrbitTrap.MPL_ORBIT_TRAP_LINE:
if ((p.c_pointA.Real) == 1)
{
kernel.Append("newd = ABS(p->c_x.x);");
//newd = Math.Abs(p.c_x.Real);
}
else
{
kernel.Append("newd = ABS(p->c_x.y);");
//newd = Math.Abs(p.c_x.Imaginary);
}
break;
case OrbitTrap.MPL_ORBIT_TRAP_GAUSS:
{
kernel.Append("newd = ABSC((round(p->c_x.x),round(p->c_x.y)) - p->c_x);");
//Complex gauss = new Complex(Math.Round(p.c_x.Real),Math.Round(p.c_x.Imaginary));
//newd = Complex.Abs(gauss - p.c_x);
}
break;
}
break;
}
switch (p.c_checkseqtype)
{
case SeqType.MPL_SEQ_NORMAL: kernel.Append("p->c_calc = newd;"); break;
case SeqType.MPL_SEQ_SUM: kernel.Append("p->c_calc += newd;"); break; // p.c_calc += newd; break;
case SeqType.MPL_SEQ_MEAN: kernel.Append("p->c_calc += newd;"); break; // p.c_calc += newd; break;
case SeqType.MPL_SEQ_VARSX:
{
kernel.Append("delta = newd - p->c_cmean;");
//double delta = newd - p.c_cmean;
kernel.Append("p->c_cmean = p->c_cmean + delta / p->c_n;");
//p.c_cmean = p.c_cmean+delta/p.c_n;
kernel.Append("p->c_calc += delta * (newd - p->c_cmean);");
//p.c_calc += delta*(newd-p.c_cmean);
}
break;
case SeqType.MPL_SEQ_VARIANCE:
{
kernel.Append("delta = newd - p->c_cmean;");
//double delta = newd - p.c_cmean;
kernel.Append("p->c_cmean = p->c_cmean + delta / p->c_n;");
//p.c_cmean = p.c_cmean+delta/p.c_n;
kernel.Append("p->c_cvarsx += delta * (newd - p->c_cmean);");
//p.c_cvarsx = p.c_cvarsx + delta*(newd-p.c_cmean);
kernel.Append("if (p->c_n!=1) { p->c_calc = p->c_cvarsx/(p->c_n-1.0f); }");
/*if (p.c_n!=1) {
p.c_calc = p.c_cvarsx/(p.c_n-1.0);
}*/
}
break;
case SeqType.MPL_SEQ_STDDEV:
{
kernel.Append("delta = newd - p->c_cmean;");
//double delta = newd - p.c_cmean;
kernel.Append("p->c_cmean = p->c_cmean + delta / p->c_n;");
//p.c_cmean = p.c_cmean+delta/p.c_n;
kernel.Append("p->c_cvarsx += delta * (newd - p->c_cmean);");
//p.c_cvarsx = p.c_cvarsx + delta*(newd-p.c_cmean);
kernel.Append("if (p->c_n!=1) { p->c_cvariance = p->c_cvarsx/((float)p->c_n-1.0f);");
/*if (p.c_n!=1) {
p.c_cvariance = p.c_cvarsx/(p.c_n-1.0);
}*/
kernel.Append("p->c_calc = sqrt(p->c_cvariance);");
//p.c_calc = Math.Sqrt(p.c_cvariance);
kernel.Append("}");
}
break;
case SeqType.MPL_SEQ_MIN:
kernel.Append("if (p->c_n==1) p->c_calc = newd; else if (p->c_calc>newd) { p->c_calc = newd; p->c_resx = p->c_x; p->c_resn = p->c_n; };");
//if (p.c_n==1) p.c_calc=newd; else if (p.c_calc>newd) { p.c_calc = newd; p.c_resx = p.c_x; p.c_resn = p.c_n; }
break;
case SeqType.MPL_SEQ_MAX:
kernel.Append("if (p->c_n==1) p->c_calc = newd; else if (p->c_calc<newd) { p->c_calc = newd; p->c_resx = p->c_x; p->c_resn = p->c_n; };");
// if (p.c_n==1) p.c_calc=newd; else if (p.c_calc<newd) { p.c_calc = newd; p.c_resx = p.c_x; p.c_resn = p.c_n; }
break;
case SeqType.MPL_SEQ_DELTA:
kernel.Append("p->c_calc = newd-p->c_calc;");
//p.c_calc = newd-p.c_calc;
break;
default:
kernel.Append("p->c_calc = newd;");
//p.c_calc = newd;
break;
}
if (p.c_convchktype == ConvCheck.MPL_CONVCHK_REAL)
{
kernel.AppendFormat(CultureInfo.InvariantCulture,"if (p->c_x.x*p->c_x.x " + (fractdiv ? ">" : "<") + " {0:E}f) p->c_active = 0;", p.c_bailout);
/*double ddd = p.c_x.Real*p.c_x.Real;
if ((fractdiv) && ( ddd>p.c_bailout)) p.c_active = false;
if (!(fractdiv) && ( ddd<p.c_bailout)) p.c_active = false;*/
}
else if (p.c_convchktype == ConvCheck.MPL_CONVCHK_IMAG)
{
kernel.AppendFormat(CultureInfo.InvariantCulture, "if (p->c_x.y*p->c_x.y " + (fractdiv ? ">" : "<") + " {0:E}f) p->c_active = 0;", p.c_bailout);
/*double ddd = p.c_x.Imaginary*p.c_x.Imaginary;
if ((fractdiv) && ( ddd>p.c_bailout)) p.c_active = false;
if (!(fractdiv) && ( ddd<p.c_bailout)) p.c_active = false;*/
}
else if (p.c_convchktype == ConvCheck.MPL_CONVCHK_OR)
{
kernel.AppendFormat(CultureInfo.InvariantCulture, "if ((p->c_x.y*p->c_x.y " + (fractdiv ? ">" : "<") + " {0:E}f) || (p->c_x.x*p->c_x.x " + (fractdiv ? ">" : "<") + " {0:E}f)) p->c_active = 0;", p.c_bailout);
/*if ((fractdiv) && ((p.c_x.Real*p.c_x.Real>p.c_bailout) || (p.c_x.Imaginary*p.c_x.Imaginary>p.c_bailout))) p.c_active = false;
if (!(fractdiv) && ((p.c_x.Real*p.c_x.Real<p.c_bailout) || (p.c_x.Imaginary*p.c_x.Imaginary<p.c_bailout))) p.c_active = false;*/
}
else if (p.c_convchktype == ConvCheck.MPL_CONVCHK_AND)
{
kernel.AppendFormat(CultureInfo.InvariantCulture, "if ((p->c_x.y*p->c_x.y " + (fractdiv ? ">" : "<") + " {0:E}f) && (p->c_x.x*p->c_x.x " + (fractdiv ? ">" : "<") + " {0:E}f)) p->c_active = 0;", p.c_bailout);
/*if ((fractdiv) && ((p.c_x.Real*p.c_x.Real>p.c_bailout) && (p.c_x.Imaginary*p.c_x.Imaginary>p.c_bailout))) p.c_active = false;
if (!(fractdiv) && ((p.c_x.Real*p.c_x.Real<p.c_bailout) && (p.c_x.Imaginary*p.c_x.Imaginary<p.c_bailout))) p.c_active = false;*/
}
else if (p.c_convchktype == ConvCheck.MPL_CONVCHK_MANH)
{
kernel.AppendFormat(CultureInfo.InvariantCulture, "if ( ((ABS(p->c_x.y)+ABS(p->c_x.x))*((ABS(p->c_x.y)+ABS(p->c_x.x))) " + (fractdiv ? ">" : "<") + " {0:G}f)) p->c_active = 0;", p.c_bailout);
/*double ddd = (absval(p.c_x.Imaginary)+absval(p.c_x.Real))*(absval(p.c_x.Imaginary)+absval(p.c_x.Real));
if ((fractdiv) && ( ddd>p.c_bailout)) p.c_active = false;
if (!(fractdiv) && ( ddd<p.c_bailout)) p.c_active = false;*/
}
else if (p.c_convchktype == ConvCheck.MPL_CONVCHK_MANR)
{
kernel.AppendFormat(CultureInfo.InvariantCulture, "if ( ((p->c_x.y+p->c_x.x)*(p->c_x.y+p->c_x.x)) " + (fractdiv ? ">" : "<") + " {0:E}f)) p->c_active = 0;", p.c_bailout);
/*double ddd = (p.c_x.Real+p.c_x.Imaginary)*(p.c_x.Real+p.c_x.Imaginary);
if ((fractdiv) && ( ddd>p.c_bailout)) p.c_active = false;
if (!(fractdiv) && ( ddd<p.c_bailout)) p.c_active = false; */
}
else
{
kernel.AppendFormat(CultureInfo.InvariantCulture, "if (NORM(p->c_x) " + (fractdiv ? ">" : "<") + " {0:E}f) p->c_active = 0;", p.c_bailout);
/*double ddd = p.c_x.Norm();
if ((fractdiv) && ( ddd>p.c_bailout)) p.c_active = false;
if (!(fractdiv) && ( ddd<p.c_bailout)) p.c_active = false;*/
}
kernel.AppendFormat(CultureInfo.InvariantCulture, "if (p->c_n>{0}) {{ p->c_active = 0; p->c_isin = 1; }}", p.c_nlimit);
//if (p.c_n>p.c_nlimit) { p.c_active = false; p.c_isin = true; }
if (p.c_checktype == SeqCheck.MPL_CHECK_TRIANGLE_SMOOTH)
{
throw new NotImplementedException("Smooth triangle algorithm is unavailable in this CalculatorFactory");
/*if (p.c_active == false)
if (!p.c_isin) {
p.c_oldx = p.c_x;
p.c_x = Fractal_Mandel(p.c_x,c);
p.c_n++;
double newxnorm = p.c_oldx.Norm();
double lowbound = absval(newxnorm-trinorm);
if ((newxnorm+trinorm-lowbound)==0) newd=0; else
newd = (p.c_x.Magnitude-lowbound)/(newxnorm+trinorm-lowbound);
p.c_calc += newd;
double oldsum = p.c_calc/(p.c_n+1);
double il2=1/Math.Log(2);
double lp=Math.Log(Math.Log(p.c_bailout));
double f=il2*lp-il2*Math.Log(Math.Log(Complex.Abs(p.c_x)))+2;
double az2 = p.c_x.Norm();
p.c_oldx = p.c_x;
p.c_x = Fractal_Mandel(p.c_oldx,c);
lowbound = absval(az2-trinorm);
if ((az2+trinorm-lowbound)!=0) p.c_calc+=(Complex.Abs(p.c_x)-lowbound)/(az2+trinorm-lowbound);
p.c_n++;
p.c_calc = p.c_calc/(p.c_n+1);
p.c_calc = oldsum+(p.c_calc-oldsum)*(f-1);
} else {
p.c_calc /= p.c_n+1;
}*/
}
else if (p.c_checkseqtype == SeqType.MPL_SEQ_MEAN)
{
kernel.Append("if (p->c_active == 0) p->c_calc /= (float)p->c_n+1.0f;");
//if (p.c_active == false) p.c_calc /= p.c_n+1;
}
if (p == deflayer)
{
kernel.Append("if (p->c_active == 0) end = 1;");
/*if (!deflayer.c_active) end = true; */
}
kernel.Append("}");
}
kernel.Append("x = newx; }");
for (int i = 0; i < LayerData.Count; i++)
{
kernel.Append("out_p"+i+"[index] = p"+i+";");
//kernel.Append("out_p" + i + "[index].c_calc = 52.0f;");
}
kernel.Append("}");
//System.Console.WriteLine(kernel.Replace(";", ";\n").Replace("}","}\n"));
//kernel.Clear();
//kernel.Append(@"kernel void VectorAdd(global read_only float* a,global read_only float* b,global write_only float* c ){int index = get_global_id(0);c[index] = a[index] + b[index];}");
ComputeProgram prg = new ComputeProgram(_context, kernel.Replace(";", ";\n").Replace("}","}\n").ToString());
try
{
prg.Build(null, null, null, IntPtr.Zero);
}
catch (ComputeException e)
{
throw new Exception("Error while building: " + prg.GetBuildLog(_context.Devices[0]), e);
}
ComputeKernel krnl = prg.CreateKernel("FractalCalc");
return new OpenCLCalculator(_context,prg,krnl);
}
public static void InitializeOpenCL()
{
string source = File.ReadAllText("MonteCarloSimulate.cl");
//Choose Device
ComputePlatform platform = ComputePlatform.Platforms[0];
openCLDevice = platform.QueryDevices()[0];
ComputeContextPropertyList properties =
new ComputeContextPropertyList(platform);
//Setup of stuff on our side
openCLContext = new ComputeContext(ComputeDeviceTypes.All,
properties, null, IntPtr.Zero);
//Build the program, which gets us the kernel
openCLProgram = new ComputeProgram(openCLContext, source);
openCLProgram.Build(null, null, null, IntPtr.Zero);
//can use notify as the 3rd command... if you want this to be non-blocking
openCLKernel = openCLProgram.CreateKernel("MonteCarloSimulate");
}
public void InitGPU(int platformIdx)
{
platform = ComputePlatform.Platforms[platformIdx];
context = new ComputeContext(ComputeDeviceTypes.Gpu, new ComputeContextPropertyList(platform), null, IntPtr.Zero);
StreamReader streamReader = new StreamReader("../../program.cl");
string clSource = streamReader.ReadToEnd();
streamReader.Close();
ComputeProgram program = new ComputeProgram(context, clSource);
program.Build(null, null, null, IntPtr.Zero);
ComputeKernel kernelInit = program.CreateKernel("init");
ComputeKernel kernelUpdate = program.CreateKernel("update");
var flags = ComputeMemoryFlags.ReadWrite | ComputeMemoryFlags.UseHostPointer;
int aantalPixels = screen.height * screen.width;
rays = new GPURay[aantalPixels];
pixels = new Vector3[aantalPixels];
setGPUCameraToCamera();
GPUCamera[] gpuCamArray = { gpuCamera };
ComputeBuffer<GPURay> bufferRays = new ComputeBuffer<GPURay>(context, flags, rays);
ComputeBuffer<Vector3> bufferPixels = new ComputeBuffer<Vector3>(context, flags, pixels);
ComputeBuffer<GPUCamera> bufferCamera = new ComputeBuffer<GPUCamera>(context, flags, gpuCamArray);
kernelUpdate.SetMemoryArgument(0, bufferRays);
kernelUpdate.SetMemoryArgument(1, bufferPixels);
kernelUpdate.SetMemoryArgument(2, bufferCamera);
ComputeCommandQueue queue = new ComputeCommandQueue(context, context.Devices[0], 0);
}
public static void Test()
{
string source = File.ReadAllText("MonteCarloSimulate.cl");
//Choose Device
ComputePlatform platform = ComputePlatform.Platforms[0];
ComputeDevice device = platform.QueryDevices()[0];
ComputeContextPropertyList properties =
new ComputeContextPropertyList(platform);
//Setup of stuff on our side
ComputeContext context = new ComputeContext(ComputeDeviceTypes.All,
properties, null, IntPtr.Zero);
//Build the program, which gets us the kernel
ComputeProgram program = new ComputeProgram(context, source);
program.Build(null, null, null, IntPtr.Zero);
//can use notify as the 3rd command... if you want this to be non-blocking
ComputeKernel kernel = program.CreateKernel("MonteCarloSimulate");
//Create arguments
int sideSize = 4096;
int[] inMatrixA = new int[sideSize * sideSize];
int[] inMatrixB = new int[sideSize * sideSize];
int[] outMatrixC = new int[sideSize * sideSize];
Random random = new Random((int)DateTime.Now.Ticks);
if (sideSize <= 32)
for (int y = 0; y < sideSize; y++)
for (int x = 0; x < sideSize; x++)
{
inMatrixA[y * sideSize + x] = random.Next(3);
inMatrixB[y * sideSize + x] = random.Next(3);
outMatrixC[y * sideSize + x] = 0;
}
ComputeBuffer<int> bufferMatrixA = new ComputeBuffer<int>(context,
ComputeMemoryFlags.UseHostPointer, inMatrixA);
ComputeBuffer<int> bufferMatrixB = new ComputeBuffer<int>(context,
ComputeMemoryFlags.UseHostPointer, inMatrixB);
ComputeBuffer<int> bufferMatrixC = new ComputeBuffer<int>(context,
ComputeMemoryFlags.UseHostPointer, outMatrixC);
long localWorkSize = Math.Min(device.MaxComputeUnits, sideSize);
//Sets arguments
kernel.SetMemoryArgument(0, bufferMatrixA);
kernel.SetMemoryArgument(1, bufferMatrixB);
kernel.SetMemoryArgument(2, bufferMatrixC);
kernel.SetLocalArgument(3, sideSize * 2);
kernel.SetValueArgument<int>(4, sideSize);
//kernel.SetLocalArgument(1, localWorkSize);
string offset = " ";
for (int x = 0; x < sideSize; x++)
offset += " ";
if (sideSize <= 32)
for (int y = 0; y < sideSize; y++)
{
Console.Write(offset);
for (int x = 0; x < sideSize; x++)
Console.Write(inMatrixA[y * sideSize + x] + " ");
Console.WriteLine();
}
//Runs commands
ComputeCommandQueue commands = new ComputeCommandQueue(context,
context.Devices[0], ComputeCommandQueueFlags.None);
long executionTime = DateTime.Now.Ticks;
//Execute kernel
//globalWorkSize in increments of localWorkSize (max of device.MaxComputeUnits or kernel.GetWorkGroupSize())
commands.Execute(kernel, null,
new long[] { Math.Min(sideSize, 16), Math.Min(sideSize, 16) },
new long[] { localWorkSize, 1 }, null);
//globalWorkSize can be any size
//localWorkSize product much not be greater than device.MaxComputeUnits
//and it must not be greater than kernel.GetWorkGroupSize()
//ESSENTIALLY, the program iterates through globalWorkSize
//in increments of localWorkSize. Both are multidimensional,
//but this just saves us the time of doing that
//(1 dimension can be put to multiple if the max dimension lengths
//are known very easily with remainder).
//Also, you should probably use this
//kernel.GetPreferredWorkGroupSizeMultiple(device);
commands.Finish();
commands.ReadFromBuffer(bufferMatrixC, ref outMatrixC, true, null);
commands.Finish();
executionTime = DateTime.Now.Ticks - executionTime;
GC.Collect();
program.Dispose();
Console.WriteLine();
if (sideSize <= 32)
for (int y = 0; y < sideSize; y++)
{
for (int x = 0; x < sideSize; x++)
Console.Write(inMatrixB[y * sideSize + x] + " ");
Console.Write(" ");
for (int x = 0; x < sideSize; x++)
Console.Write(outMatrixC[y * sideSize + x] + " ");
Console.WriteLine();
}
int testY = random.Next(sideSize);
int testX = random.Next(sideSize);
int sum = 0;
for (int q = 0; q < sideSize; q++)
sum += inMatrixA[q * sideSize + testX] *
inMatrixB[testY * sideSize + q];
Console.WriteLine(sum == outMatrixC[testY * sideSize + testX]);
Console.WriteLine(executionTime / 10000.0);
}
public ComputeKernel CreateKernel(object kernelInstance)
{
string kernelName = kernelInstance.GetType().Name;
if (HardwareAccelerationEnabled)
{
IKernel program = KernelManager.LoadKernel(kernelName);
// Create and build the opencl program.
var computeProgram = new ComputeProgram(_context, program.Code);
computeProgram.Build(null, null, null, IntPtr.Zero);
// Create the kernel function and set its arguments.
ComputeKernel kernel = computeProgram.CreateKernel("Run");
int index = 0;
foreach (string key in _intComputeBuffers.Keys)
{
kernel.SetMemoryArgument(index, _intComputeBuffers[key]);
index++;
}
foreach (string key in _floatComputeBuffers.Keys)
{
kernel.SetMemoryArgument(index, _floatComputeBuffers[key]);
index++;
}
return kernel;
}
return null;
}
private void CalculateConvolution(ComputeContext computeContext)
{
Stopwatch stopwatch = new Stopwatch();
stopwatch.Start();
float dx;
bool shiftXParse = float.TryParse(textBoxShiftX.Text, NumberStyles.Float, CultureInfo.InvariantCulture.NumberFormat, out dx);
if (!shiftXParse)
throw new SyntaxErrorException(", needs to be .");
float dy;
bool shiftYParse = float.TryParse(textBoxShiftX.Text, NumberStyles.Float, CultureInfo.InvariantCulture.NumberFormat, out dy);
if (!shiftYParse)
throw new SyntaxErrorException(", needs to be .");
float dz;
bool shiftZParse = float.TryParse(textBoxShiftX.Text, NumberStyles.Float, CultureInfo.InvariantCulture.NumberFormat, out dz);
if (!shiftZParse)
throw new SyntaxErrorException(", needs to be .");
int pixelCount = _imageDimensionX*_imageDimensionY*_imageDimensionZ;
Console.WriteLine("Computing...");
Console.WriteLine("Reading kernel...");
String kernelPath = Directory.GetParent(Directory.GetCurrentDirectory()).Parent.Parent.FullName;
String kernelString;
using (var sr = new StreamReader(kernelPath + "\\convolution.cl"))
kernelString = sr.ReadToEnd();
Console.WriteLine("Reading kernel... done");
float[] selectedTransformation = Transformations.GetTransformation((TransformationType)comboBoxTransform.SelectedItem, 1.0f / float.Parse(textBoxPixelSize.Text), 1.0f / float.Parse(textBoxPixelSize.Text), 1.0f / float.Parse(textBoxPixelSize.Text), dx, dy, dz);
//create openCL program
ComputeProgram computeProgram = new ComputeProgram(computeContext, kernelString);
computeProgram.Build(computeContext.Devices, null, null, IntPtr.Zero);
ComputeProgramBuildStatus computeProgramBuildStatus = computeProgram.GetBuildStatus(_selectedComputeDevice);
Console.WriteLine("computeProgramBuildStatus\n\t"+computeProgramBuildStatus);
String buildLog = computeProgram.GetBuildLog(_selectedComputeDevice);
Console.WriteLine("buildLog");
if (buildLog.Equals("\n"))
Console.WriteLine("\tbuildLog is empty...");
else
Console.WriteLine("\t" + buildLog);
float[] fluorophores = CsvData.ReadFluorophores(_sourceFilename);
/////////////////////////////////////////////
// Create a Command Queue & Event List
/////////////////////////////////////////////
ComputeCommandQueue computeCommandQueue = new ComputeCommandQueue(computeContext, _selectedComputeDevice, ComputeCommandQueueFlags.None);
////////////////////////////////////////////////////////////////
// Create Buffers Transform
////////////////////////////////////////////////////////////////
ComputeBuffer<float> fluorophoresCoords = new ComputeBuffer<float>(computeContext, ComputeMemoryFlags.ReadWrite, fluorophores.LongLength);
ComputeBuffer<float> transformationMatrix = new ComputeBuffer<float>(computeContext, ComputeMemoryFlags.ReadOnly, selectedTransformation.LongLength);
/////////////////////////////////////////////
// Create the transformFluorophoresKernel
///////////////////////////////////////////////////////////
ComputeKernel transformFluorophoresKernel = computeProgram.CreateKernel("transform_fluorophores");
/////////////////////////////////////////////
// Set the transformFluorophoresKernel arguments
/////////////////////////////////////////////
transformFluorophoresKernel.SetMemoryArgument(0, fluorophoresCoords);
transformFluorophoresKernel.SetMemoryArgument(1, transformationMatrix);
/////////////////////////////////////////////
// Configure the work-item structure
/////////////////////////////////////////////
long[] globalWorkOffsetTransformFluorophoresKernel = null;
long[] globalWorkSizeTransformFluorophoresKernel = new long[] { fluorophores.Length / 4 };
long[] localWorkSizeTransformFluorophoresKernel = null;
////////////////////////////////////////////////////////
// Enqueue the transformFluorophoresKernel for execution
////////////////////////////////////////////////////////
computeCommandQueue.WriteToBuffer(fluorophores, fluorophoresCoords, true, null);
computeCommandQueue.WriteToBuffer(selectedTransformation, transformationMatrix, true, null);
computeCommandQueue.Execute(transformFluorophoresKernel, globalWorkOffsetTransformFluorophoresKernel, globalWorkSizeTransformFluorophoresKernel, localWorkSizeTransformFluorophoresKernel, null);
// computeCommandQueue.ExecuteTask(transformFluorophoresKernel, transformFluorophoresEvents);
float[] transformedFluorophores = new float[fluorophores.Length];
computeCommandQueue.ReadFromBuffer(fluorophoresCoords, ref transformedFluorophores, true, null);
computeCommandQueue.Finish();
//TODO remove, only for testing
// for (int i = 0; i < transformedFluorophores.Length; i++)
// {
// Console.WriteLine(transformedFluorophores[i]);
// }
// /TODO remove, only for testing
stopwatch.Stop();
Console.WriteLine("Transform fluophores duration:\n\t" + stopwatch.Elapsed);
stopwatch.Reset();
stopwatch.Start();
// fluorophoresCoords are now transformed (done in place)
////////////////////////////////////////////////////////////////
// Create Buffers Convolve Fluorophores
////////////////////////////////////////////////////////////////
const int convolve_kernel_lwgs = 16;
int totalBuffer = (int) Math.Ceiling(pixelCount / (float)convolve_kernel_lwgs) * convolve_kernel_lwgs;
ComputeBuffer<float> resultImage = new ComputeBuffer<float>(computeContext, ComputeMemoryFlags.WriteOnly, totalBuffer);
/////////////////////////////////////////////
// Create the transformFluorophoresKernel
/////////////////////////////////////////////
ComputeKernel convolveFluorophoresKernel = computeProgram.CreateKernel("convolve_fluorophores");
/////////////////////////////////////////////
// Set the convolveFluorophoresKernel arguments
/////////////////////////////////////////////
convolveFluorophoresKernel.SetMemoryArgument(0, resultImage);
convolveFluorophoresKernel.SetValueArgument(1, _imageDimensionX);
convolveFluorophoresKernel.SetValueArgument(2, _imageDimensionY);
convolveFluorophoresKernel.SetMemoryArgument(3, fluorophoresCoords);
convolveFluorophoresKernel.SetLocalArgument(4, convolve_kernel_lwgs);
convolveFluorophoresKernel.SetValueArgument(5, fluorophores.Length / 4);
/////////////////////////////////////////////
// Configure the work-item structure
/////////////////////////////////////////////
long[] globalWorkOffsetTransformConvolveFluorophoresKernel = null;
long[] globalWorkSizeTransformConvolveFluorophoresKernel = new long[] { pixelCount };
long[] localWorkSizeTransformConvolveFluorophoresKernel = new long[] {convolve_kernel_lwgs};
////////////////////////////////////////////////////////
// Enqueue the convolveFluorophoresKernel for execution
////////////////////////////////////////////////////////
computeCommandQueue.Execute(convolveFluorophoresKernel, globalWorkOffsetTransformConvolveFluorophoresKernel, globalWorkSizeTransformConvolveFluorophoresKernel, localWorkSizeTransformConvolveFluorophoresKernel, null);
float[] resultImageData = new float[totalBuffer];
computeCommandQueue.ReadFromBuffer(resultImage, ref resultImageData, true, null);
computeCommandQueue.Finish();
for (int i = 0; i < pixelCount; i++)
{
Console.WriteLine(resultImageData[i]);
}
Console.WriteLine("Writing data to file...");
// CsvData.WriteToDisk("..\\..\\..\\output.csv", resultImageData);
TiffData.WriteToDisk(resultImageData, _saveFilename, _imageDimensionX, _imageDimensionY);
Bitmap bitmap = new Bitmap(_imageDimensionX, _imageDimensionY);
float max = resultImageData.Max();
float scale = 255/(float)max;
// for (int r = 0; r < _imageDimensionY; r++)
// {
// for (int c = 0; c < _imageDimensionX; c++)
// {
// float value = resultImageData[c*(r + 1)];
// Color newColor = Color.FromArgb((int)(value * scale), (int)(value * scale), (int)(value * scale));
// bitmap.SetPixel(c,r, newColor);
// }
// }
ushort[] ushortdata = new ushort[resultImageData.Length];
for (int i = 0; i < resultImageData.Length; i++)
{
ushortdata[i] = (ushort)resultImageData[i];
}
uint[] convertGray16ToRgb = ConvertGray16ToRGB(ushortdata, 16);
byte[] bytes = new byte[convertGray16ToRgb.Length * 4];
//
// int[] resultImageData2 = new int[resultImageData.Length];
//
for (int index = 0; index < convertGray16ToRgb.Length; index++)
{
// resultImageData2[index] = (int)(scale*resultImageData[index]);
byte[] bytes1 = BitConverter.GetBytes(convertGray16ToRgb[index]);
bytes[index] = bytes1[0];
bytes[4 * index + 1] = bytes1[1];
bytes[4 * index + 2] = bytes1[2];
bytes[4 * index + 3] = bytes1[3];
}
//
// for (int r = 0; r < _imageDimensionY; r++)
// {
// for (int c = 0; c < _imageDimensionX; c++)
// {
// float value = resultImageData2[c*(r + 1)];
// Color newColor = Color.FromArgb((int)(value), (int)(value), (int)(value));
// bitmap.SetPixel(c,r, newColor);
// }
// }
// bitmap.Save("c:\\temp.bmp");
using (MemoryStream ms = new MemoryStream(bytes))
{
Image image = Bitmap.FromStream(ms);
image.Save("c:\\temp.bmp");
}
Console.WriteLine("Writing data to file... done");
stopwatch.Stop();
Console.WriteLine("Convolve fluophores duration:\n\t" + stopwatch.Elapsed);
Console.WriteLine("Computing... done");
}
/// <summary>
/// Compile the kernel with a map of preprocessor defines, a collection of
/// name-value pairs.
/// </summary>
///
/// <param name="options">A map of preprocessor defines.</param>
public void Compile(IDictionary<String, String> options)
{
// clear out any old program
if (this.program != null)
{
this.program.Dispose();
this.kernel.Dispose();
}
// Create the program from the source code
this.program = new ComputeProgram(this.context, this.cl);
if (options.Count > 0)
{
StringBuilder builder = new StringBuilder();
/* foreach */
foreach (KeyValuePair<String, String> obj in options)
{
if (builder.Length > 0)
{
builder.Append(" ");
}
builder.Append("-D ");
builder.Append(obj.Key);
builder.Append("=");
builder.Append(obj.Value);
}
program.Build(null, builder.ToString(), null, IntPtr.Zero);
}
else
{
program.Build(null, null, null, IntPtr.Zero);
}
// Create the kernel
this.kernel = Program.CreateKernel(this.kernelName);
}
/// <summary>
/// Attempts to initialize OpenCL for the selected GPU.
/// </summary>
private void InitializeOpenCL()
{
// only initialize once
if (clKernel != null)
return;
// select the device we've been instructed to use
clDevice = ComputePlatform.Platforms
.SelectMany(i => i.Devices)
.SingleOrDefault(i => i.Handle.Value == Gpu.CLDeviceHandle.Value);
// context we'll be working underneath
clContext = new ComputeContext(new ComputeDevice[] { clDevice }, new ComputeContextPropertyList(clDevice.Platform), null, IntPtr.Zero);
// queue to control device
clQueue = new ComputeCommandQueue(clContext, clDevice, ComputeCommandQueueFlags.None);
// buffers to store kernel output
clBuffer0 = new ComputeBuffer<uint>(clContext, ComputeMemoryFlags.ReadOnly, 16);
clBuffer1 = new ComputeBuffer<uint>(clContext, ComputeMemoryFlags.ReadOnly, 16);
// kernel code
string kernelCode;
using (var rdr = new StreamReader(GetType().Assembly.GetManifestResourceStream("BitMaker.Miner.Gpu.DiabloMiner.cl")))
kernelCode = rdr.ReadToEnd();
clProgram = new ComputeProgram(clContext, kernelCode);
try
{
// build kernel for device
clProgram.Build(new ComputeDevice[] { clDevice }, "-D WORKSIZE=" + clDevice.MaxWorkGroupSize, null, IntPtr.Zero);
}
catch (ComputeException)
{
throw new Exception(clProgram.GetBuildLog(clDevice));
}
clKernel = clProgram.CreateKernel("search");
}
private void buildProgramMenuItem_Click(object sender, EventArgs e)
{
if (editorTextBox.Text.Length == 0)
{
logTextBox.Text = "No source.";
return;
}
string[] logContent;
ComputeContextPropertyList properties = new ComputeContextPropertyList(configForm.Platform);
ComputeContext context = new ComputeContext(configForm.Devices, properties, null, IntPtr.Zero);
ComputeProgram program = new ComputeProgram(context, editorTextBox.Text);
try
{
program.Build(configForm.Devices, configForm.Options, null, IntPtr.Zero);
logContent = new string[] { "Build succeeded." };
}
catch (Exception exception)
{
List<string> lineList = new List<string>();
foreach (ComputeDevice device in context.Devices)
{
string header = "PLATFORM: " + configForm.Platform.Name + ", DEVICE: " + device.Name;
lineList.Add(header);
StringReader reader = new StringReader(program.GetBuildLog(device));
string line = reader.ReadLine();
while (line != null)
{
lineList.Add(line);
line = reader.ReadLine();
}
lineList.Add("");
lineList.Add(exception.Message);
}
logContent = lineList.ToArray();
}
logTextBox.Lines = logContent;
}
protected override void RunInternal()
{
int count = 10;
float[] arrA = new float[count];
float[] arrB = new float[count];
float[] arrC = new float[count];
Random rand = new Random();
for (int i = 0; i < count; i++)
{
arrA[i] = (float)(rand.NextDouble() * 100);
arrB[i] = (float)(rand.NextDouble() * 100);
}
ComputeBuffer<float> a = new ComputeBuffer<float>(context, ComputeMemoryFlags.ReadOnly | ComputeMemoryFlags.CopyHostPointer, arrA);
ComputeBuffer<float> b = new ComputeBuffer<float>(context, ComputeMemoryFlags.ReadOnly | ComputeMemoryFlags.CopyHostPointer, arrB);
ComputeBuffer<float> c = new ComputeBuffer<float>(context, ComputeMemoryFlags.WriteOnly, arrC.Length);
ComputeProgram program = new ComputeProgram(context, new string[] { kernelSource });
program.Build(null, null, null, IntPtr.Zero);
ComputeKernel kernel = program.CreateKernel("VectorAdd");
kernel.SetMemoryArgument(0, a);
kernel.SetMemoryArgument(1, b);
kernel.SetMemoryArgument(2, c);
ComputeCommandQueue commands = new ComputeCommandQueue(context, context.Devices[0], ComputeCommandQueueFlags.None);
ComputeEventList events = new ComputeEventList();
commands.Execute(kernel, null, new long[] { count }, null, events);
arrC = new float[count];
GCHandle arrCHandle = GCHandle.Alloc(arrC, GCHandleType.Pinned);
commands.Read(c, false, 0, count, arrCHandle.AddrOfPinnedObject(), events);
commands.Finish();
arrCHandle.Free();
for (int i = 0; i < count; i++)
Console.WriteLine("{0} + {1} = {2}", arrA[i], arrB[i], arrC[i]);
}
public TerrainGen()
{
#if CPU_DEBUG
var platform = ComputePlatform.Platforms[1];
#else
var platform = ComputePlatform.Platforms[0];
#endif
_devices = new List<ComputeDevice>();
_devices.Add(platform.Devices[0]);
_properties = new ComputeContextPropertyList(platform);
_context = new ComputeContext(_devices, _properties, null, IntPtr.Zero);
_cmdQueue = new ComputeCommandQueue(_context, _devices[0], ComputeCommandQueueFlags.None);
#region setup generator kernel
bool loadFromSource = Gbl.HasRawHashChanged[Gbl.RawDir.Scripts];
loadFromSource = true;
_chunkWidthInBlocks = Gbl.LoadContent<int>("TGen_ChunkWidthInBlocks");
_chunkWidthInVerts = _chunkWidthInBlocks + 1;
_blockWidth = Gbl.LoadContent<int>("TGen_BlockWidthInMeters");
float lacunarity = Gbl.LoadContent<float>("TGen_Lacunarity");
float gain = Gbl.LoadContent<float>("TGen_Gain");
int octaves = Gbl.LoadContent<int>("TGen_Octaves");
float offset = Gbl.LoadContent<float>("TGen_Offset");
float hScale = Gbl.LoadContent<float>("TGen_HScale");
float vScale = Gbl.LoadContent<float>("TGen_VScale");
_genConstants = new ComputeBuffer<float>(_context, ComputeMemoryFlags.ReadOnly, 8);
var genArr = new[]{
lacunarity,
gain,
offset,
octaves,
hScale,
vScale,
_blockWidth,
_chunkWidthInBlocks
};
_cmdQueue.WriteToBuffer(genArr, _genConstants, false, null);
if (loadFromSource){
_generationPrgm = new ComputeProgram(_context, Gbl.LoadScript("TGen_Generator"));
#if CPU_DEBUG
_generationPrgm.Build(null, @"-g -s D:\Projects\Gondola\Scripts\GenTerrain.cl", null, IntPtr.Zero); //use option -I + scriptDir for header search
#else
_generationPrgm.Build(null, "", null, IntPtr.Zero);//use option -I + scriptDir for header search
#endif
Gbl.SaveBinary(_generationPrgm.Binaries, "TGen_Generator");
}
else{
var binary = Gbl.LoadBinary("TGen_Generator");
_generationPrgm = new ComputeProgram(_context, binary, _devices);
_generationPrgm.Build(null, "", null, IntPtr.Zero);
}
//loadFromSource = false;
_terrainGenKernel = _generationPrgm.CreateKernel("GenTerrain");
_normalGenKernel = _generationPrgm.CreateKernel("GenNormals");
//despite the script using float3 for these fields, we need to consider it to be float4 because the
//implementation is basically a float4 wrapper that uses zero for the last variable
_geometry = new ComputeBuffer<float>(_context, ComputeMemoryFlags.None, _chunkWidthInVerts*_chunkWidthInVerts*4);
_normals = new ComputeBuffer<ushort>(_context, ComputeMemoryFlags.None, _chunkWidthInVerts * _chunkWidthInVerts * 4);
_binormals = new ComputeBuffer<byte>(_context, ComputeMemoryFlags.None, _chunkWidthInVerts*_chunkWidthInVerts*4);
_tangents = new ComputeBuffer<byte>(_context, ComputeMemoryFlags.None, _chunkWidthInVerts*_chunkWidthInVerts*4);
_uvCoords = new ComputeBuffer<float>(_context, ComputeMemoryFlags.None, _chunkWidthInVerts*_chunkWidthInVerts*2);
_terrainGenKernel.SetMemoryArgument(0, _genConstants);
_terrainGenKernel.SetMemoryArgument(3, _geometry);
_terrainGenKernel.SetMemoryArgument(4, _uvCoords);
_normalGenKernel.SetMemoryArgument(0, _genConstants);
_normalGenKernel.SetMemoryArgument(3, _geometry);
_normalGenKernel.SetMemoryArgument(4, _normals);
_normalGenKernel.SetMemoryArgument(5, _binormals);
_normalGenKernel.SetMemoryArgument(6, _tangents);
#endregion
#region setup quadtree kernel
if (loadFromSource){
_qTreePrgm = new ComputeProgram(_context, Gbl.LoadScript("TGen_QTree"));
#if CPU_DEBUG
_qTreePrgm.Build(null, @"-g -s D:\Projects\Gondola\Scripts\Quadtree.cl", null, IntPtr.Zero);
#else
_qTreePrgm.Build(null, "", null, IntPtr.Zero);
#endif
Gbl.SaveBinary(_qTreePrgm.Binaries, "TGen_QTree");
}
else{
var binary = Gbl.LoadBinary("TGen_QTree");
_qTreePrgm = new ComputeProgram(_context, binary, _devices);
_qTreePrgm.Build(null, "", null, IntPtr.Zero);
}
_qTreeKernel = _qTreePrgm.CreateKernel("QuadTree");
_crossCullKernel = _qTreePrgm.CreateKernel("CrossCull");
_activeVerts = new ComputeBuffer<byte>(_context, ComputeMemoryFlags.None, _chunkWidthInVerts*_chunkWidthInVerts);
_dummy = new ComputeBuffer<int>(_context, ComputeMemoryFlags.None, 50);
var rawNormals = new ushort[_chunkWidthInVerts * _chunkWidthInVerts * 4];
_emptyVerts = new byte[_chunkWidthInVerts*_chunkWidthInVerts];
for (int i = 0; i < _emptyVerts.Length; i++){
_emptyVerts[i] = 1;
}
_cmdQueue.WriteToBuffer(rawNormals, _normals, true, null);
_cmdQueue.WriteToBuffer(_emptyVerts, _activeVerts, true, null);
_qTreeKernel.SetValueArgument(1, _chunkWidthInBlocks);
_qTreeKernel.SetMemoryArgument(2, _normals);
_qTreeKernel.SetMemoryArgument(3, _activeVerts);
_qTreeKernel.SetMemoryArgument(4, _dummy);
_crossCullKernel.SetValueArgument(1, _chunkWidthInBlocks);
_crossCullKernel.SetMemoryArgument(2, _normals);
_crossCullKernel.SetMemoryArgument(3, _activeVerts);
_crossCullKernel.SetMemoryArgument(4, _dummy);
#endregion
#region setup winding kernel
if (loadFromSource){
_winderPrgm = new ComputeProgram(_context, Gbl.LoadScript("TGen_VertexWinder"));
#if CPU_DEBUG
_winderPrgm.Build(null, @"-g -s D:\Projects\Gondola\Scripts\VertexWinder.cl", null, IntPtr.Zero);
#else
_winderPrgm.Build(null, "", null, IntPtr.Zero);
#endif
Gbl.SaveBinary(_winderPrgm.Binaries, "TGen_VertexWinder");
}
else{
var binary = Gbl.LoadBinary("TGen_VertexWinder");
_winderPrgm = new ComputeProgram(_context, binary, _devices);
_winderPrgm.Build(null, "", null, IntPtr.Zero);
}
_winderKernel = _winderPrgm.CreateKernel("VertexWinder");
_indicies = new ComputeBuffer<int>(_context, ComputeMemoryFlags.None, (_chunkWidthInBlocks)*(_chunkWidthInBlocks)*8);
_winderKernel.SetMemoryArgument(0, _activeVerts);
_winderKernel.SetMemoryArgument(1, _indicies);
_emptyIndices = new int[(_chunkWidthInBlocks)*(_chunkWidthInBlocks)*8];
for (int i = 0; i < (_chunkWidthInBlocks)*(_chunkWidthInBlocks)*8; i++){
_emptyIndices[i] = 0;
}
_cmdQueue.WriteToBuffer(_emptyIndices, _indicies, true, null);
#endregion
if (loadFromSource){
Gbl.AllowMD5Refresh[Gbl.RawDir.Scripts] = true;
}
_cmdQueue.Finish();
}
public void Start(IMinerContext context)
{
string code;
var kernelRes = Assembly.GetExecutingAssembly().GetManifestResourceStream("BitMaker.Miner.Cloo.Miner.cl");
using (var rdr = new StreamReader(kernelRes))
code = clProgramSource;
var platform = ComputePlatform.Platforms[0];
var properties = new ComputeContextPropertyList(platform);
device = platform.Devices[0];
ccontext = new ComputeContext(platform.Devices, properties, null, IntPtr.Zero);
program = new ComputeProgram(ccontext, clProgramSource);
program.Build(null, null, notify, IntPtr.Zero);
}