public static void complexDiv(GThread thread, ComplexD[,] a, ComplexD[,] b, ComplexD[,] c)
{
int x = thread.blockIdx.x;
int y = 0;
while (y < YSIZE)
{
c[x, y] = ComplexD.Divide(a[x, y], b[x, y]);
y++;
}
}
private static ComplexD cth(ComplexD c)
{
ComplexD sn = sh(c);
ComplexD cs = ch(c);
if (sn.x == 0 && sn.y == 0)
{
return(new ComplexD(double.MaxValue, 0));
}
return(ComplexD.Divide(cs, sn));
}
private static ComplexD th(ComplexD c)
{
ComplexD sn = sh(c);
ComplexD cs = ch(c);
if (cs.x == 0 && cs.y == 0)
{
return(new ComplexD(double.MaxValue, 0));
}
return(ComplexD.Divide(sn, cs));
}
private static ComplexD tg(ComplexD c)
{
ComplexD sn = sin(c);
ComplexD cs = cos(c);
if (cs.x == 0 && cs.y == 0)
{
return(new ComplexD(1e37F, 0));
}
return(ComplexD.Divide(sn, cs));
}
public static void Execute()
{
CudafyModule km = CudafyTranslator.Cudafy();
GPGPU gpu = CudafyHost.GetDevice(CudafyModes.Target);
gpu.LoadModule(km);
// 2D
Console.WriteLine("kernel");
ComplexD[,] host_A = new ComplexD[XSIZE, YSIZE];
ComplexD[,] host_B = new ComplexD[XSIZE, YSIZE];
ComplexD[,] host_C = new ComplexD[XSIZE, YSIZE];
int i = 0;
for (int x = 0; x < XSIZE; x++)
{
for (int y = 0; y < YSIZE; y++)
{
host_A[x, y] = new ComplexD(i, i);
host_B[x, y] = new ComplexD(2, 0);
i++;
}
}
ComplexD[,] dev_A = gpu.CopyToDevice(host_A);
ComplexD[,] dev_B = gpu.CopyToDevice(host_B);
ComplexD[,] dev_C = gpu.Allocate <ComplexD>(XSIZE, YSIZE);
Console.WriteLine("complexAdd");
gpu.Launch(XSIZE, 1, "complexAdd", dev_A, dev_B, dev_C);
gpu.CopyFromDevice(dev_C, host_C);
i = 0;
bool pass = true;
for (int x = 0; x < XSIZE; x++)
{
for (int y = 0; y < YSIZE && pass; y++)
{
ComplexD expected = ComplexD.Add(host_A[x, y], host_B[x, y]);
pass = host_C[x, y].x == expected.x && host_C[x, y].y == expected.y;
}
}
Console.WriteLine(pass ? "Pass" : "Fail");
Console.WriteLine("complexSub");
gpu.Launch(XSIZE, 1, "complexSub", dev_A, dev_B, dev_C);
gpu.CopyFromDevice(dev_C, host_C);
i = 0;
pass = true;
for (int x = 0; x < XSIZE; x++)
{
for (int y = 0; y < YSIZE && pass; y++)
{
ComplexD expected = ComplexD.Subtract(host_A[x, y], host_B[x, y]);
pass = host_C[x, y].x == expected.x && host_C[x, y].y == expected.y;
}
}
Console.WriteLine(pass ? "Pass" : "Fail");
Console.WriteLine("complexMpy");
gpu.Launch(XSIZE, 1, "complexMpy", dev_A, dev_B, dev_C);
gpu.CopyFromDevice(dev_C, host_C);
i = 0;
pass = true;
for (int x = 0; x < XSIZE; x++)
{
for (int y = 0; y < YSIZE && pass; y++)
{
ComplexD expected = ComplexD.Multiply(host_A[x, y], host_B[x, y]);
//Console.WriteLine("{0} {1} : {2} {3}", host_C[x, y].R, host_C[x, y].I, expected.R, expected.I);
pass = Verify(host_C[x, y], expected, 1e-14F);
i++;
}
}
Console.WriteLine(pass ? "Pass" : "Fail");
Console.WriteLine("complexDiv");
gpu.Launch(XSIZE, 1, "complexDiv", dev_A, dev_B, dev_C);
gpu.CopyFromDevice(dev_C, host_C);
i = 0;
pass = true;
for (int x = 0; x < XSIZE; x++)
{
for (int y = 0; y < YSIZE && pass; y++)
{
ComplexD expected = ComplexD.Divide(host_A[x, y], host_B[x, y]);
//Console.WriteLine("{0} {1} : {2} {3}", host_C[x, y].R, host_C[x, y].I, expected.R, expected.I);
if (i > 0)
{
pass = Verify(host_C[x, y], expected, 1e-13F);
}
i++;
}
}
Console.WriteLine(pass ? "Pass" : "Fail");
Console.WriteLine("complexAbs");
gpu.Launch(XSIZE, 1, "complexAbs", dev_A, dev_C);
gpu.CopyFromDevice(dev_C, host_C);
i = 0;
pass = true;
for (int x = 0; x < XSIZE; x++)
{
for (int y = 0; y < YSIZE && pass; y++)
{
double expected = ComplexD.Abs(host_A[x, y]);
pass = Verify(host_C[x, y].x, expected, 1e-2F);
//Console.WriteLine("{0} {1} : {2}", host_C[x, y].x, host_C[x, y].y, expected);
i++;
}
}
Console.WriteLine(pass ? "Pass" : "Fail");
gpu.FreeAll();
}
private static ComplexD eval(int rpnLength, byte[] rpnFormula, ComplexD[] rpnKoef, ComplexD c, ComplexD z, ComplexD[] stack, int stackOffset)
{
//return ComplexD.Add( ComplexD.Multiply(ComplexD.Multiply(z,z),z),c);
int sPtr = stackOffset - 1;
int vPtr = 0;
for (int i = 0; i < rpnLength; i++)
{
byte v = rpnFormula[i];
if (v == 0)
{
stack[++sPtr] = rpnKoef[vPtr++];
}
else if (v <= 5)
{
if (v == 1)
{
stack[sPtr - 1] = ComplexD.Add(stack[sPtr - 1], stack[sPtr]);
}
else if (v == 2)
{
stack[sPtr - 1] = ComplexD.Subtract(stack[sPtr - 1], stack[sPtr]);
}
else if (v == 3)
{
stack[sPtr - 1] = ComplexD.Multiply(stack[sPtr - 1], stack[sPtr]);
}
else if (v == 4)
{
stack[sPtr - 1] = ComplexD.Divide(stack[sPtr - 1], stack[sPtr]);
}
else if (v == 5)
{
stack[sPtr - 1] = pow(stack[sPtr - 1], stack[sPtr]);
}
sPtr--;
}
else if (v <= 10)
{
if (v == 6)
{
stack[sPtr] = log(stack[sPtr]);
}
else if (v == 7)
{
stack[sPtr] = exp(stack[sPtr]);
}
else if (v == 8)
{
stack[sPtr] = sin(stack[sPtr]);
}
else if (v == 9)
{
stack[sPtr] = cos(stack[sPtr]);
}
else if (v == 10)
{
stack[sPtr] = tg(stack[sPtr]);
}
}
else if (v <= 19)
{
if (v == 11)
{
stack[sPtr] = ctg(stack[sPtr]);
}
else if (v == 12)
{
stack[sPtr] = arcsin(stack[sPtr]);
}
else if (v == 13)
{
stack[sPtr] = arccos(stack[sPtr]);
}
else if (v == 14)
{
stack[sPtr] = arctg(stack[sPtr]);
}
else if (v == 15)
{
stack[sPtr] = arcctg(stack[sPtr]);
}
else if (v == 16)
{
stack[sPtr] = sh(stack[sPtr]);
}
else if (v == 17)
{
stack[sPtr] = ch(stack[sPtr]);
}
else if (v == 18)
{
stack[sPtr] = th(stack[sPtr]);
}
else if (v == 19)
{
stack[sPtr] = cth(stack[sPtr]);
}
}
else if (v <= 25)
{
if (v == 20)
{
stack[sPtr] = new ComplexD(abs(stack[sPtr]), 0);
}
if (v == 21)
{
stack[sPtr] = new ComplexD(stack[sPtr].x, 0);
}
if (v == 22)
{
stack[sPtr] = new ComplexD(stack[sPtr].y, 0);
}
if (v == 23)
{
stack[sPtr] = new ComplexD(arg(stack[sPtr]), 0);
}
if (v == 24)
{
stack[sPtr] = sqrt(stack[sPtr]);
}
if (v == 25)
{
stack[sPtr] = new ComplexD(-stack[sPtr].x, -stack[sPtr].y);
}
}
else if (v == 64)
{
stack[++sPtr] = c;
}
else if (v == 65)
{
stack[++sPtr] = z;
}
}
return(stack[stackOffset]);
}
private static ComplexD arctg(ComplexD c)
{
ComplexD x = ComplexD.Divide(new ComplexD(1 - c.y, c.x), new ComplexD(1 + c.y, -c.x));
return(ComplexD.Multiply(new ComplexD(0, -0.5), log(x)));
}
