private static int[] ExecuteStepDpq(Polynomial x, Polynomial y)
{
if (x.Degree != y.Degree)
{
throw new InvalidOperationException("Only works for polynomials of same degree!");
}
var d = Math.Min(x.Degree, y.Degree);
var n = d + 1;
var dpq = new int[n * n];
CLCalc.InitCL();
CLCalc.Program.Compile(new[] { KernelCodeStepDpq });
var kernel = new CLCalc.Program.Kernel("StepDpq");
var nCl = new CLCalc.Program.Variable(new[] { n });
var xCl = new CLCalc.Program.Variable(x.Coefficients);
var yCl = new CLCalc.Program.Variable(y.Coefficients);
var dpqCl = new CLCalc.Program.Variable(dpq);
CLCalc.Program.MemoryObject[] args = { nCl, xCl, yCl, dpqCl };
var workers = new[] { 2 * n - 2 };
kernel.Execute(args, workers);
dpqCl.ReadFromDeviceTo(dpq);
return(dpq);
}
private void button6_Click(object sender, EventArgs e)
{
float[] x = new float[] { 0 };
float[] y = new float[] { 0 };
CLCalc.InitCL();
if (!string.IsNullOrWhiteSpace(textBox1.Text) && !string.IsNullOrWhiteSpace(textBox2.Text))
{
x[0] = (float)Convert.ToDouble(textBox1.Text.Replace('.', ','));
y[0] = (float)Convert.ToDouble(textBox2.Text.Replace('.', ','));
}
string s = @"
__kernel void
sum(global float4 *x, global float4 *y) {
x[0] = x[0] + y[0];
}";
CLCalc.Program.Compile(new string[] { s });
CLCalc.Program.Kernel sum = new CLCalc.Program.Kernel("sum");
CLCalc.Program.Variable varx = new CLCalc.Program.Variable(x);
CLCalc.Program.Variable vary = new CLCalc.Program.Variable(y);
CLCalc.Program.Variable[] args = { varx, vary };
int[] max = new int[] { 1 };
sum.Execute(args, max);
varx.ReadFromDeviceTo(x);
textBox3.Text = Convert.ToString(x[0]);
}
/// <summary>Static Constructor. Builds kernels</summary>
static SparseLinalg()
{
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.Unknown)
{
try { CLCalc.InitCL(); }
catch
{
}
}
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.UsingCL)
{
//Kernel
CLLinalgSrc src = new CLLinalgSrc();
CLCalc.Program.Compile(new string[] { src.srcDotProd, src.srcMatVecMult, src.srcLinConjGrad });
kernelDotProduct = new CLCalc.Program.Kernel("dotProd");
kernelSum = new CLCalc.Program.Kernel("sumDotProd");
kernelGetDotSum = new CLCalc.Program.Kernel("GetResp");
kernelSparseMatrixVecMult = new CLCalc.Program.Kernel("SparseMatrixVecMult");
//Linear solving
kernelInitRP = new CLCalc.Program.Kernel("InitRP");
kernelMultiplyAdd = new CLCalc.Program.Kernel("MultiplyAdd");
kernelCopyToTemp = new CLCalc.Program.Kernel("CopyToTemp");
}
}
/// <summary>Compiles code and initializes kernel for this svm stance</summary>
private void CLSVMInit()
{
if (CLResource == null)
{
CLResource = new int[0];
}
lock (CLResource)
{
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.Unknown)
{
CLCalc.InitCL();
}
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.UsingCL)
{
if (kernelComputeKernelRBF == null)
{
CLSVMSrc s = new CLSVMSrc();
CLCalc.Program.Compile(new string[] { s.srcKernels, s.srcFindMaxMinErr, s.srcMultClass });
//Kernel computation
kernelComputeKernelRBF = new CLCalc.Program.Kernel("ComputeKernelRBF");
kernelMaxErr = new CLCalc.Program.Kernel("maxErr");
kernelComputeMax = new CLCalc.Program.Kernel("computeMax");
kernelMinErr = new CLCalc.Program.Kernel("minErr");
kernelComputeMin = new CLCalc.Program.Kernel("computeMin");
kernelGetResp = new CLCalc.Program.Kernel("getResp");
//Update error
kernelUpdateErr = new CLCalc.Program.Kernel("UpdateErr");
//Multiple classification
kernelComputeMultiKernelRBF = new CLCalc.Program.Kernel("ComputeMultiKernelRBF");
kernelSumKernels = new CLCalc.Program.Kernel("SumKernels");
}
//Memory obbjects
//Find max/min
CLErrLen = new CLCalc.Program.Variable(new int[1]);
HostResp = new int[1];
CLResp = new CLCalc.Program.Variable(HostResp);
CLMaxMinErrs = new CLCalc.Program.Variable(new float[MAXMINWORKSIZE]);
CLMaxMinInds = new CLCalc.Program.Variable(new int[MAXMINWORKSIZE]);
//Update error
CLUpdtErrParams = new CLCalc.Program.Variable(new float[3]);
}
}
}
/// <summary>Creates a new filter from a given compiled kernel.</summary>
/// <param name="filterCode">Code to compile</param>
/// <param name="filterName">Filter name</param>
public CLFilter(CLCalc.Program.Kernel filterKernel, string filterName)
{
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.Unknown)
{
CLCalc.InitCL();
}
if (CLCalc.CLAcceleration != CLCalc.CLAccelerationType.UsingCL)
{
throw new Exception("OpenCL not available");
}
this.FilterName = filterName;
this.FilterKernel = filterKernel;
}
private void InitKernel()
{
if (kernelExtractFeatures == null)
{
CLExtractFeatSrc src = new CLExtractFeatSrc();
CLBracketRegionsSrc src2 = new CLBracketRegionsSrc();
CLCalc.Program.Compile(new string[] { src2.src, src.src });
kernelExtractFeatures = new CLCalc.Program.Kernel("ExtractFeatures");
kernelComputeFrameDiff = new CLCalc.Program.Kernel("ComputeFrameDiff");
kernelSegregateSkin = new CLCalc.Program.Kernel("SegregateSkin");
}
}
private static void InitKernels()
{
string s = new CLFFTSrc().s;
CLCalc.InitCL();
try
{
CLCalc.Program.Compile(s);
}
catch
{
}
kernelfft_radix16 = new CLCalc.Program.Kernel("fft_radix16");
kernelfft_radix4 = new CLCalc.Program.Kernel("fft_radix4");
kernelConjugate = new CLCalc.Program.Kernel("Conjugate");
CLp = new CLCalc.Program.Variable(new int[1]);
}
/// <summary>Creates a new filter from a given code. Filter kernel name has to be the same as kernel name
/// (disregarding spaces - Ex: Kernel is GaussianBlur and filter name is Gaussian Blur)</summary>
/// <param name="filterCode">Complete OpenCL kernel code to compile</param>
/// <param name="filterName">Filter name</param>
public CLFilter(string filterCode, string filterName)
{
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.Unknown)
{
CLCalc.InitCL(Cloo.ComputeDeviceTypes.Gpu);
}
if (CLCalc.CLAcceleration != CLCalc.CLAccelerationType.UsingCL)
{
throw new Exception("OpenCL not available");
}
CLCalc.Program.Compile(filterCode);
this.FilterName = filterName;
string kernelname = FilterName.Replace(" ", "").ToLower();
FilterKernel = new CLCalc.Program.Kernel(kernelname);
}
public void Initialize(BaseCameraApplication app)
{
try
{
CLCalc.Program.Compile(app.GetPrimaryDevice().GetPreprocessCode() + src);
}
catch (BuildProgramFailureComputeException ex)
{
System.Console.WriteLine(ex.Message);
Environment.Exit(1);
}
int width = app.GetPrimaryDevice().GetDepthImage().Width;
int height = app.GetPrimaryDevice().GetDepthImage().Height;
pointBuffer = CLCalc.Program.Variable.Create(new ComputeBuffer<float>(CLCalc.Program.Context, ComputeMemoryFlags.ReadWrite | ComputeMemoryFlags.CopyHostPointer, points = new float[width * height * 4]));
colorBuffer = CLCalc.Program.Variable.Create(new ComputeBuffer<float>(CLCalc.Program.Context, ComputeMemoryFlags.ReadWrite | ComputeMemoryFlags.CopyHostPointer, colors = new float[width * height * 4]));
kernelCopyImage = new CLCalc.Program.Kernel("CopyToPoinCloud");
}
private void button3_Click(object sender, EventArgs e)
{
float[] x = new float[] { 1, 2, 3, 0.123f };
float[] y = new float[] { 1, 2, 1, 1 };
string s = @"
kernel void
sum (global float4 * x, global float4 * y)
{
x[0] = x[0] + y[0];
}
";
CLCalc.Program.Compile(new string[] { s });
CLCalc.Program.Kernel sum = new CLCalc.Program.Kernel("sum");
CLCalc.Program.Variable varx = new CLCalc.Program.Variable(x);
CLCalc.Program.Variable vary = new CLCalc.Program.Variable(y);
CLCalc.Program.Variable[] args = { varx, vary };
int[] max = new int[] { 1 };
sum.Execute(args, max);
varx.ReadFromDeviceTo(x);
//float[] t = new float[1] { 0 };
//float[] pos = new float[] { 1, 2, 3, 4, 5, 6 };
//float[] vel = new float[] { 1, 0, 0, 0, 0, 0 };
//float[] forces = new float[] { 0, 1, 0, 0, 0, 0 };
//float[] masses = new float[] { 1, 1 };
//float[] colSizes = new float[] { 0.1f, 0.1f };
//CLCalc.InitCL();
//CLCalc.CLPrograms.floatBodyPhysics phys = new CLCalc.CLPrograms.floatBodyPhysics(10);
//CLCalc.CLPrograms.floatBodyPhysics phys2 = new CLCalc.CLPrograms.floatBodyPhysics(20);
//CLCalc.FinishCL();
}
/// <summary>Initialize kernels</summary>
private void InitKernels()
{
if (kernelComputeHistograms == null || CLN == null)
{
CLSrc src = new CLSrc();
CLCalc.Program.Compile(src.src);
kernelComputeHistograms = new CLCalc.Program.Kernel("ComputeHistogram");
kernelConsolidateHist = new CLCalc.Program.Kernel("ConsolidateHist");
kernelPerformNormalization = new CLCalc.Program.Kernel("PerformNormalization");
kernelComputeHue = new CLCalc.Program.Kernel("ComputeHue");
CLN = new CLCalc.Program.Variable(new int[] { NLumIntens });
CLWidth = new CLCalc.Program.Variable(new int[] { bmp.Width });
CLHeight = new CLCalc.Program.Variable(new int[] { bmp.Height });
CLbmp = new CLCalc.Program.Image2D(bmp);
CLNewBmp = new CLCalc.Program.Image2D(bmp);
}
}
private static void GPU()
{
string vecSum = @"
__kernel void
floatVectorSum(__global float * v1,
__global float * v2)
{
// Vector element index
int i = get_global_id(0);
v1[i] = v1[i] + v2[i];
}";
//Initializes OpenCL Platforms and Devices and sets everything up
CLCalc.InitCL();
//Compiles the source codes. The source is a string array because the user may want
//to split the source into many strings.
CLCalc.Program.Compile(new string[] { vecSum });
//Gets host access to the OpenCL floatVectorSum kernel
CLCalc.Program.Kernel VectorSum = new CLCalc.Program.Kernel("floatVectorSum");
//We want to sum 2000 numbers
int n = 20000000;
//Create vectors with 2000 numbers
float[] v1 = new float[n], v2 = new float[n];
//Creates population for v1 and v2
for (int i = 0; i < n; i++)
{
v1[i] = (float)i / 10;
v2[i] = -(float)i / 8;
}
//Creates vectors v1 and v2 in the device memory
CLCalc.Program.Variable varV1 = new CLCalc.Program.Variable(v1);
CLCalc.Program.Variable varV2 = new CLCalc.Program.Variable(v2);
//Arguments of VectorSum kernel
CLCalc.Program.Variable[] args = new CLCalc.Program.Variable[] { varV1, varV2 };
//How many workers will there be? We need "n", one for each element
//int[] workers = new int[1] { n };
//Execute the kernel
VectorSum.Execute(args, 20000000);
//Read device memory varV1 to host memory v1
varV1.ReadFromDeviceTo(v1);
}
private void button1_Click(object sender, EventArgs e)
{
CLCalc.InitCL();
double[] a = new double[] { 2, 147483647, 2, 7 };
double[] b = new double[] { 1, 2, 7, 4 };
double[] c = new double[4];
CLCalc.Program.Variable v1 = new CLCalc.Program.Variable(a);
CLCalc.Program.Variable v2 = new CLCalc.Program.Variable(b);
CLCalc.Program.Variable v3 = new CLCalc.Program.Variable(c);
CLCalc.CLPrograms.VectorSum VecSum = new CLCalc.CLPrograms.VectorSum();
CLCalc.CLPrograms.MinDifs Mdifs = new CLCalc.CLPrograms.MinDifs();
//string[] s = new string[] { VecSum.intVectorSum, VecSum.floatVectorSum };
string[] s = new string[] { VecSum.doubleVectorSum };
CLCalc.Program.Compile(s);
CLCalc.Program.Kernel k = new CLCalc.Program.Kernel("doubleVectorSum");
//CLCalc.Program.Kernel k2 = new CLCalc.Program.Kernel("intVectorSum");
//CLCalc.Program.Kernel k = new CLCalc.Program.Kernel("floatMinDifs");
CLCalc.Program.Variable[] vv = new CLCalc.Program.Variable[3] {
v1, v2, v3
};
int[] max = new int[1] {
a.Length
};
k.Execute(vv, max);
CLCalc.Program.Sync();
v3.ReadFromDeviceTo(c);
CLCalc.FinishCL();
}
private void button4_Click(object sender, EventArgs e)
{
nMassas = new int[1] {
10
};
//Integrador de EDO
float x0 = 0;
float[] y0 = new float[2 * nMassas[0]];
//y0[0]=posicao, y0[1] = velocidade
for (int i = 0; i < 2 * nMassas[0]; i += 2)
{
y0[i] = 1;
}
CLCalc.CLPrograms.floatODE46 ode46 = new CLCalc.CLPrograms.floatODE46(x0, 0.005f, y0, MassaMola);
//CLCalc.Program.DefaultCQ = 0;
//Retorno de derivadas
string[] s = new string[] { CLCalc.EnableDblSupport, floatDerivsMMola };
CLCalc.Program.Compile(s);
KernelDerivs = new CLCalc.Program.Kernel("floatDerivs");
System.Diagnostics.Stopwatch sw = new System.Diagnostics.Stopwatch();
sw.Start();
ode46.Integrate(20);
sw.Stop();
this.Text = sw.ElapsedMilliseconds.ToString();
Application.DoEvents();
float[] y = ode46.State;
ode46.Integrate(25);
y = ode46.State;
float[] yerr = ode46.AbsError;
float x = ode46.IndepVar;
}
public CollisionDetector()
{
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.Unknown)
{
try
{
CLCalc.InitCL();
}
catch
{
}
}
try
{
if (CLCalc.CLAcceleration != CLCalc.CLAccelerationType.UsingCL)
{
return;
}
CollisionDetector.DisplacementSource displacementSource = new CollisionDetector.DisplacementSource();
CollisionDetector.ExactCollisionSource exactCollisionSource = new CollisionDetector.ExactCollisionSource();
CollisionDetector.VertexCollisionSource vertexCollisionSource = new CollisionDetector.VertexCollisionSource();
CLCalc.Program.Compile(new string[4]
{
displacementSource.srcCalcRotacoes,
displacementSource.srcCalcTransl,
exactCollisionSource.srcExactCollision,
vertexCollisionSource.srcVertexCollision
});
this.kernelCalcRotacoes = new CLCalc.Program.Kernel("CalcRotacoes");
this.kernelCalcTransl = new CLCalc.Program.Kernel("CalcTransl");
this.kernelCalcExactCollision = new CLCalc.Program.Kernel("CalcExactCollision");
this.kernelCalcVertexCollision = new CLCalc.Program.Kernel("CalcVertexCollision");
}
catch (Exception ex)
{
int num = (int)MessageBox.Show(ex.ToString());
}
}
public void DoubleSumTest()
{
string text = File.ReadAllText("Examples/SumTest.cl");
CLCalc.InitCL();
CLCalc.Program.Compile(new string[] { text });
int count = 2000;
var a = new double[count];
var b = new double[count];
var ab = new double[count];
for (int i = 0; i < count; i++)
{
a[i] = i / 10.0;
b[i] = -i / 9.0;
}
using (CLCalc.Program.Kernel Kernel = new CLCalc.Program.Kernel("doubleVectorSum"))
{
using (CLCalc.Program.Variable
varA = new CLCalc.Program.Variable(a),
varB = new CLCalc.Program.Variable(b))
{
var args = new CLCalc.Program.Variable[] { varA, varB };
var workers = new int[1] {
count
};
Kernel.Execute(args, workers);
varA.ReadFromDeviceTo(ab);
}
}
for (int i = 0; i < count; i++)
{
Assert.AreEqual(-i / 90.0, ab[i], 1E-13);
}
}
public static Polynomial MultiplyKaratsubaOpenCl(Polynomial x, Polynomial y)
{
if (x.Degree != y.Degree)
{
throw new InvalidOperationException("Only works for polynomials of same degree!");
}
var d = Math.Min(x.Degree, y.Degree);
var n = d + 1;
var di = ExecuteStepDi(x, y);
var dpq = ExecuteStepDpq(x, y);
var z = new int[2 * n - 1];
z[0] = di[0];
z[2 * n - 2] = di[n - 1];
CLCalc.InitCL();
CLCalc.Program.Compile(new[] { KernelCodeStepKaratsuba });
var kernel = new CLCalc.Program.Kernel("StepKaratsuba");
var nCl = new CLCalc.Program.Variable(new[] { n });
var xCl = new CLCalc.Program.Variable(x.Coefficients);
var yCl = new CLCalc.Program.Variable(y.Coefficients);
var diCl = new CLCalc.Program.Variable(di);
var dpqCl = new CLCalc.Program.Variable(dpq);
var zCl = new CLCalc.Program.Variable(z);
CLCalc.Program.MemoryObject[] args = { nCl, xCl, yCl, diCl, dpqCl, zCl };
var workers = new[] { 2 * n - 3 };
kernel.Execute(args, workers);
zCl.ReadFromDeviceTo(z);
return(new Polynomial(z));
}
static Kernels()
{
try
{
CLCalc.Program.Compile(src);
CLCalc.Program.MemoryObject[] Args = new CLCalc.Program.MemoryObject[100];;
int globalWorkSize = 4;
// compile the kernels
KernelStart = new CLCalc.Program.Kernel("KernelStart");
vetTransl = new CLCalc.Program.Kernel("vetTransl");
// run kernel start
KernelStart.Execute(Args, globalWorkSize);
}
catch (NullReferenceException nre)
{
System.Console.WriteLine("" + nre);
}
// System.Diagnostics.Debug.WriteLine("Hello");
}
static DynamicShading()
{
#region OpenCL Source
#region Thresholding
string srcThresh = @"
__kernel void Threshold(__constant int * cfg,
__read_only image2d_t imgSrc,
__global uchar * byteInfo,
__write_only image2d_t imgThresh)
{
const sampler_t smp = CLK_NORMALIZED_COORDS_FALSE | //Natural coordinates
CLK_ADDRESS_CLAMP | //Clamp to zeros
CLK_FILTER_NEAREST; //Don't interpolate
int thresh = cfg[0];
int2 coord = (int2)(get_global_id(0),get_global_id(1));
uint4 pix = read_imageui(imgSrc, smp, coord);
int pixBW = (int)pix.x+(int)pix.y+(int)pix.z;
pixBW = pixBW > 3*thresh ? 255 : 0;
byteInfo[coord.x+get_global_size(0)*coord.y] = pixBW;
pix = (uint4)((uint)pixBW,(uint)pixBW,(uint)pixBW,255);
write_imageui(imgThresh,coord,pix);
}
__kernel void RestoreBlackPixels(__constant int * cfg,
__read_only image2d_t imgSrc,
__read_only image2d_t imgRender,
__write_only image2d_t dst)
{
const sampler_t smp = CLK_NORMALIZED_COORDS_FALSE | //Natural coordinates
CLK_ADDRESS_CLAMP | //Clamp to zeros
CLK_FILTER_NEAREST; //Don't interpolate
int thresh = cfg[0];
int2 coord = (int2)(get_global_id(0),get_global_id(1));
uint4 pix = read_imageui(imgSrc, smp, coord);
int pixBW = (int)pix.x+(int)pix.y+(int)pix.z;
if (pixBW <= 3*thresh) pix = (uint4)(0,0,0,255);
else pix = read_imageui(imgRender, smp, coord);
write_imageui(dst,coord,pix);
}
";
#endregion
#region Propagate distance to line
string srcPropag = @"
__kernel void initWeight(__global float * weight)
{
int2 coord = (int2)(get_global_id(0),get_global_id(1));
int w = get_global_size(0);
int idx = coord.x+w*coord.y;
weight[idx] = 1e20f;
}
__kernel void initTotalWeight(__global float * weight)
{
int2 coord = (int2)(get_global_id(0),get_global_id(1));
int w = get_global_size(0);
int idx = coord.x+w*coord.y;
weight[idx] = 0.0f;
}
__kernel void AddToTotalWeight(__global float * totalWeight,
__global const float * weight,
__constant float * color,
__read_only image2d_t curImg,
__write_only image2d_t dstImg,
__global const uchar* byteInfo)
{
const sampler_t smp = CLK_NORMALIZED_COORDS_FALSE | //Natural coordinates
CLK_ADDRESS_CLAMP | //Clamp to zeros
CLK_FILTER_NEAREST; //Don't interpolate
int x = get_global_id(0);
int y = get_global_id(1);
int w = get_global_size(0);
int idx = x+w*y;
float totWeight = totalWeight[idx];
float myWeight = native_recip(weight[idx]);
myWeight = powr(myWeight, 1.6f);
//myWeight = native_log(1.0f + myWeight);
int2 coord = (int2)(x,y);
uint4 pix = read_imageui(curImg, smp, coord);
float4 curColor = (float4)((float)pix.x,(float)pix.y,(float)pix.z,255.0f);
float4 newColor = (float4)((float)color[0],(float)color[1],(float)color[2],255.0f);
newColor = (newColor * myWeight + curColor * totWeight)*native_recip(myWeight+totWeight);
newColor = clamp(newColor, 0.0f, 255.0f);
pix = (uint4)((uint)newColor.x,(uint)newColor.y,(uint)newColor.z,255);
if (byteInfo[idx] == 0) pix = (uint4)(0,0,0,255);
write_imageui(dstImg,coord,pix);
totalWeight[idx] = totWeight + myWeight;
}
__kernel void PropagateDist(__global int * changed,
__read_only image2d_t imgStroke,
__global const uchar* byteInfo,
__global float * weight,
__write_only image2d_t imgDists)
{
const sampler_t smp = CLK_NORMALIZED_COORDS_FALSE | //Natural coordinates
CLK_ADDRESS_CLAMP | //Clamp to zeros
CLK_FILTER_NEAREST; //Don't interpolate
int2 coord = (int2)(get_global_id(0),get_global_id(1));
int w = get_global_size(0);
uint4 pix = read_imageui(imgStroke, smp, coord);
int val = max((int)pix.z,max((int)pix.x,(int)pix.y));
int idx = coord.x+w*coord.y;
float curW = 0.0f;
if (val > 0) curW == 1E-9f;
else if (coord.x == 0 || coord.y == 0 || coord.x == w-1 || coord.y == get_global_size(1)-1 || byteInfo[idx]==0) curW = 1e20f;
else
{
curW = weight[idx-1]+1;
curW = fmin(curW, weight[idx+1] + 1);
curW = fmin(curW, weight[idx+w] + 1);
curW = fmin(curW, weight[idx-w] + 1);
curW = fmin(curW, weight[idx-w-1] + 1.41421356237f);
curW = fmin(curW, weight[idx-w+1] + 1.41421356237f);
curW = fmin(curW, weight[idx+w-1] + 1.41421356237f);
curW = fmin(curW, weight[idx+w+1] + 1.41421356237f);
}
if (weight[idx] != curW) changed[0] = 1;
weight[idx] = curW;
//float pixBW = clamp(curW,0.0f,255.0f);
//uint4 pix2 = (uint4)((uint)pixBW,(uint)0,255-(uint)pixBW,255);
//if (byteInfo[idx]==0) pix2 = (uint4)((uint)0,(uint)0,(uint)0,255);
//write_imageui(imgDists,coord,pix2);
}
";
#endregion
#endregion
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.Unknown)
{
CLCalc.InitCL();
}
CLCalc.Program.Compile(new string[] { srcThresh, srcPropag });
kernelThreshold = new CLCalc.Program.Kernel("Threshold");
kernelPropagateDist = new CLCalc.Program.Kernel("PropagateDist");
kernelinitWeight = new CLCalc.Program.Kernel("initWeight");
kernelinitTotalWeight = new CLCalc.Program.Kernel("initTotalWeight");
kernelAddToTotalWeight = new CLCalc.Program.Kernel("AddToTotalWeight");
kernelRestoreBlackPixels = new CLCalc.Program.Kernel("RestoreBlackPixels");
}
/// <summary>Static Constructor. Builds kernels</summary>
static SparseLinalg()
{
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.Unknown)
{
try { CLCalc.InitCL(); }
catch
{
}
}
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.UsingCL)
{
//Kernel
CLLinalgSrc src = new CLLinalgSrc();
CLCalc.Program.Compile(new string[] { src.srcDotProd, src.srcMatVecMult, src.srcLinConjGrad });
kernelDotProduct = new CLCalc.Program.Kernel("dotProd");
kernelSum = new CLCalc.Program.Kernel("sumDotProd");
kernelGetDotSum = new CLCalc.Program.Kernel("GetResp");
kernelSparseMatrixVecMult = new CLCalc.Program.Kernel("SparseMatrixVecMult");
//Linear solving
kernelInitRP = new CLCalc.Program.Kernel("InitRP");
kernelMultiplyAdd = new CLCalc.Program.Kernel("MultiplyAdd");
kernelCopyToTemp = new CLCalc.Program.Kernel("CopyToTemp");
}
}
public void Initialize(BaseCameraApplication app, OpenTKWrapper.CLGLInterop.GLAdvancedRender glw)
{
try
{
CLCalc.Program.Compile(app.GetPrimaryDevice().GetPreprocessCode() + src);
}
catch (BuildProgramFailureComputeException ex)
{
System.Console.WriteLine(ex.Message);
Environment.Exit(1);
}
kernelCopyImage = new CLCalc.Program.Kernel("CopyImageToMesh");
BoundingBox bbox = app.GetPrimaryDevice().GetBoundingBox();
int w = app.GetPrimaryDevice().GetDepthImage().Width;
int h = app.GetPrimaryDevice().GetDepthImage().Height;
int size = w * h;
ColorData = new float[16 * size];
PositionData = new float[16 * size];
NormalData = new float[12 * size];
for (int i = 0; i < size; i++)
{
PositionData[4 * i] = (i / w) - w / 2;
PositionData[4 * i + 2] = i % w - h / 2;
PositionData[4 * i + 1] = i % 7;
PositionData[4 * i + 3] = 1.0f;
}
GL.GenBuffers(3, QuadMeshBufs);
GL.BindBuffer(BufferTarget.ArrayBuffer, QuadMeshBufs[0]);
GL.BufferData(BufferTarget.ArrayBuffer, (IntPtr)(ColorData.Length * sizeof(float)), ColorData, BufferUsageHint.StreamDraw);
GL.BindBuffer(BufferTarget.ArrayBuffer, QuadMeshBufs[1]);
GL.BufferData(BufferTarget.ArrayBuffer, (IntPtr)(PositionData.Length * sizeof(float)), PositionData, BufferUsageHint.StreamDraw);//Notice STREAM DRAW
GL.BindBuffer(BufferTarget.ArrayBuffer, QuadMeshBufs[2]);
GL.BufferData(BufferTarget.ArrayBuffer, (IntPtr)(NormalData.Length * sizeof(float)), NormalData, BufferUsageHint.StreamDraw);//Notice STREAM DRAW
colorBuffer = new CLCalc.Program.Variable(QuadMeshBufs[0], typeof(float));
positionBuffer = new CLCalc.Program.Variable(QuadMeshBufs[1], typeof(float));
normalBuffer = new CLCalc.Program.Variable(QuadMeshBufs[2], typeof(float));
GL.BindBuffer(BufferTarget.ArrayBuffer, 0);
GL.Enable(EnableCap.Blend);
}
/// <summary>Constructor.</summary>
/// <param name="InitialState">Initial state of system</param>
/// <param name="StepSize">Desired step per integration pass</param>
/// <param name="InitialIndepVarValue">Initial independent variable value</param>
/// <param name="DerivativeCalculator">Function to calculate derivatives vector</param>
public doubleODE46(double InitialIndepVarValue, double StepSize, double[] InitialState, DerivCalcDeleg DerivativeCalculator)
{
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.Unknown)
{
CLCalc.InitCL();
}
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.NotUsingCL)
throw new Exception("OpenCL not available");
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.UsingCL)
{
ODE46Source Source = new ODE46Source();
string[] s = new string[] { @"
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
", Source.doubleStep2, Source.doubleStep3, Source.doubleStep4, Source.doubleStep5, Source.doubleStep6, Source.doubleFinalizeCalc };
CLCalc.Program.Compile(s);
//Calculador de derivada
Derivs = DerivativeCalculator;
//Scalars
double[] xx = new double[1] { InitialIndepVarValue };
x = new CLCalc.Program.Variable(xx);
xsav = new CLCalc.Program.Variable(xx);
//Sets initial values to Device and local variables
hdid = new CLCalc.Program.Variable(xx);
currentX = InitialIndepVarValue;
SetStep(StepSize);
//Vectors
yy = new double[InitialState.Length];
for (int i = 0; i < InitialState.Length; i++) yy[i] = InitialState[i];
ysav = new CLCalc.Program.Variable(yy);
k1 = new CLCalc.Program.Variable(InitialState);
k2 = new CLCalc.Program.Variable(InitialState);
k3 = new CLCalc.Program.Variable(InitialState);
k4 = new CLCalc.Program.Variable(InitialState);
k5 = new CLCalc.Program.Variable(InitialState);
k6 = new CLCalc.Program.Variable(InitialState);
absError = new CLCalc.Program.Variable(new double[InitialState.Length]);
y = new CLCalc.Program.Variable(yy);
//Kernels
KernelFinalizeCalc = new CLCalc.Program.Kernel("doubleFinalizeCalc");
KernelUpdateX = new CLCalc.Program.Kernel("doubleUpdateX");
KernelRK46YStep2 = new CLCalc.Program.Kernel("doubleYStep2");
KernelRK46XStep2 = new CLCalc.Program.Kernel("doubleXStep2");
KernelRK46YStep3 = new CLCalc.Program.Kernel("doubleYStep3");
KernelRK46XStep3 = new CLCalc.Program.Kernel("doubleXStep3");
KernelRK46YStep4 = new CLCalc.Program.Kernel("doubleYStep4");
KernelRK46XStep4 = new CLCalc.Program.Kernel("doubleXStep4");
KernelRK46YStep5 = new CLCalc.Program.Kernel("doubleYStep5");
KernelRK46XStep5 = new CLCalc.Program.Kernel("doubleXStep5");
KernelRK46YStep6 = new CLCalc.Program.Kernel("doubleYStep6");
KernelRK46XStep6 = new CLCalc.Program.Kernel("doubleXStep6");
//Kernel arguments
ArgsFinalize = new CLCalc.Program.Variable[] { x, hdid, y, ysav, absError, k1, k2, k3, k4, k5, k6 };
ArgsRK46Y = new CLCalc.Program.Variable[] { x, hdid, y, ysav, k1, k2, k3, k4, k5, k6 };
ArgsRK46X = new CLCalc.Program.Variable[] { x, hdid, xsav };
NStates = new int[1] { InitialState.Length };
NScalar = new int[1] { 1 };
//Data retrieving
yerr = new double[NStates[0]];
xRet = new double[NScalar[0]];
}
}
/// <summary>Creates a new isosurface calculator. You may pass variables created from a OpenGL context to the CL variables if you are using interop or NULL
/// if not using OpenCL/GL interop.</summary>
/// <param name="FuncValues">Values of the evaluated 3D function f(x,y,z). FuncValues=float[maxX,maxY,maxZ]</param>
/// <param name="CLEdgeCoords">OpenCL variable (float) to hold edge coordinates. Dimension has to be 9 * maxX * maxY * maxZ</param>
/// <param name="CLEdgeNormals">OpenCL variable (float) to hold edge normals. Dimension has to be 9 * maxX * maxY * maxZ</param>
/// <param name="CLElementArrayIndex">OpenCL variable (int) to hold element array index. Dimension has to be 5 * 3 * (maxX - 1) * (maxY - 1) * (maxZ - 1)</param>
private void InitMarchingCubes(float[, ,] FuncValues, CLCalc.Program.Variable CLEdgeCoords, CLCalc.Program.Variable CLEdgeNormals, CLCalc.Program.Variable CLElementArrayIndex)
{
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.Unknown)
{
CLCalc.InitCL();
}
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.UsingCL)
{
//Reads maximum lengths
int maxX = FuncValues.GetLength(0);
int maxY = FuncValues.GetLength(1);
int maxZ = FuncValues.GetLength(2);
max = new int[] { maxX, maxY, maxZ };
#region Creating variables
//Isolevel
isoLevel = new float[1] {
1.32746E-5f
};
varIsoLevel = new CLCalc.Program.Variable(isoLevel);
//Step size and x0,y0,z0
varStep = new CLCalc.Program.Variable(step);
varInitVals = new CLCalc.Program.Variable(initVals);
//Create and copy function values
funcVals = new float[maxX * maxY * maxZ];
CLFuncVals = new CLCalc.Program.Variable(funcVals);
SetFuncVals(FuncValues);
//Edge coordinates - 3 coords * 3 possible directions * number of points
edgeCoords = new float[9 * maxX * maxY * maxZ];
if (CLEdgeCoords != null)
{
varEdgeCoords = CLEdgeCoords;
varEdgeCoords.WriteToDevice(edgeCoords);
}
else
{
varEdgeCoords = new CLCalc.Program.Variable(edgeCoords);
}
//4 preliminary normals per edge - has to be averaged afterwards
edgePrelimNormals = new float[36 * maxX * maxY * maxZ];
varEdgePrelimNormals = new CLCalc.Program.Variable(edgePrelimNormals);
//Edge normals
edgeNormals = new float[9 * maxX * maxY * maxZ];
if (CLEdgeNormals != null)
{
varEdgeNormals = CLEdgeNormals;
varEdgeNormals.WriteToDevice(edgeNormals);
}
else
{
varEdgeNormals = new CLCalc.Program.Variable(edgeNormals);
}
//Number of cubes: (maxX-1)*(maxY-1)*(maxZ-1)
//Marching cube algorithm: each cube can have 5 triangles drawn, 3 vertexes per triangle
//q-th vertex of p-th triangle of the ijk-th cube: [(5*(i+(maxX-1)*j+k*(maxX-1)*(maxY-1))+p)*3+q]
elementIndex = new int[5 * 3 * (maxX - 1) * (maxY - 1) * (maxZ - 1)];
if (CLElementArrayIndex != null)
{
varElemIndex = CLElementArrayIndex;
varElemIndex.WriteToDevice(elementIndex);
}
else
{
varElemIndex = new CLCalc.Program.Variable(elementIndex);
}
//Edge remapping to build output
edges = new int[edgeCoords.Length / 3];
for (int i = 0; i < edges.Length; i++)
{
edges[i] = -1;
}
#endregion
#region Compile code and create kernels
CLMarchingCubesSrc cmsrc = new CLMarchingCubesSrc();
CLCalc.Program.Compile(new string[] { cmsrc.definitions, cmsrc.src });
kernelInterpPts = new CLCalc.Program.Kernel("interpPts");
kernelPolygonize = new CLCalc.Program.Kernel("Polygonize");
kernelSmoothNormals = new CLCalc.Program.Kernel("SmoothNormals");
kernelPolygonizeNoNormals = new CLCalc.Program.Kernel("PolygonizeNoNormals");
#endregion
}
else
{
throw new Exception("OpenCL not available");
}
}
public void Initialize(BaseCameraApplication capture)
{
DepthCameraFrame frame = capture.GetDevices()[0].GetDepthImage();
try
{
StreamReader reader = new StreamReader(new MemoryStream(Perceptual.Foundation.Properties.Resources.AdaptiveTemporalFilter));
string text = reader.ReadToEnd();
CLCalc.Program.Compile(capture.GetPrimaryDevice().GetPreprocessCode() + "\n#define HISTORY_SIZE " + historySize + "\n" + text);
reader.Close();
}
catch (Exception ex)
{
System.Console.WriteLine(ex.Message);
System.Console.WriteLine("Could not find AdaptiveTemporalFilter.cl");
Environment.Exit(1);
}
historyIndex = new CLCalc.Program.Value<int>(historySize - 1);
updateBuffer = new CLCalc.Program.Kernel("UpdateFilter");
copyToTemporalBuffer = new CLCalc.Program.Kernel("CopyToTemporalBuffer");
depthBuffer = CLCalc.Program.Variable.Create(new ComputeBuffer<float>(CLCalc.Program.Context, ComputeMemoryFlags.ReadWrite, 4 * frame.Width * frame.Height));
depthCopyBuffer = new CLCalc.Program.Variable(new float[4 * frame.Width * frame.Height]);
depthTemporalBuffer = CLCalc.Program.Variable.Create(new ComputeBuffer<float>(CLCalc.Program.Context, ComputeMemoryFlags.ReadWrite, historySize * 4 * frame.Width * frame.Height));
depthImage = new CLCalc.Program.Image2D(new float[frame.Height * frame.Width * 4], frame.Width, frame.Height);
uvImage = new CLCalc.Program.Image2D(new float[frame.Height * frame.Width * 4], frame.Width, frame.Height);
kernelErodeFilter = new CLCalc.Program.Kernel("ErodeFilter");
kernelDilateFilter = new CLCalc.Program.Kernel("DilateFilter");
kernelCopyImage = new CLCalc.Program.Kernel("CopyImage");
kernelCopyBuffer = new CLCalc.Program.Kernel("CopyDepth");
kernelMedianFilter1 = new CLCalc.Program.Kernel("SmallFilter");
kernelMedianFilter2 = new CLCalc.Program.Kernel("LargeFilter");
}
/// <summary>Initializes CL kernels</summary>
public static void Init()
{
if (kernelCholeskyDiagBlock == null)
{
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.Unknown)
{
CLCalc.InitCL();
}
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.UsingCL)
{
if (kernelCholeskyDiagBlock == null)
{
SUBMATRIXSIZE = (int)Math.Sqrt((double)CLCalc.Program.CommQueues[CLCalc.Program.DefaultCQ].Device.MaxWorkGroupSize);
SUBMATRIXSIZE = Math.Min(16, SUBMATRIXSIZE);
string strSubSize = SUBMATRIXSIZE.ToString();
string strTotSize = (SUBMATRIXSIZE * (SUBMATRIXSIZE + 1) / 2).ToString();
LinalgSrc src = new LinalgSrc();
string srcBlockChol = src.srcBlockCholesky.Replace("CONSTSUBMATRIXSIZE", strSubSize).Replace("CONSTGLOBALSIZE", strTotSize);
CLCalc.Program.Compile(new string[] { srcBlockChol, src.srcBkSubs, src.srcOperations, src.srcVecSum,
src.srcpNorm, src.strFeasibFunc, src.srcLogistReg });
kernelCholeskyDiagBlock = new CLCalc.Program.Kernel("CholeskyDiagBlock");
kernelCholeskyComputePanel = new CLCalc.Program.Kernel("CholeskyComputePanel");
kernelCholeskyForwardProp = new CLCalc.Program.Kernel("CholeskyForwardProp");
kernelFwdUpperBackSubs = new CLCalc.Program.Kernel("FwdUpperBackSubs");
kernelBkLowerBackSubs = new CLCalc.Program.Kernel("BkLowerBackSubs");
kernelFwdPropag = new CLCalc.Program.Kernel("FwdPropag");
kernelFwdPropag2 = new CLCalc.Program.Kernel("FwdPropag2");
kernelBackPropag = new CLCalc.Program.Kernel("BackPropag");
kernelBackPropag2 = new CLCalc.Program.Kernel("BackPropag2");
kernelInPlaceSubtract = new CLCalc.Program.Kernel("InPlaceSubtract");
kernelElemWiseAbs = new CLCalc.Program.Kernel("ElemWiseAbs");
kernelInnerProd = new CLCalc.Program.Kernel("InnerProd");
//Linear algebra
kernelSymMatrVecMultiply = new CLCalc.Program.Kernel("SymMatrVecMultiply");
kernelSymMatrMatrMultiply = new CLCalc.Program.Kernel("SymMatrMatrMultiply");
kernelComputeAtWA = new CLCalc.Program.Kernel("ComputeAtWA");
kernelComputeAinvHAt = new CLCalc.Program.Kernel("ComputeAinvHAt");
kernelRegularMatrTranspMatrProd = new CLCalc.Program.Kernel("RegularMatrTranspMatrProd");
kernelRegularMatrMatrProd = new CLCalc.Program.Kernel("RegularMatrMatrProd");
kernelCopyBuffer = new CLCalc.Program.Kernel("CopyBuffer");
kernelLinearComb = new CLCalc.Program.Kernel("LinearComb");
kernelMatrVecProd = new CLCalc.Program.Kernel("MatrVecProd");
kernelTranspMatrVecProdW = new CLCalc.Program.Kernel("TranspMatrVecProdW");
kernelMatrVecProdSumVec = new CLCalc.Program.Kernel("MatrVecProdSumVec");
kernelDiagVecProd = new CLCalc.Program.Kernel("DiagVecProd");
kernelDiagTranspMatProd = new CLCalc.Program.Kernel("DiagTranspMatProd");
kernelElemWiseProd = new CLCalc.Program.Kernel("ElemWiseProd");
kernelElemWiseInv = new CLCalc.Program.Kernel("ElemWiseInv");
kernelElemWiseInv2 = new CLCalc.Program.Kernel("ElemWiseInv2");
kernelClear = new CLCalc.Program.Kernel("ClearResps");
kernelPreSum = new CLCalc.Program.Kernel("PreSum");
kernelCoalLocalSum = new CLCalc.Program.Kernel("CoalLocalSum");
kernelHasPositiveEntry = new CLCalc.Program.Kernel("HasPositiveEntry");
//pNorm minimization
floatOptimization.CurveFitting.kernelpNorm = new CLCalc.Program.Kernel("pNorm");
floatOptimization.CurveFitting.kerneldpNorm = new CLCalc.Program.Kernel("dpNorm");
//Logistic regression
floatOptimization.LogisticRegression.kernelComputeLogistRegParams = new CLCalc.Program.Kernel("ComputeLogistRegParams");
floatOptimization.LogisticRegression.kernelpNorm = floatOptimization.CurveFitting.kernelpNorm;
floatOptimization.LogisticRegression.kerneldpNorm = floatOptimization.CurveFitting.kerneldpNorm;
//Feasibility
floatOptimization.QuadraticProgramming.kernelgetLast = new CLCalc.Program.Kernel("getLast");
}
}
}
}
/// <summary>Constructor.</summary>
/// <param name="InitialState">Initial state of system</param>
/// <param name="StepSize">Desired step per integration pass</param>
/// <param name="InitialIndepVarValue">Initial independent variable value</param>
/// <param name="DerivativeCalculator">Function to calculate derivatives vector</param>
public floatODE46(float InitialIndepVarValue, float StepSize, float[] InitialState, DerivCalcDeleg DerivativeCalculator)
{
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.Unknown)
{
CLCalc.InitCL();
}
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.NotUsingCL)
{
throw new Exception("OpenCL not available");
}
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.UsingCL)
{
ODE46Source Source = new ODE46Source();
string[] s = new string[] { Source.floatStep2, Source.floatStep3, Source.floatStep4, Source.floatStep5, Source.floatStep6, Source.floatFinalizeCalc };
CLCalc.Program.Compile(s);
//Calculador de derivada
Derivs = DerivativeCalculator;
//Scalars
float[] xx = new float[1] {
InitialIndepVarValue
};
x = new CLCalc.Program.Variable(xx);
xsav = new CLCalc.Program.Variable(xx);
//Sets initial values to Device and local variables
hdid = new CLCalc.Program.Variable(xx);
currentX = InitialIndepVarValue;
SetStep(StepSize);
//Vectors
yy = new float[InitialState.Length];
for (int i = 0; i < InitialState.Length; i++)
{
yy[i] = InitialState[i];
}
ysav = new CLCalc.Program.Variable(yy);
k1 = new CLCalc.Program.Variable(InitialState);
k2 = new CLCalc.Program.Variable(InitialState);
k3 = new CLCalc.Program.Variable(InitialState);
k4 = new CLCalc.Program.Variable(InitialState);
k5 = new CLCalc.Program.Variable(InitialState);
k6 = new CLCalc.Program.Variable(InitialState);
absError = new CLCalc.Program.Variable(new float[InitialState.Length]);
y = new CLCalc.Program.Variable(yy);
//Kernels
KernelFinalizeCalc = new CLCalc.Program.Kernel("floatFinalizeCalc");
KernelUpdateX = new CLCalc.Program.Kernel("floatUpdateX");
KernelRK46YStep2 = new CLCalc.Program.Kernel("floatYStep2");
KernelRK46XStep2 = new CLCalc.Program.Kernel("floatXStep2");
KernelRK46YStep3 = new CLCalc.Program.Kernel("floatYStep3");
KernelRK46XStep3 = new CLCalc.Program.Kernel("floatXStep3");
KernelRK46YStep4 = new CLCalc.Program.Kernel("floatYStep4");
KernelRK46XStep4 = new CLCalc.Program.Kernel("floatXStep4");
KernelRK46YStep5 = new CLCalc.Program.Kernel("floatYStep5");
KernelRK46XStep5 = new CLCalc.Program.Kernel("floatXStep5");
KernelRK46YStep6 = new CLCalc.Program.Kernel("floatYStep6");
KernelRK46XStep6 = new CLCalc.Program.Kernel("floatXStep6");
//Kernel arguments
ArgsFinalize = new CLCalc.Program.Variable[] { x, hdid, y, ysav, absError, k1, k2, k3, k4, k5, k6 };
ArgsRK46Y = new CLCalc.Program.Variable[] { x, hdid, y, ysav, k1, k2, k3, k4, k5, k6 };
ArgsRK46X = new CLCalc.Program.Variable[] { x, hdid, xsav };
NStates = new int[1] {
InitialState.Length
};
NScalar = new int[1] {
1
};
//Data retrieving
yerr = new float[NStates[0]];
xRet = new float[NScalar[0]];
}
}
/// <summary>Creates a new isosurface calculator. You may pass variables created from a OpenGL context to the CL variables if you are using interop or NULL
/// if not using OpenCL/GL interop.</summary>
/// <param name="FuncValues">Values of the evaluated 3D function f(x,y,z). FuncValues=float[maxX,maxY,maxZ]</param>
/// <param name="CLEdgeCoords">OpenCL variable (float) to hold edge coordinates. Dimension has to be 9 * maxX * maxY * maxZ</param>
/// <param name="CLEdgeNormals">OpenCL variable (float) to hold edge normals. Dimension has to be 9 * maxX * maxY * maxZ</param>
/// <param name="CLElementArrayIndex">OpenCL variable (int) to hold element array index. Dimension has to be 5 * 3 * (maxX - 1) * (maxY - 1) * (maxZ - 1)</param>
private void InitMarchingCubes(float[, ,] FuncValues, CLCalc.Program.Variable CLEdgeCoords, CLCalc.Program.Variable CLEdgeNormals, CLCalc.Program.Variable CLElementArrayIndex)
{
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.Unknown) CLCalc.InitCL();
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.UsingCL)
{
//Reads maximum lengths
int maxX = FuncValues.GetLength(0);
int maxY = FuncValues.GetLength(1);
int maxZ = FuncValues.GetLength(2);
max = new int[] { maxX, maxY, maxZ };
#region Creating variables
//Isolevel
isoLevel = new float[1] { 1.32746E-5f };
varIsoLevel = new CLCalc.Program.Variable(isoLevel);
//Step size and x0,y0,z0
varStep = new CLCalc.Program.Variable(step);
varInitVals = new CLCalc.Program.Variable(initVals);
//Create and copy function values
funcVals = new float[maxX * maxY * maxZ];
CLFuncVals = new CLCalc.Program.Variable(funcVals);
SetFuncVals(FuncValues);
//Edge coordinates - 3 coords * 3 possible directions * number of points
edgeCoords = new float[9 * maxX * maxY * maxZ];
if (CLEdgeCoords != null)
{
varEdgeCoords = CLEdgeCoords;
varEdgeCoords.WriteToDevice(edgeCoords);
}
else varEdgeCoords = new CLCalc.Program.Variable(edgeCoords);
//4 preliminary normals per edge - has to be averaged afterwards
edgePrelimNormals = new float[36 * maxX * maxY * maxZ];
varEdgePrelimNormals = new CLCalc.Program.Variable(edgePrelimNormals);
//Edge normals
edgeNormals = new float[9 * maxX * maxY * maxZ];
if (CLEdgeNormals != null)
{
varEdgeNormals = CLEdgeNormals;
varEdgeNormals.WriteToDevice(edgeNormals);
}
else varEdgeNormals = new CLCalc.Program.Variable(edgeNormals);
//Number of cubes: (maxX-1)*(maxY-1)*(maxZ-1)
//Marching cube algorithm: each cube can have 5 triangles drawn, 3 vertexes per triangle
//q-th vertex of p-th triangle of the ijk-th cube: [(5*(i+(maxX-1)*j+k*(maxX-1)*(maxY-1))+p)*3+q]
elementIndex = new int[5 * 3 * (maxX - 1) * (maxY - 1) * (maxZ - 1)];
if (CLElementArrayIndex != null)
{
varElemIndex = CLElementArrayIndex;
varElemIndex.WriteToDevice(elementIndex);
}
else varElemIndex = new CLCalc.Program.Variable(elementIndex);
//Edge remapping to build output
edges = new int[edgeCoords.Length / 3];
for (int i = 0; i < edges.Length; i++) edges[i] = -1;
#endregion
#region Compile code and create kernels
CLMarchingCubesSrc cmsrc = new CLMarchingCubesSrc();
CLCalc.Program.Compile(new string[] { cmsrc.definitions, cmsrc.src });
kernelInterpPts = new CLCalc.Program.Kernel("interpPts");
kernelPolygonize = new CLCalc.Program.Kernel("Polygonize");
kernelSmoothNormals = new CLCalc.Program.Kernel("SmoothNormals");
kernelPolygonizeNoNormals = new CLCalc.Program.Kernel("PolygonizeNoNormals");
#endregion
}
else throw new Exception("OpenCL not available");
}
public void Initialize(BaseCameraApplication capture)
{
DepthCameraFrame depthImage = capture.GetPrimaryDevice().GetDepthImage();
this.width = depthImage.Width;
this.height = depthImage.Height;
this.filter = ((AdaptiveTemporalFilter)capture.GetImageFilter());
CvSize sz = new CvSize(depthImage.Width, depthImage.Height);
gray = new IplImage(sz, BitDepth.U8, 1);
erode = new IplImage(sz, BitDepth.U8, 1);
dilate = new IplImage(sz, BitDepth.U8, 1);
tmp = new IplImage(sz, BitDepth.U8, 1);
mask = new IplImage(sz, BitDepth.U8, 1);
imgLabel = new IplImage(sz, BitDepth.F32, 1);
faceDetectionBuffer = CLCalc.Program.Variable.Create(new ComputeBuffer<FaceLandmarks>(CLCalc.Program.Context, ComputeMemoryFlags.ReadWrite, 1));
try
{
CLCalc.Program.Compile(capture.GetPrimaryDevice().GetPreprocessCode() + src);
}
catch (BuildProgramFailureComputeException ex)
{
System.Console.WriteLine(ex.Message);
Environment.Exit(1);
}
irImageBuffer = CLCalc.Program.Variable.Create(new ComputeBuffer<byte>(CLCalc.Program.Context, ComputeMemoryFlags.ReadWrite | ComputeMemoryFlags.CopyHostPointer, ir = new byte[width * height]));
kernelCopyIRImage = new CLCalc.Program.Kernel("CopyIRImage");
kernelFindFaceLandmarks = new CLCalc.Program.Kernel("FindFaceLandmarks");
}
/// <summary>Compiles code and initializes kernel for this svm stance</summary>
private void CLSVMInit()
{
if (CLResource == null) CLResource = new int[0];
lock (CLResource)
{
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.Unknown) CLCalc.InitCL();
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.UsingCL)
{
if (kernelComputeKernelRBF == null)
{
CLSVMSrc s = new CLSVMSrc();
CLCalc.Program.Compile(new string[] { s.srcKernels, s.srcFindMaxMinErr, s.srcMultClass });
//Kernel computation
kernelComputeKernelRBF = new CLCalc.Program.Kernel("ComputeKernelRBF");
kernelMaxErr = new CLCalc.Program.Kernel("maxErr");
kernelComputeMax = new CLCalc.Program.Kernel("computeMax");
kernelMinErr = new CLCalc.Program.Kernel("minErr");
kernelComputeMin = new CLCalc.Program.Kernel("computeMin");
kernelGetResp = new CLCalc.Program.Kernel("getResp");
//Update error
kernelUpdateErr = new CLCalc.Program.Kernel("UpdateErr");
//Multiple classification
kernelComputeMultiKernelRBF = new CLCalc.Program.Kernel("ComputeMultiKernelRBF");
kernelSumKernels=new CLCalc.Program.Kernel("SumKernels");
}
//Memory obbjects
//Find max/min
CLErrLen = new CLCalc.Program.Variable(new int[1]);
HostResp = new int[1];
CLResp = new CLCalc.Program.Variable(HostResp);
CLMaxMinErrs = new CLCalc.Program.Variable(new float[MAXMINWORKSIZE]);
CLMaxMinInds = new CLCalc.Program.Variable(new int[MAXMINWORKSIZE]);
//Update error
CLUpdtErrParams = new CLCalc.Program.Variable(new float[3]);
}
}
}
private static void InitKernels()
{
string s = new CLFFTSrc().s;
CLCalc.InitCL();
try
{
CLCalc.Program.Compile(s);
}
catch
{
}
kernelfft_radix16 = new CLCalc.Program.Kernel("fft_radix16");
kernelfft_radix4 = new CLCalc.Program.Kernel("fft_radix4");
kernelConjugate = new CLCalc.Program.Kernel("Conjugate");
CLp = new CLCalc.Program.Variable(new int[1]);
}
private void InitKernel()
{
if (kernelExtractFeatures == null)
{
CLExtractFeatSrc src = new CLExtractFeatSrc();
CLBracketRegionsSrc src2 = new CLBracketRegionsSrc();
CLCalc.Program.Compile(new string[] { src2.src, src.src });
kernelExtractFeatures = new CLCalc.Program.Kernel("ExtractFeatures");
kernelComputeFrameDiff = new CLCalc.Program.Kernel("ComputeFrameDiff");
kernelSegregateSkin = new CLCalc.Program.Kernel("SegregateSkin");
}
}
public void Initialize(BaseCameraApplication capture, GLAdvancedRender glw)
{
try
{
CLCalc.Program.Compile(capture.GetPrimaryDevice().GetPreprocessCode() + src);
}
catch (BuildProgramFailureComputeException ex)
{
System.Console.WriteLine(ex.Message);
Environment.Exit(1);
}
DepthCameraFrame frame = capture.GetDevices()[0].GetDepthImage();
kernelCopyBmp = new CLCalc.Program.Kernel("CopyImageToPointCloud");
int size = frame.Width * frame.Height;
bufs = new int[4];
ColorData = new float[4 * size];
PositionData = new float[4 * size];
GL.GenBuffers(2, bufs);
GL.BindBuffer(BufferTarget.ArrayBuffer, bufs[0]);
GL.BufferData(BufferTarget.ArrayBuffer, (IntPtr)(ColorData.Length * sizeof(float)), ColorData, BufferUsageHint.StreamDraw);
GL.BindBuffer(BufferTarget.ArrayBuffer, bufs[1]);
GL.BufferData(BufferTarget.ArrayBuffer, (IntPtr)(PositionData.Length * sizeof(float)), PositionData, BufferUsageHint.StreamDraw);//Notice STREAM DRAW
GL.Enable(EnableCap.PointSmooth);
GL.PointSize(4.0f);
positions = new CLCalc.Program.Variable(bufs[1], typeof(float));
colors = new CLCalc.Program.Variable(bufs[0], typeof(float));
}
/// <summary>Initializes CL kernels</summary>
public static void Init()
{
if (kernelCholeskyDiagBlock == null)
{
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.Unknown) CLCalc.InitCL();
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.UsingCL)
{
if (kernelCholeskyDiagBlock == null)
{
SUBMATRIXSIZE = (int)Math.Sqrt((double)CLCalc.Program.CommQueues[CLCalc.Program.DefaultCQ].Device.MaxWorkGroupSize);
SUBMATRIXSIZE = Math.Min(16, SUBMATRIXSIZE);
string strSubSize = SUBMATRIXSIZE.ToString();
string strTotSize = (SUBMATRIXSIZE * (SUBMATRIXSIZE + 1) / 2).ToString();
LinalgSrc src = new LinalgSrc();
string srcBlockChol = src.srcBlockCholesky.Replace("CONSTSUBMATRIXSIZE", strSubSize).Replace("CONSTGLOBALSIZE", strTotSize);
CLCalc.Program.Compile(new string[] { srcBlockChol, src.srcBkSubs, src.srcOperations, src.srcVecSum,
src.srcpNorm, src.strFeasibFunc, src.srcLogistReg });
kernelCholeskyDiagBlock = new CLCalc.Program.Kernel("CholeskyDiagBlock");
kernelCholeskyComputePanel = new CLCalc.Program.Kernel("CholeskyComputePanel");
kernelCholeskyForwardProp = new CLCalc.Program.Kernel("CholeskyForwardProp");
kernelFwdUpperBackSubs = new CLCalc.Program.Kernel("FwdUpperBackSubs");
kernelBkLowerBackSubs = new CLCalc.Program.Kernel("BkLowerBackSubs");
kernelFwdPropag = new CLCalc.Program.Kernel("FwdPropag");
kernelFwdPropag2 = new CLCalc.Program.Kernel("FwdPropag2");
kernelBackPropag = new CLCalc.Program.Kernel("BackPropag");
kernelBackPropag2 = new CLCalc.Program.Kernel("BackPropag2");
kernelInPlaceSubtract = new CLCalc.Program.Kernel("InPlaceSubtract");
kernelElemWiseAbs = new CLCalc.Program.Kernel("ElemWiseAbs");
kernelInnerProd = new CLCalc.Program.Kernel("InnerProd");
//Linear algebra
kernelSymMatrVecMultiply = new CLCalc.Program.Kernel("SymMatrVecMultiply");
kernelSymMatrMatrMultiply = new CLCalc.Program.Kernel("SymMatrMatrMultiply");
kernelComputeAtWA = new CLCalc.Program.Kernel("ComputeAtWA");
kernelComputeAinvHAt = new CLCalc.Program.Kernel("ComputeAinvHAt");
kernelRegularMatrTranspMatrProd = new CLCalc.Program.Kernel("RegularMatrTranspMatrProd");
kernelCopyBuffer = new CLCalc.Program.Kernel("CopyBuffer");
kernelLinearComb = new CLCalc.Program.Kernel("LinearComb");
kernelMatrVecProd = new CLCalc.Program.Kernel("MatrVecProd");
kernelTranspMatrVecProdW = new CLCalc.Program.Kernel("TranspMatrVecProdW");
kernelMatrVecProdSumVec = new CLCalc.Program.Kernel("MatrVecProdSumVec");
kernelDiagVecProd = new CLCalc.Program.Kernel("DiagVecProd");
kernelDiagTranspMatProd = new CLCalc.Program.Kernel("DiagTranspMatProd");
kernelElemWiseProd = new CLCalc.Program.Kernel("ElemWiseProd");
kernelElemWiseInv = new CLCalc.Program.Kernel("ElemWiseInv");
kernelElemWiseInv2 = new CLCalc.Program.Kernel("ElemWiseInv2");
kernelClear = new CLCalc.Program.Kernel("ClearResps");
kernelPreSum = new CLCalc.Program.Kernel("PreSum");
kernelCoalLocalSum = new CLCalc.Program.Kernel("CoalLocalSum");
kernelHasPositiveEntry = new CLCalc.Program.Kernel("HasPositiveEntry");
//pNorm minimization
floatOptimization.CurveFitting.kernelpNorm = new CLCalc.Program.Kernel("pNorm");
floatOptimization.CurveFitting.kerneldpNorm = new CLCalc.Program.Kernel("dpNorm");
//Logistic regression
floatOptimization.LogisticRegression.kernelComputeLogistRegParams = new CLCalc.Program.Kernel("ComputeLogistRegParams");
floatOptimization.LogisticRegression.kernelpNorm = floatOptimization.CurveFitting.kernelpNorm;
floatOptimization.LogisticRegression.kerneldpNorm = floatOptimization.CurveFitting.kerneldpNorm;
//Feasibility
floatOptimization.QuadraticProgramming.kernelgetLast = new CLCalc.Program.Kernel("getLast");
}
}
}
}
private void frmCLInfo_Load(object sender, EventArgs e)
{
CLCalc.InitCL(ComputeDeviceTypes.All);
if (CLCalc.CLAcceleration != CLCalc.CLAccelerationType.UsingCL)
{
cmbPlat.Items.Add("OpenCL ERROR");
if (cmbPlat.Items.Count > 0) cmbPlat.SelectedIndex = 0;
}
else
{
foreach(ComputePlatform p in CLCalc.CLPlatforms)
{
cmbPlat.Items.Add(p.Name + " " + p.Profile + " " + p.Vendor + " " + p.Version);
}
if (cmbPlat.Items.Count > 0) cmbPlat.SelectedIndex = 0;
int i=0;
foreach (ComputeDevice d in CLCalc.CLDevices)
{
//if (d.CLDeviceAvailable)
//{
cmbDevices.Items.Add(d.Name + " " + d.Type + " " + d.Vendor + " " + d.Version);
cmbCurDevice.Items.Add(d.Name + " " + d.Type + " " + d.Vendor + " " + d.Version);
//}
//else
//{
// cmbDevices.Items.Add("NOT AVAILABLE: " + d.CLDeviceName + " " + d.CLDeviceType + " " + d.CLDeviceVendor + " " + d.CLDeviceVersion);
// cmbCurDevice.Items.Add("NOT AVAILABLE: " + d.CLDeviceName + " " + d.CLDeviceType + " " + d.CLDeviceVendor + " " + d.CLDeviceVersion);
//}
i++;
}
if (cmbDevices.Items.Count > 0)
{
cmbDevices.SelectedIndex = 0;
cmbCurDevice.SelectedIndex = CLCalc.Program.DefaultCQ;
}
}
ReadImportantRegistryEntries();
//int[] n = new int[3] {1,1,1};
//int[] nn = new int[3];
//CLCalc.Program.Variable v = new CLCalc.Program.Variable(n);
//v.WriteToDevice(n);
//v.ReadFromDeviceTo(nn);
string s = @" kernel void teste() {}";
CLCalc.Program.Compile(s);
try
{
CLCalc.Program.Kernel k = new CLCalc.Program.Kernel("teste");
}
catch
{
MessageBox.Show("");
}
}
public void Initialize()
{
try
{
CLCalc.Program.Compile(src);
}
catch (Exception ex)
{
System.Console.WriteLine(ex.Message);
Environment.Exit(1);
}
kernelUpdateImageGradients = new CLCalc.Program.Kernel("UpdateGradients");
}
private void frmCLInfo_Load(object sender, EventArgs e)
{
CLCalc.InitCL(ComputeDeviceTypes.All);
if (CLCalc.CLAcceleration != CLCalc.CLAccelerationType.UsingCL)
{
cmbPlat.Items.Add("OpenCL ERROR");
if (cmbPlat.Items.Count > 0)
{
cmbPlat.SelectedIndex = 0;
}
}
else
{
foreach (ComputePlatform p in CLCalc.CLPlatforms)
{
cmbPlat.Items.Add(p.Name + " " + p.Profile + " " + p.Vendor + " " + p.Version);
}
if (cmbPlat.Items.Count > 0)
{
cmbPlat.SelectedIndex = 0;
}
int i = 0;
foreach (ComputeDevice d in CLCalc.CLDevices)
{
//if (d.CLDeviceAvailable)
//{
cmbDevices.Items.Add(d.Name + " " + d.Type + " " + d.Vendor + " " + d.Version);
cmbCurDevice.Items.Add(d.Name + " " + d.Type + " " + d.Vendor + " " + d.Version);
//}
//else
//{
// cmbDevices.Items.Add("NOT AVAILABLE: " + d.CLDeviceName + " " + d.CLDeviceType + " " + d.CLDeviceVendor + " " + d.CLDeviceVersion);
// cmbCurDevice.Items.Add("NOT AVAILABLE: " + d.CLDeviceName + " " + d.CLDeviceType + " " + d.CLDeviceVendor + " " + d.CLDeviceVersion);
//}
i++;
}
if (cmbDevices.Items.Count > 0)
{
cmbDevices.SelectedIndex = 0;
cmbCurDevice.SelectedIndex = CLCalc.Program.DefaultCQ;
}
}
ReadImportantRegistryEntries();
//int[] n = new int[3] {1,1,1};
//int[] nn = new int[3];
//CLCalc.Program.Variable v = new CLCalc.Program.Variable(n);
//v.WriteToDevice(n);
//v.ReadFromDeviceTo(nn);
string s = @" kernel void teste() {}";
CLCalc.Program.Compile(s);
try
{
CLCalc.Program.Kernel k = new CLCalc.Program.Kernel("teste");
}
catch
{
MessageBox.Show("");
}
}
static LaserLineTrack()
{
if (CLCalc.CLAcceleration == CLCalc.CLAccelerationType.Unknown)
{
CLCalc.InitCL();
}
#region Source
string src = @"
constant int PIXELSTOSEARCH = PXtoSearch;
const sampler_t smp = CLK_NORMALIZED_COORDS_FALSE | //Natural coordinates
CLK_ADDRESS_CLAMP | //Clamp to zeros
CLK_FILTER_NEAREST; //Don't interpolate
__kernel void initIntM(__read_only image2d_t img,
__global float * intM,
__constant float * threshold)
{
int y = get_global_id(0);
uint4 imgColor = read_imageui(img, smp, (int2)(0,y));
float intens = ((float)imgColor.x + (float)imgColor.y + (float)imgColor.z) * 0.001307189542f; // *1.0f/(3*255)
intM[y] = intens > threshold[0] ? intens : 0.0f;
}
//penalties[0] = distancePenalty, penalties[1] = changePenalty
__kernel void integrateM(__read_only image2d_t img,
__global float * intM,
__constant float * threshold,
__constant int * x,
__constant float * penalties)
{
int y = get_global_id(0);
int H = get_global_size(0);
uint4 imgColor = read_imageui(img, smp, (int2)(x[0],y));
float intens = ((float)imgColor.x + (float)imgColor.y + (float)imgColor.z) * 0.001307189542f; // *1.0f/(3*255)
float maxv = -1e-8f; //intM[H*(x[0] - 1) + y] - penalties[0];
for (int k = -PIXELSTOSEARCH; k <= PIXELSTOSEARCH; k++)
{
if (y + k >= 0 && y + k < H)
maxv = fmax(maxv, intM[H*(x[0] - 1) + y + k] - sqrt(1.0f + k * k) * penalties[0] - (k != 0 ? penalties[1] : 0.0f));
}
intM[y + H*x[0]] = maxv + (intens > threshold[0] ? intens : 0.0f);
}
__kernel void incCounter(__global int * x)
{
x[0]++;
}
__kernel void backTrack( __global float * intM,
__global int * path,
__constant int * Dimensions)
{
int W = Dimensions[0];
int H = Dimensions[1];
int idMax = 0;
float valMax = intM[H*(W - 1)];
//find most amplified path
for (int y = 0; y < H; y++)
{
if (valMax < intM[H*(W - 1) + y])
{
valMax = intM[H*(W - 1) + y];
idMax = y;
}
}
path[W - 1] = idMax;
for (int x = W - 2; x >= 0; x--)
{
int y = path[x + 1];
float maxv = -1e8f;
int idv = -1;
for (int k = -PIXELSTOSEARCH; k <= PIXELSTOSEARCH; k++)
{
if (y + k >= 0 && y + k < H)
{
if (intM[H*x + y + k] > maxv)
{
maxv = intM[H* x + y + k];
idv = y + k;
}
}
}
path[x] = idv;
}
}
";
#endregion
CLCalc.Program.Compile(src.Replace("PXtoSearch", PIXELSTOSEARCH.ToString()));
CLthresh = new CLCalc.Program.Variable(new float[1]);
CLx = new CLCalc.Program.Variable(new int[1]);
CLDim = new CLCalc.Program.Variable(new int[2]);
CLpenalties = new CLCalc.Program.Variable(new float[2]);
kernelinitIntM = new CLCalc.Program.Kernel("initIntM");
kernelintM = new CLCalc.Program.Kernel("integrateM");
kernelBackTrack = new CLCalc.Program.Kernel("backTrack");
kernelincCounter = new CLCalc.Program.Kernel("incCounter");
}
