void MxPropSlim::discontinuity_constraint(MxVertexID i, MxVertexID j, MxFaceID f)
{
Vec3 org(m->vertex(i)), dest(m->vertex(j));
Vec3 e = dest - org;
Vec3 v1(m->vertex(m->face(f)(0)));
Vec3 v2(m->vertex(m->face(f)(1)));
Vec3 v3(m->vertex(m->face(f)(2)));
Vec3 n = triangle_normal(v1,v2,v3);
Vec3 n2 = e ^ n;
unitize(n2);
MxQuadric3 Q3(n2, -(n2*org));
Q3 *= boundary_weight;
if( weighting_policy == MX_WEIGHT_AREA ||
weighting_policy == MX_WEIGHT_AREA_AVG )
{
Q3.set_area(norm2(e));
Q3 *= Q3.area();
}
MxQuadric Q(Q3, dim());
quadric(i) += Q;
quadric(j) += Q;
}
int main ()
{
struct node *Q1head;
printf("Q1 OUTPUT:\n\n");
Q1head=Q1();
printf("Q2 OUTPUT:\n\n");
Q2(Q1head);
printf("Q3 OUTPUT:\n\n");
Q3();
system("pause>nul");
}
static int
aoegen(Chan *c, char *, Dirtab *, int, int s, Dir *dp)
{
int i;
Aoedev *d;
Qid q;
if(c->qid.path == 0){
switch(s){
case DEVDOTDOT:
q.path = 0;
q.type = QTDIR;
devdir(c, q, "#æ", 0, eve, 0555, dp);
break;
case 0:
q.path = Qtopdir;
q.type = QTDIR;
devdir(c, q, "aoe", 0, eve, 0555, dp);
break;
default:
return -1;
}
return 1;
}
switch(TYPE(c->qid)){
default:
return -1;
case Qtopdir:
if(s == DEVDOTDOT){
mkqid(&q, Qzero, 0, QTDIR);
devdir(c, q, "aoe", 0, eve, 0555, dp);
return 1;
}
if(s < Qtopfiles)
return topgen(c, Qtopbase + s, dp);
s -= Qtopfiles;
if(s >= units.ref)
return -1;
mkqid(&q, QID(s, Qunitdir), 0, QTDIR);
d = unit2dev(s);
assert(d != nil);
devdir(c, q, unitname(d), 0, eve, 0555, dp);
return 1;
case Qtopctl:
case Qtoplog:
return topgen(c, TYPE(c->qid), dp);
case Qunitdir:
if(s == DEVDOTDOT){
mkqid(&q, QID(0, Qtopdir), 0, QTDIR);
uprint("%uld", UNIT(c->qid));
devdir(c, q, up->genbuf, 0, eve, 0555, dp);
return 1;
}
return unitgen(c, Qunitbase+s, dp);
case Qctl:
case Qdata:
case Qconfig:
case Qident:
return unitgen(c, TYPE(c->qid), dp);
case Qdevlinkdir:
i = UNIT(c->qid);
if(s == DEVDOTDOT){
mkqid(&q, QID(i, Qunitdir), 0, QTDIR);
devdir(c, q, "devlink", 0, eve, 0555, dp);
return 1;
}
if(i >= units.ref)
return -1;
d = unit2dev(i);
if(s >= d->ndl)
return -1;
uprint("%d", s);
mkqid(&q, Q3(s, i, Qdevlink), 0, QTFILE);
devdir(c, q, up->genbuf, 0, eve, 0755, dp);
return 1;
case Qdevlink:
uprint("%d", s);
mkqid(&q, Q3(s, UNIT(c->qid), Qdevlink), 0, QTFILE);
devdir(c, q, up->genbuf, 0, eve, 0755, dp);
return 1;
}
}
int main( int argc, char* argv[] ) {
printf("Test quaternions.\n");
{
printf("\n=== 1 =============================================\n");
CQuaternion Q1;
Dump(Q1);
}
{
printf("\n=== 2 =============================================\n");
CQuaternion Q2( CVector(0.0f,0.0f,666.0f), (float)CONST_PI );
Dump(Q2);
}
{
printf("\n=== 3 =============================================\n");
CQuaternion Q3( 1, 1, 1, 1 );
Dump(Q3);
printf("%f\n", Q3.Norm() );
printf("%f\n", Q3.Length() );
Q3.Normalize();
Dump(Q3);
}
{
printf("\n=== 4 =============================================\n");
CQuaternion Q4( 1, 1, 1, 1 );
Dump(Q4);
CQuaternion Q41 = !Q4;
Dump(Q41);
CQuaternion Q42 = -Q4;
Dump(Q42);
}
{
printf("\n=== 5 =============================================\n");
CQuaternion Q51( 1, 2, 3, 4 );
Dump(Q51);
CQuaternion Q52( 1, 2, 3, -4 );
Dump(Q52);
printf("%f\n", Q51^Q52 );
}
{
printf("\n=== 6 =============================================\n");
CQuaternion Q61( 1, 2, 3, 4 );
Dump(Q61);
CQuaternion Q62( 5, 6, 7, 8 );
Dump(Q62);
Dump( Q61 + Q62 );
Dump( Q61 - Q62 );
Dump( Q61 += Q62 );
Dump( Q61 -= Q62 );
}
{
printf("\n=== 7 =============================================\n");
CQuaternion Q71( 1, 2, 3, 4 );
Q71.Normalize();
Dump(Q71);
printf("%f\n",Q71.Length());
//CQuaternion Q72( 5, 6, 7, 8 );
//Q72.Normalize();
//Dump(Q72);
//Dump( Q71 * Q72 );
CQuaternion Q73 = -Q71;
Dump( Q73 );
printf("%f\n",Q73.Length());
CQuaternion Product = Q71 * Q73;
Dump( Product );
printf("%f\n",Product.Length());
Dump( Q71*=Q73 );
}
{
printf("\n=== 8 =============================================\n");
CQuaternion Q8( CVector(0.0f,0.0f,1.0f), (float)CONST_PI_2 );
//CQuaternion Q8( CVector(0,0,1), CONST_PI );
Dump(Q8);
CVector Src(1,0,0);
CVector Dst = RotateVectorByQuaternion( Src, Q8 );
printf("%f %f %f\n",Dst.x,Dst.y,Dst.z);
}
{
printf("\n=== 9 =============================================\n");
CQuaternion Q81( CVector(0,0,1), 0 );
Dump(Q81);
CQuaternion Q82( CVector(0,0,1), CONST_PI );
Dump(Q82);
for( float t=0.0f; t<1.0f; t+=0.1f ) {
CQuaternion R = SLerp(Q81,Q82,t);
Dump( R );
}
}
{
printf("\n=== 10 ============================================\n");
CQuaternion Q10( CVector(0,0,1), CONST_PI_2 );
//CQuaternion Q10( CVector(0,0,1), CONST_PI );
Dump(Q10);
CVector V;
float A;
Q10.ToAxisAngle( V, A );
printf("%f %f %f\n",V.x,V.y,V.z);
printf("%f\n",A);
CVector Src(1,0,0);
CVector Dst1 = RotateVectorByQuaternion( Src, Q10 );
printf("%f %f %f\n",Dst1.x,Dst1.y,Dst1.z);
CVector Dst11 = Src*Q10;
printf("%f %f %f\n",Dst11.x,Dst11.y,Dst11.z);
CMatrix M( Q10.ToMatrix() );
CVector Dst2 = Src*M;
printf("%f %f %f\n",Dst2.x,Dst2.y,Dst2.z);
CMatrix M3;
M3.ConstructRotation( CVector(0,0,1), CONST_PI_2 );
CVector Dst3 = Src*M3;
printf("%f %f %f\n",Dst3.x,Dst3.y,Dst3.z);
CMatrix M4;
M4.ConstructRotationZ( CONST_PI_2 );
CVector Dst4 = Src*M4;
printf("%f %f %f\n",Dst4.x,Dst4.y,Dst4.z);
}
{
printf("\n=== 11 =============================================\n");
CMatrix M11;
M11.ConstructRotation( CVector(0.3f,-0.78f,-0.9f), 0.666 );
CQuaternion Q11 = CreateNonUnitQuaternionFromRotationMatrix( M11 );
Q11.Normalize();
CQuaternion Q111 = CreateUnitQuaternionFromRotationMatrix( M11 );
Q111.Normalize();
CVector V11( 0.0f, 0.0f, 1.0f );
CVector R1 = V11*M11;
printf("%f %f %f\n",R1.x,R1.y,R1.z);
CVector R2 = RotateVectorByQuaternion( V11, Q11 );
printf("%f %f %f\n",R2.x,R2.y,R2.z);
CVector R3 = RotateVectorByQuaternion( V11, Q111 );
printf("%f %f %f\n",R3.x,R3.y,R3.z);
}
{
printf("\n=== 12 =============================================\n");
CQuaternion Q1( CVector(1.0f,1.0f,1.0f), (float)CONST_PI_3 );
Q1.Normalize();
CQuaternion Q2( CVector(1.0f,1.0f,1.0f), (float)CONST_PI_3 );
Q2.Normalize();
CQuaternion Q3 = SLerp(Q1,Q2,0.5f);
Dump(Q3);
}
}
Func ColorMgather(Func stBasis, float angle, uint8_t * orders, Expr filterthreshold, Expr divisionthreshold, Expr divisionthreshold2) {
uint8_t x_order = orders[0];
uint8_t y_order = orders[1];
uint8_t t_order = orders[2];
uint8_t c_order = orders[3];
Func X("X"),Y("Y"),T("T"),Xrg("Xrg"),Yrg("Yrg"),Trg("Trg");
uint8_t max_order = x_order;
// std::vector<Expr>Xk_expr (max_order,cast<float>(0.0f));
// std::vector<Expr>Yk_expr (max_order,cast<float>(0.0f));
// std::vector<Expr>Tk_expr (max_order,cast<float>(0.0f));
uint8_t Xk_uI[max_order];
uint8_t Yk_uI[max_order];
uint8_t Tk_uI[max_order];
Func Xk[max_order]; Func Yk[max_order]; Func Tk[max_order];
// Expr Xk[max_order],Yk[max_order],Tk[max_order];
for (int iO=0; iO < x_order; iO++) {
Xk[iO](x,y,t) = Expr(0.0f);
Yk[iO](x,y,t) = Expr(0.0f);
Tk[iO](x,y,t) = Expr(0.0f);
Xk_uI[iO] = 0;
Yk_uI[iO] = 0;
Tk_uI[iO] = 0;
}
int k = 0;
for (int iXo = 0; iXo < x_order; iXo++) // x_order
for (int iYo = 0; iYo < y_order; iYo++) // y_oder
for (int iTo = 0; iTo < t_order; iTo++) // t_order
for (int iCo = 0; iCo < c_order; iCo ++ ) // c_order: index of color channel
{
if ((iYo+iTo+iCo == 0 || iYo+iTo+iCo == 1)
&& ((iXo+iYo+iTo+iCo+1) < (x_order + 1))) {
X = ColorMgetfilter(stBasis, angle, iXo+1, iYo, iTo, iCo);
Y = ColorMgetfilter(stBasis, angle, iXo, iYo+1, iTo, iCo);
T = ColorMgetfilter(stBasis, angle, iXo, iYo, iTo+1, iCo);
Xrg = ColorMgetfilter(stBasis, angle, iXo+1, iYo, iTo, iCo+1);
Yrg = ColorMgetfilter(stBasis, angle, iXo, iYo+1, iTo, iCo+1);
Trg = ColorMgetfilter(stBasis, angle, iXo, iYo, iTo+1, iCo+1);
k = iXo + iYo + iTo + iCo;
Xk[k](x,y,t) += X(x,y,t) + Xrg(x,y,t);
Yk[k](x,y,t) += Y(x,y,t) + Yrg(x,y,t);
Tk[k](x,y,t) += T(x,y,t) + Trg(x,y,t);
Xk[k].update(Xk_uI[k]); Xk_uI[k]++;
Yk[k].update(Yk_uI[k]); Yk_uI[k]++;
Tk[k].update(Tk_uI[k]); Tk_uI[k]++;
}
}
// Scheduling
for (int iO = 0; iO <= k; iO++) {
Xk[iO].compute_root();
Yk[iO].compute_root();
Tk[iO].compute_root();
}
std::vector<Expr> st_expr(6,cast<float>(0.0f));
for (int iK=0; iK <= k; iK++) {
st_expr[0] += Xk[iK](x,y,t)*Tk[iK](x,y,t);
st_expr[1] += Tk[iK](x,y,t)*Tk[iK](x,y,t);
st_expr[2] += Xk[iK](x,y,t)*Xk[iK](x,y,t);
st_expr[3] += Yk[iK](x,y,t)*Tk[iK](x,y,t);
st_expr[4] += Yk[iK](x,y,t)*Yk[iK](x,y,t);
st_expr[5] += Xk[iK](x,y,t)*Yk[iK](x,y,t);
}
Func st("st"); st(x,y,t) = Tuple(st_expr);
st.compute_root();
Expr x_clamped = clamp(x,0,width-1);
Expr y_clamped = clamp(y,0,height-1);
Func st_clamped("st_clamped"); st_clamped(x,y,t) = st(x_clamped,y_clamped,t);
// float win = 7.0;
// Image<float> meanfilter(7,7,"meanfilter_data");
// meanfilter(x,y) = Expr(1.0f/(win*win));
// RDom rMF(meanfilter);
uint8_t win = 7;
RDom rMF(0,win,0,win);
Func st_filtered[6];
for (uint8_t iPc=0; iPc<6; iPc++) {
// iPc: index of product component
// Apply average filter
st_filtered[iPc](x,y,t) = sum(rMF,st_clamped(x + rMF.x,y + rMF.y,t)[iPc]/Expr(float(win*win)),"mean_filter");
st_filtered[iPc].compute_root();
}
// Tuple st_tuple = Tuple(st_expr);
// 4 debug
// Func tmpOut("tmpOut"); tmpOut(x,y,t) = Tuple(st_filtered[0](x,y,t),st_filtered[1](x,y,t),st_filtered[2](x,y,t),st_filtered[3](x,y,t),st_filtered[4](x,y,t),st_filtered[5](x,y,t));
// return tmpOut;
Tuple pbx = Tuple(st_filtered[2](x,y,t),st_filtered[5](x,y,t),st_filtered[0](x,y,t));
Tuple pby = Tuple(st_filtered[5](x,y,t),st_filtered[4](x,y,t),st_filtered[3](x,y,t));
Tuple pbt = Tuple(st_filtered[0](x,y,t),st_filtered[3](x,y,t),st_filtered[1](x,y,t));
Func pbxy("pbxy"); pbxy = cross(pby,pbx); pbxy.compute_root();
Func pbxt("pbxt"); pbxt = cross(pbx,pbt); pbxt.compute_root();
Func pbyt("pbyt"); pbyt = cross(pby,pbt); pbyt.compute_root();
Func pbxyd("pbxyd"); pbxyd = dot(pby,pbx); pbxyd.compute_root();
Func pbxtd("pbxtd"); pbxtd = dot(pbx,pbt); pbxtd.compute_root();
Func pbytd("pbytd"); pbytd = dot(pby,pbt); pbytd.compute_root();
// 4 debug
// Func tmpOut("tmpOut"); tmpOut(x,y,t) = Tuple(pbxy(x,y,t)[0],pbxt(x,y,t)[0],pbyt(x,y,t)[0],pbxyd(x,y,t),pbxtd(x,y,t),pbytd(x,y,t));
// return tmpOut;
Func yt_xy("yt_xy"); yt_xy = dot(pbyt(x,y,t),pbxy(x,y,t)); yt_xy.compute_root();
Func xt_yt("xt_yt"); xt_yt = dot(pbxt(x,y,t),pbyt(x,y,t)); xt_yt.compute_root();
Func xt_xy("xt_xy"); xt_xy = dot(pbxt(x,y,t),pbxy(x,y,t)); xt_xy.compute_root();
Func yt_yt("yt_yt"); yt_yt = dot(pbyt(x,y,t),pbyt(x,y,t)); yt_yt.compute_root();
Func xt_xt("xt_xt"); xt_xt = dot(pbxt(x,y,t),pbxt(x,y,t)); xt_xt.compute_root();
Func xy_xy("xy_xy"); xy_xy = dot(pbxy(x,y,t),pbxy(x,y,t)); xy_xy.compute_root();
Tuple Tk_tuple = Tuple(Tk[0](x,y,t),Tk[1](x,y,t),Tk[2](x,y,t),
Tk[3](x,y,t),Tk[4](x,y,t));
Func Tkd("Tkd"); Tkd = dot(Tk_tuple,Tk_tuple); Tkd.compute_root();
// Expr Dimen = pbxyd/xy_xy;
Expr kill(1.0f);
Func Oxy; Oxy(x,y,t) = Mdefdiv(st_filtered[5](x,y,t) - Mdefdivang(yt_xy(x,y,t),yt_yt(x,y,t),pbxyd(x,y,t),divisionthreshold2)*st_filtered[3](x,y,t)*kill,st_filtered[4](x,y,t),divisionthreshold);
Oxy.compute_root();
Func Oyx; Oyx(x,y,t) = Mdefdiv(st_filtered[5](x,y,t) + Mdefdivang(xt_xy(x,y,t),xt_xt(x,y,t),pbxyd(x,y,t),divisionthreshold2)*st_filtered[0](x,y,t)*kill,st_filtered[2](x,y,t),divisionthreshold);
Oyx.compute_root();
Func C0; C0(x,y,t) = st_filtered[3](x,y,t) * Mdefdivang(Expr(-1.0f)*xt_yt(x,y,t),yt_yt(x,y,t),pbxyd(x,y,t),divisionthreshold2)*kill;
C0.compute_root();
Func M0; M0(x,y,t) = Mdefdiv(st_filtered[0](x,y,t) + C0(x,y,t), st_filtered[1](x,y,t)*pow(Mdefdivang(xt_yt(x,y,t),yt_yt(x,y,t),pbxyd(x,y,t),divisionthreshold2),Expr(2.0f)),divisionthreshold);
M0.compute_root();
Func C1; C1(x,y,t) = st_filtered[5](x,y,t) * Mdefdivang(Expr(-1.0f)*xt_xy(x,y,t),xy_xy(x,y,t),pbxyd(x,y,t),divisionthreshold2)*kill;
C1.compute_root();
Func P1; P1(x,y,t) = pow(Mdefdivang(xt_yt(x,y,t),xt_xt(x,y,t),pbxyd(x,y,t),divisionthreshold2),Expr(2.0f))*kill + 1.0f;
P1.compute_root();
// 4 debug
// Func tmpOut("tmpOut"); tmpOut(x,y,t) = Tuple(Oxy(x,y,t),Oyx(x,y,t),C0(x,y,t),M0(x,y,t),C1(x,y,t),P1(x,y,t));
// return tmpOut;
Func Q1; Q1(x,y,t) = st_filtered[2](x,y,t) * (pow(Oyx(x,y,t),Expr(2.0f))+Expr(1.0f));
Q1.compute_root();
Func M1; M1(x,y,t) = Mdefdiv(((st_filtered[0](x,y,t)-C1(x,y,t))*P1(x,y,t)),Q1(x,y,t),divisionthreshold);
M1.compute_root();
Func C2; C2(x,y,t) = st_filtered[0](x,y,t) * Mdefdivang(Expr(-1.0f)*xt_yt(x,y,t),xt_xt(x,y,t),pbxyd(x,y,t),divisionthreshold2)*kill;
C2.compute_root();
Func M2; M2(x,y,t) = Mdefdiv(st_filtered[3](x,y,t)+C2(x,y,t),st_filtered[1](x,y,t)*(pow(Mdefdivang(xt_yt(x,y,t),xt_xt(x,y,t),pbxyd(x,y,t),divisionthreshold2),Expr(2.0f))*kill+Expr(1.0f)),divisionthreshold);
M2.compute_root();
Func C3; C3(x,y,t) = st_filtered[5](x,y,t) * Mdefdivang(yt_xy(x,y,t),xy_xy(x,y,t),pbxyd(x,y,t),divisionthreshold2)*kill;
C3.compute_root();
Func P3; P3(x,y,t) = pow(Mdefdivang(xt_yt(x,y,t),yt_yt(x,y,t),pbxyd(x,y,t),divisionthreshold2),Expr(2.0f))*kill + Expr(1.0f);
P3.compute_root();
Func Q3; Q3(x,y,t) = st_filtered[4](x,y,t) * (pow(Oxy(x,y,t),Expr(2.0f))+Expr(1.0f));
Q3.compute_root();
Func M3; M3(x,y,t) = Mdefdiv(((st_filtered[3](x,y,t)-C3(x,y,t))*P3(x,y,t)),Q3(x,y,t),divisionthreshold);
M3.compute_root();
Func basisAtAngle;
basisAtAngle(x,y,t) = Tuple(M0(x,y,t),M1(x,y,t),M2(x,y,t),M3(x,y,t),Tkd(x,y,t));
return basisAtAngle;
// Func hsv2rgb(Func colorImage) { // Took this function
// Var x, y, c, t;
// Func output;
// output(x,y,c,t) = cast <float> (0.0f);
// Expr fR, fG, fB; // R,G & B values
// Expr fH = (colorImage(x,y,0,t)); //H value [0-360)
// Expr fS = (colorImage(x,y,1,t)); //S value
// Expr fV = (colorImage(x,y,2,t)); //V value
// //Conversion (I took the one on Wikipedia)
// // https://fr.wikipedia.org/wiki/Teinte_Saturation_Valeur#Conversion_de_TSV_vers_RVB
// Expr fHi = floor(fH / Expr(60.0f));
// Expr fF = fH / 60.0f - fHi;
// Expr fL = fV * (1 - fS);
// Expr fM = fV * (1 - fF * fS) ;
// Expr fN = fV * (1 - (1 - fF) * fS);
// fR = select((0 == fHi),fV,
// (1 == fHi),fM,
// (2 == fHi),fL,
// (3 == fHi),fL,
// (4 == fHi),fN,
// (5 == fHi),fV,
// 0.0f);
// fG = select((0 == fHi),fN,
// (1 == fHi),fV,
// (2 == fHi),fV,
// (3 == fHi),fM,
// (4 == fHi),fL,
// (5 == fHi),fL,
// 0.0f);
// fB = select((0 == fHi),fL,
// (1 == fHi),fL,
// (2 == fHi),fN,
// (3 == fHi),fV,
// (4 == fHi),fV,
// (5 == fHi),fM,
// 0.0f);
// output(x,y,0,t) = fR;
// output(x,y,1,t) = fG;
// output(x,y,2,t) = fB;
// return output;
// }
// Func angle2rgb (Func v) {
// Var x, y, c, t;
// Func ov, a;
// ov(x,y,c,t) = cast <float> (0.0f);
// Expr pi2(2*M_PI);
// a(x,y,c,t) = v(x,y,c,t) / pi2;
// ov(x,y,0,t) = a(x,y,c,t);
// ov(x,y,1,t) = 1;
// ov(x,y,2,t) = 1;
// return ov;
// }
// Func outputvelocity(Func Blur, Func Speed, Func Angle, int border, Expr speedthreshold, Expr filterthreshold) {
// extern Expr width;
// extern Expr height;
// Func Blur3, Speed3;
// Blur3(x,y,c,t) = cast <float> (0.0f);
// Speed3(x,y,c,t) = cast <float> (0.0f);
// //Scale the grey level images
// Blur(x,y,0,t) = (Blur(x,y,0,t) - minimum(Blur(x,y,0,t))) / (maximum(Blur(x,y,0,t)) - minimum(Blur(x,y,0,t)));
// //Concatenation along the third dimension
// Blur3(x,y,0,t) = Blur(x,y,0,t);
// Blur3(x,y,1,t) = Blur(x,y,0,t);
// Blur3(x,y,2,t) = Blur(x,y,0,t);
// //Speed scaled to 1
// //Concatenation along the third dimension
// Speed3(x,y,1,t) = Speed(x,y,0,t);
// Speed3(x,y,2,t) = Speed(x,y,0,t);
// //Use the log speed to visualise speed
// Func LogSpeed;
// LogSpeed(x,y,c,t) = fast_log(Speed3(x,y,c,t) + Expr(0.0000001f))/fast_log(Expr(10.0f));
// LogSpeed(x,y,c,t) = (LogSpeed(x,y,c,t) - minimum(LogSpeed(x,y,c,t))) / (maximum(LogSpeed(x,y,c,t)) - minimum(LogSpeed(x,y,c,t)));
// //Make a colour image
// // uint16_t rows = height;
// // uint16_t cols = width;
// // int depth = Angle.channels();
// //Do it the HSV way
// Func colorImage;
// colorImage(x,y,0,t) = Angle(x,y,0,t);
// //Do hsv to rgb
// Func colorImage1;
// colorImage1 = hsv2rgb(colorImage);
// // Assume the border equals to the size of spatial filter
// //Make the border
// // int bir = rows + 2 * border;
// // int bic = cols + 2 * border;
// Expr orows = height / Expr(2);
// Expr ocols = width / Expr(2);
// //Rotation matrix
// int ph = 0;
// Func mb, sb;
// // if (rx < border - 1 || rx >= rows+border -1 || ry < border - 1 || ry >= cols+border - 1) {
// Expr co1 = x - orows;
// Expr co2 = - (y - ocols);
// Expr cosPh(cos(ph));
// Expr sinPh(sin(ph));
// Expr rco1 = cosPh * co1 - sinPh * co2; //Using rotation matrix
// Expr rco2 = sinPh * co1 + cosPh * co2;
// // Expr justPi (M_PI);
// mb(x,y,c,t) =
// select (((x < (border - 1)) ||
// (x >= (height+border -1)) ||
// (y < (border - 1)) ||
// (y >= (width+border - 1))),
// atan2(rco1,rco2) + Expr(M_PI),mb(x,y,c,t));
// sb(x,y,c,t) =
// select (((x < (border - 1)) ||
// (x >= (height+border -1)) ||
// (y < (border - 1) ) ||
// (y >= (width+border - 1))),
// 1, sb (x,y,c,t));
// Func cb;
// cb = angle2rgb(mb);
// //Get the old data
// // Expr pi2(2*M_PI);
// colorImage1(x,y,0,t)=colorImage(x,y,0,t) * Expr(2*M_PI);
// colorImage1=angle2rgb(colorImage1);
// colorImage1(x,y,c,t)=select(abs(Speed3(x,y,c,t))<speedthreshold,Expr(0.0f),colorImage1(x,y,c,t));
// Func colorImage2;
// colorImage2(x,y,c,t) = colorImage1(x,y,c,t) * Speed(x,y,c,t);
// //Put the data in the border
// RDom bordx (border,rows + border);
// RDom bordy (border,cols + border);
// Func ang1, ang2;
// ang1 (x,y,c,t) = cast <float> (0.0f);
// ang2 (x,y,c,t) = cast <float> (0.0f);
// cb(bordx, bordy,c,t) = colorImage1(x,y,c,t);
// ang1 = cb;
// cb(bordx, bordy,c,t) = colorImage2(x,y,c,t);
// ang2 = cb;
// sb(bordx, bordy,c,t) = Speed3(x,y,c,t);
// Speed3 = sb;
// sb(bordx, bordy,c,t) = Blur3(x,y,c,t);
// Blur3 = sb;
// // Func I;
// // I (x,y,c,t) = Blur3(x,y,c,t) + Speed3(x,y - height,c,t) + ang1(x - width,y,c,t) + ang2(x - width,y - height,c,t);
// //I = cat(2,cat(1,Blur,Speed),cat(1,ang1,ang2));
// return I;
// }
}
int main(int argc,char **argv){
// Print GPU properties
//print_properties();
// Files to print the result after the last time step
FILE *rho_file;
FILE *E_file;
rho_file = fopen("rho_final.txt", "w");
E_file = fopen("E_final.txt", "w");
// Construct initial condition for problem
ICsinus Config(-1.0, 1.0, -1.0, 1.0);
//ICsquare Config(0.5,0.5,gasGam);
// Set initial values for Configuration 1
/*
Config.set_rho(rhoConfig19);
Config.set_pressure(pressureConfig19);
Config.set_u(uConfig19);
Config.set_v(vConfig19);
*/
// Determining global border based on left over tiles (a little hack)
int globalPadding;
globalPadding = (nx+2*border+16)/16;
globalPadding = 16*globalPadding - (nx+2*border);
//printf("Globalpad: %i\n", globalPadding);
// Change border to add padding
//border = border + globalPadding/2;
// Initiate the matrices for the unknowns in the Euler equations
cpu_ptr_2D rho(nx, ny, border,1);
cpu_ptr_2D E(nx, ny, border,1);
cpu_ptr_2D rho_u(nx, ny, border,1);
cpu_ptr_2D rho_v(nx, ny, border,1);
cpu_ptr_2D zeros(nx, ny, border,1);
// Set initial condition
Config.setIC(rho, rho_u, rho_v, E);
double timeStart = get_wall_time();
// Test
cpu_ptr_2D rho_dummy(nx, ny, border);
cpu_ptr_2D E_dummy(nx, ny, border);
/*
rho_dummy.xmin = -1.0;
rho_dummy.ymin = -1.0;
E_dummy.xmin = -1.0;
E_dummy.ymin = -1.0;
*/
// Set block and grid sizes
dim3 gridBC = dim3(1, 1, 1);
dim3 blockBC = dim3(BLOCKDIM_BC,1,1);
dim3 gridBlockFlux;
dim3 threadBlockFlux;
dim3 gridBlockRK;
dim3 threadBlockRK;
computeGridBlock(gridBlockFlux, threadBlockFlux, nx + 2*border, ny + 2*border, INNERTILEDIM_X, INNERTILEDIM_Y, BLOCKDIM_X, BLOCKDIM_Y);
computeGridBlock(gridBlockRK, threadBlockRK, nx + 2*border, ny + 2*border, BLOCKDIM_X_RK, BLOCKDIM_Y_RK, BLOCKDIM_X_RK, BLOCKDIM_Y_RK);
int nElements = gridBlockFlux.x*gridBlockFlux.y;
// Allocate memory for the GPU pointers
gpu_ptr_1D L_device(nElements);
gpu_ptr_1D dt_device(1);
gpu_ptr_2D rho_device(nx, ny, border);
gpu_ptr_2D E_device(nx, ny, border);
gpu_ptr_2D rho_u_device(nx, ny, border);
gpu_ptr_2D rho_v_device(nx, ny, border);
gpu_ptr_2D R0(nx, ny, border);
gpu_ptr_2D R1(nx, ny, border);
gpu_ptr_2D R2(nx, ny, border);
gpu_ptr_2D R3(nx, ny, border);
gpu_ptr_2D Q0(nx, ny, border);
gpu_ptr_2D Q1(nx, ny, border);
gpu_ptr_2D Q2(nx, ny, border);
gpu_ptr_2D Q3(nx, ny, border);
// Allocate pinned memory on host
init_allocate();
// Set BC arguments
set_bc_args(BCArgs[0], rho_device.getRawPtr(), rho_u_device.getRawPtr(), rho_v_device.getRawPtr(), E_device.getRawPtr(), nx+2*border, ny+2*border, border);
set_bc_args(BCArgs[1], Q0.getRawPtr(), Q1.getRawPtr(), Q2.getRawPtr(), Q3.getRawPtr(), nx+2*border, ny+2*border, border);
set_bc_args(BCArgs[2], rho_device.getRawPtr(), rho_u_device.getRawPtr(), rho_v_device.getRawPtr(), E_device.getRawPtr(), nx+2*border, ny+2*border, border);
// Set FLUX arguments
set_flux_args(fluxArgs[0], L_device.getRawPtr(), rho_device.getRawPtr(), rho_u_device.getRawPtr(), rho_v_device.getRawPtr(), E_device.getRawPtr(), R0.getRawPtr(),R1.getRawPtr(), R2.getRawPtr(), R3.getRawPtr(), nx, ny, border, rho.get_dx(), rho.get_dy(), theta, gasGam, INNERTILEDIM_X, INNERTILEDIM_Y);
set_flux_args(fluxArgs[1], L_device.getRawPtr(), Q0.getRawPtr(), Q1.getRawPtr(), Q2.getRawPtr(), Q3.getRawPtr(), R0.getRawPtr(),R1.getRawPtr(), R2.getRawPtr(), R3.getRawPtr(), nx, ny, border, rho.get_dx(), rho.get_dy(), theta, gasGam, INNERTILEDIM_X, INNERTILEDIM_Y);
// Set TIME argument
set_dt_args(dtArgs, L_device.getRawPtr(), dt_device.getRawPtr(), nElements, rho.get_dx(), rho.get_dy(), cfl_number);
// Set Rk arguments
set_rk_args(RKArgs[0], dt_device.getRawPtr(), rho_device.getRawPtr(), rho_u_device.getRawPtr(), rho_v_device.getRawPtr(), E_device.getRawPtr(), R0.getRawPtr(), R1.getRawPtr(), R2.getRawPtr(), R3.getRawPtr(), Q0.getRawPtr(), Q1.getRawPtr(), Q2.getRawPtr(), Q3.getRawPtr(), nx, ny, border);
set_rk_args(RKArgs[1], dt_device.getRawPtr(), Q0.getRawPtr(), Q1.getRawPtr(), Q2.getRawPtr(), Q3.getRawPtr(), R0.getRawPtr(), R1.getRawPtr(), R2.getRawPtr(), R3.getRawPtr(), rho_device.getRawPtr(), rho_u_device.getRawPtr(), rho_v_device.getRawPtr(), E_device.getRawPtr(), nx, ny, border);
L_device.set(FLT_MAX);
/*
R0.upload(zeros.get_ptr());
R1.upload(zeros.get_ptr());
R2.upload(zeros.get_ptr());
R3.upload(zeros.get_ptr());
Q0.upload(zeros.get_ptr());
Q1.upload(zeros.get_ptr());
Q2.upload(zeros.get_ptr());
Q3.upload(zeros.get_ptr());
*/
R0.set(0,0,0,nx,ny,border);
R1.set(0,0,0,nx,ny,border);
R2.set(0,0,0,nx,ny,border);
R3.set(0,0,0,nx,ny,border);
Q0.set(0,0,0,nx,ny,border);
Q1.set(0,0,0,nx,ny,border);
Q2.set(0,0,0,nx,ny,border);
Q3.set(0,0,0,nx,ny,border);
rho_device.upload(rho.get_ptr());
rho_u_device.upload(rho_u.get_ptr());
rho_v_device.upload(rho_v.get_ptr());
E_device.upload(E.get_ptr());
// Update boudries
callCollectiveSetBCPeriodic(gridBC, blockBC, BCArgs[0]);
//Create cuda stream
cudaStream_t stream1;
cudaStreamCreate(&stream1);
cudaEvent_t dt_complete;
cudaEventCreate(&dt_complete);
while (currentTime < timeLength && step < maxStep){
//RK1
//Compute flux
callFluxKernel(gridBlockFlux, threadBlockFlux, 0, fluxArgs[0]);
// Compute timestep (based on CFL condition)
callDtKernel(TIMETHREADS, dtArgs);
cudaMemcpyAsync(dt_host, dt_device.getRawPtr(), sizeof(float), cudaMemcpyDeviceToHost, stream1);
cudaEventRecord(dt_complete, stream1);
// Perform RK1 step
callRKKernel(gridBlockRK, threadBlockRK, 0, RKArgs[0]);
//Update boudries
callCollectiveSetBCPeriodic(gridBC, blockBC, BCArgs[1]);
//RK2
// Compute flux
callFluxKernel(gridBlockFlux, threadBlockFlux, 1, fluxArgs[1]);
//Perform RK2 step
callRKKernel(gridBlockRK, threadBlockRK, 1, RKArgs[1]);
//cudaEventRecord(srteam_sync, srteam1);
callCollectiveSetBCPeriodic(gridBC, blockBC, BCArgs[2]);
cudaEventSynchronize(dt_complete);
step++;
currentTime += *dt_host;
// printf("Step: %i, current time: %.6f dt:%.6f\n" , step,currentTime, dt_host[0]);
}
//cuProfilerStop();
//cudaProfilerStop();
printf("Elapsed time %.5f", get_wall_time() - timeStart);
E_device.download(E.get_ptr());
rho_u_device.download(rho_u.get_ptr());
rho_v_device.download(rho_v.get_ptr());
rho_device.download(rho_dummy.get_ptr());
rho_dummy.printToFile(rho_file, true, false);
Config.exactSolution(E_dummy, currentTime);
E_dummy.printToFile(E_file, true, false);
float LinfError = Linf(E_dummy, rho_dummy);
float L1Error = L1(E_dummy, rho_dummy);
float L1Error2 = L1test(E_dummy, rho_dummy);
printf("nx: %i\t Linf error %.9f\t L1 error %.7f L1test erro %.7f", nx, LinfError, L1Error, L1Error2);
printf("nx: %i step: %i, current time: %.6f dt:%.6f\n" , nx, step,currentTime, dt_host[0]);
/*
cudaMemcpy(L_host, L_device, sizeof(float)*(nElements), cudaMemcpyDeviceToHost);
for (int i =0; i < nElements; i++)
printf(" %.7f ", L_host[i]);
*/
printf("%s\n", cudaGetErrorString(cudaGetLastError()));
return(0);
}
