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) + "...");
}
internal OpenCLProgram(OpenCLContext context, ComputeProgram program)
{
Contract.Requires(context != null);
Contract.Requires(program != null);
Context = context;
ComputeProgram = program;
}
public virtual void Initialize() {
platform = ComputePlatform.Platforms[0];
device = platform.Devices[0];
properties = new ComputeContextPropertyList(platform);
context = new ComputeContext(new[] { device }, properties, null, IntPtr.Zero);
program = new ComputeProgram(context, KernelSrc);
}
internal ComputeKernel(CLKernelHandle handle, ComputeProgram program)
{
Handle = handle;
SetID(Handle.Value);
context = program.Context;
functionName = GetStringInfo <CLKernelHandle, ComputeKernelInfo>(Handle, ComputeKernelInfo.FunctionName, CLInterface.CL12.GetKernelInfo);
this.program = program;
}
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");
}
////////////////////////////////////////////////////////////////////////////////////////////////////
/// <summary> Initializes a new instance of the Cloo.ComputeKernel class. </summary>
///
/// <param name="handle"> The handle of the <see cref="ComputeKernel"/>. </param>
/// <param name="program"> Gets the <see cref="ComputeProgram"/> that the
/// <see cref="ComputeKernel"/> belongs to. </param>
////////////////////////////////////////////////////////////////////////////////////////////////////
internal ComputeKernel(CLKernelHandle handle, ComputeProgram program)
{
Handle = handle;
SetID(Handle.Value);
context = program.Context;
functionName = GetStringInfo <CLKernelHandle, ComputeKernelInfo>(Handle, ComputeKernelInfo.FunctionName, CL12.GetKernelInfo);
this.program = program;
Trace.WriteLine("Create " + this + " in Thread(" + Thread.CurrentThread.ManagedThreadId + ").", "Information");
}
////////////////////////////////////////////////////////////////////////////////////////////////////
/// <summary> Initializes a new instance of the Cloo.ComputeKernel class. </summary>
///
/// <param name="handle"> The handle of the <see cref="ComputeKernel"/>. </param>
/// <param name="program"> Gets the <see cref="ComputeProgram"/> that the
/// <see cref="ComputeKernel"/> belongs to. </param>
////////////////////////////////////////////////////////////////////////////////////////////////////
internal ComputeKernel(CLKernelHandle handle, ComputeProgram program)
{
Handle = handle;
SetID(Handle.Value);
context = program.Context;
functionName = GetStringInfo <CLKernelHandle, ComputeKernelInfo>(Handle, ComputeKernelInfo.FunctionName, CL12.GetKernelInfo);
this.program = program;
RILogManager.Default?.SendTrace(string.Intern("Create ") + this + string.Intern(" in Thread(") + Thread.CurrentThread.ManagedThreadId + string.Intern(")."), string.Intern("Information"));
}
internal ComputeKernel(string functionName, ComputeProgram program)
{
ComputeErrorCode error = ComputeErrorCode.Success;
Handle = CLInterface.CL12.CreateKernel(program.Handle, functionName, out error);
ComputeException.ThrowOnError(error);
SetID(Handle.Value);
context = program.Context;
this.functionName = functionName;
this.program = program;
}
////////////////////////////////////////////////////////////////////////////////////////////////////
/// <summary> Initializes a new instance of the Cloo.ComputeKernel class. </summary>
///
/// <param name="functionName"> Gets the function name of the <see cref="ComputeKernel"/>. </param>
/// <param name="program"> Gets the <see cref="ComputeProgram"/> that the
/// <see cref="ComputeKernel"/> belongs to. </param>
////////////////////////////////////////////////////////////////////////////////////////////////////
internal ComputeKernel(string functionName, ComputeProgram program)
{
Handle = CL12.CreateKernel(program.Handle, functionName, out var error);
ComputeException.ThrowOnError(error);
SetID(Handle.Value);
context = program.Context;
this.functionName = functionName;
this.program = program;
RILogManager.Default?.SendTrace(string.Intern("Create ") + this + string.Intern(" in Thread(") + Thread.CurrentThread.ManagedThreadId + string.Intern(")."), string.Intern("Information"));
}
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());
}
}
////////////////////////////////////////////////////////////////////////////////////////////////////
/// <summary> Initializes a new instance of the Cloo.ComputeKernel class. </summary>
///
/// <param name="functionName"> Gets the function name of the <see cref="ComputeKernel"/>. </param>
/// <param name="program"> Gets the <see cref="ComputeProgram"/> that the
/// <see cref="ComputeKernel"/> belongs to. </param>
////////////////////////////////////////////////////////////////////////////////////////////////////
internal ComputeKernel(string functionName, ComputeProgram program)
{
ComputeErrorCode error = ComputeErrorCode.Success;
Handle = CL12.CreateKernel(program.Handle, functionName, out error);
ComputeException.ThrowOnError(error);
SetID(Handle.Value);
context = program.Context;
this.functionName = functionName;
this.program = program;
Trace.WriteLine("Create " + this + " in Thread(" + Thread.CurrentThread.ManagedThreadId + ").", "Information");
}
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 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());
}
}
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);
}
/// <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 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 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");
}
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());
}
}
/// <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>
/// Entry point for a standard work thread.
/// </summary>
private void WorkThread()
{
InitializeOpenCL();
try
{
// continue working until canceled
while (!cts.IsCancellationRequested)
Work(Context.GetWork(this, GetType().Name));
}
catch (OperationCanceledException)
{
// ignore
}
clQueue.Finish();
clKernel.Dispose();
clKernel = null;
clBuffer0.Dispose();
clBuffer0 = null;
clBuffer1.Dispose();
clBuffer1 = null;
clQueue.Dispose();
clQueue = null;
clDevice = null;
clProgram.Dispose();
clProgram = null;
clContext.Dispose();
clContext = null;
}
public OpenCLCalculator(ComputeContext context, ComputeProgram prg, ComputeKernel krnl)
{
_context = context;
_prg = prg;
_krnl = krnl;
}
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;
}
}
/// <summary>
/// Initializes local fields and the underlying compute context.
/// </summary>
public void Initialize()
{
if (this.context == null)
{
var devices = ComputePlatform.Platforms.SelectMany(a => a.Devices).Where(a => a.Extensions.Contains("cl_khr_fp64")).Take(1).ToArray();
ComputeContextPropertyList list = new ComputeContextPropertyList(devices[0].Platform);
this.context = new ComputeContext(devices, list, null, IntPtr.Zero);
}
this.program = new ComputeProgram(this.context, File.ReadAllText("Mandelbrot.cl"));
this.program.Build(null, null, null, IntPtr.Zero);
this.mandelbrot = this.program.CreateKernel("Mandelbrot");
this.toBitmap = this.program.CreateKernel("ToBitmap");
this.resultBuffer = new ComputeBuffer<int>(this.context, ComputeMemoryFlags.ReadWrite, this.ImageWidth * this.ImageHeight);
this.bitmapBuffer = new ComputeBuffer<byte>(this.context, ComputeMemoryFlags.ReadWrite, this.ImageWidth * this.ImageHeight * 4);
this.mandelbrot.SetMemoryArgument(7, this.resultBuffer);
this.toBitmap.SetMemoryArgument(1, this.resultBuffer);
this.toBitmap.SetMemoryArgument(2, this.bitmapBuffer);
this.commandQueue = new ComputeCommandQueue(this.context, this.context.Devices.OrderBy(a => a.Type).Where(a => a.Extensions.Contains("cl_khr_fp64")).First(), ComputeCommandQueueFlags.None);
}
/// <summary>
/// Release this kernel.
/// </summary>
///
public virtual void Release()
{
if (this.program != null)
{
this.program.Dispose();
this.kernel.Dispose();
this.program = null;
this.kernel = null;
}
}
/// <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;
}
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);
}
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 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);
}