This commit is contained in:
2026-04-10 13:13:49 -04:00
parent 8fa5ac61a6
commit 33198f9f45
25 changed files with 236 additions and 220 deletions

Binary file not shown.

Binary file not shown.

Binary file not shown.

Binary file not shown.

Binary file not shown.

Binary file not shown.

Binary file not shown.

Binary file not shown.

Binary file not shown.

Binary file not shown.

Binary file not shown.

Binary file not shown.

Binary file not shown.

Binary file not shown.

View File

@ -92,34 +92,11 @@ namespace amscuda
__host__ __device__ cuvect3f hodge_dual(const cumat3f &min);
__host__ __device__ cumat3f rotmat_from_axisangle(const cuvect3f &axis, const float &angle);
//3x3 matrix operations
//matrix order is assumed to be mat[I,J] = mat[I+3*J]
//transposes a 3x3 (9 element) matrix
__host__ __device__ void mat3f_transpose(float *mat3inout);
//copies src to dest
__host__ __device__ void mat3f_copy(float *mat3f_dest, const float *mat3f_src);
//returns determinant of 3x3 matrix
__host__ __device__ float mat3f_det(float *mat3in);
//inverts a 3x3 (9 element) matrix
__host__ __device__ void mat3f_inverse(float *mat3inout);
__host__ __device__ cuvect3f mat3f_mult(float *mat3in, const cuvect3f &cvin);
__host__ __device__ void mat3f_mult(float *matina, float *matinb, float *matout);
__host__ __device__ void mat3f_hodgedual(const cuvect3f &vecin, float *matout);
__host__ __device__ void mat3f_hodgedual(float *matin, cuvect3f &vecout);
//returns direction cosine rotation matrix from axis and angle
__host__ __device__ void mat3f_rot_from_axisangle(cuvect3f axis, float angle, float *matout);
__host__ void test_cudavectf_logic1();
};
#include <amsculib2/cuvect3f_legacy.hpp>
#endif

View File

@ -0,0 +1,34 @@
#ifndef __CUVECT3_LEGACY_HPP__
#define __CUVECT3_LEGACY_HPP__
namespace amscuda
{
//3x3 matrix operations
//matrix order is assumed to be mat[I,J] = mat[I+3*J]
//transposes a 3x3 (9 element) matrix
__host__ __device__ void mat3f_transpose(float *mat3inout);
//copies src to dest
__host__ __device__ void mat3f_copy(float *mat3f_dest, const float *mat3f_src);
//returns determinant of 3x3 matrix
__host__ __device__ float mat3f_det(float *mat3in);
//inverts a 3x3 (9 element) matrix
__host__ __device__ void mat3f_inverse(float *mat3inout);
__host__ __device__ cuvect3f mat3f_mult(float *mat3in, const cuvect3f &cvin);
__host__ __device__ void mat3f_mult(float *matina, float *matinb, float *matout);
__host__ __device__ void mat3f_hodgedual(const cuvect3f &vecin, float *matout);
__host__ __device__ void mat3f_hodgedual(float *matin, cuvect3f &vecout);
//returns direction cosine rotation matrix from axis and angle
__host__ __device__ void mat3f_rot_from_axisangle(cuvect3f axis, float angle, float *matout);
};
#endif

View File

@ -85,10 +85,10 @@ namespace amscuda
__host__ __device__ cumat4f& operator*=(const cumat4f &rhs);
};
__host__ __device__ float cuvect4f_dot(cuvect4f a, cuvect4f b);
__host__ __device__ float cuvect4f_norm(cuvect4f a);
__host__ __device__ cuvect4f cuvect4f_normalize(cuvect4f a);
__host__ __device__ cuvect4f cuvect4f_proj(cuvect4f a, cuvect4f b);
__host__ __device__ float cuvect4f_dot(cuvect4f &a, cuvect4f &b);
__host__ __device__ float cuvect4f_norm(cuvect4f &a);
__host__ __device__ cuvect4f cuvect4f_normalize(cuvect4f &a);
__host__ __device__ cuvect4f cuvect4f_proj(cuvect4f &a, cuvect4f &b);
};

View File

@ -5,13 +5,13 @@ namespace amscuda
__host__ __device__ cuvect3f::cuvect3f()
{
x = 0.0; y = 0.0; z = 0.0;
x = 0.0f; y = 0.0f; z = 0.0f;
return;
}
__host__ __device__ cuvect3f::~cuvect3f()
{
x = 0.0; y = 0.0; z = 0.0;
x = 0.0f; y = 0.0f; z = 0.0f;
return;
}
@ -144,7 +144,7 @@ namespace amscuda
}
else
{
ret.x = 0.0; ret.y = 0.0; ret.z = 0.0;
ret.x = 0.0f; ret.y = 0.0f; ret.z = 0.0f;
}
return ret;
@ -390,7 +390,7 @@ __host__ __device__ cumat3f cumat3f::operator*(const cumat3f &rhs)
__host__ __device__ float cumat3f::det()
{
float ret = 0.0;
float ret = 0.0f;
ret += m00*m11*m22;
ret += m01*m12*m20;
@ -587,186 +587,6 @@ __host__ __device__ const float* cumat3f::data() const
return (const float*) this;
}
//transposes a 3x3 (9 element) matrix
__host__ __device__ void mat3f_transpose(float *mat3inout)
{
int I,J;
float matint[9];
for(I=0;I<9;I++)
{
matint[I] = mat3inout[I];
}
for(I=0;I<3;I++)
{
for(J=0;J<3;J++)
{
mat3inout[I+J*3] = matint[J+I*3];
}
}
return;
}
//copies src to dest
__host__ __device__ void mat3f_copy(float *mat3f_dest, const float *mat3f_src)
{
int I;
if(mat3f_dest==NULL || mat3f_src==NULL)
return;
for(I=0;I<9;I++)
mat3f_dest[I] = mat3f_src[I];
return;
}
__host__ __device__ float mat3f_det(float *mat3in)
{
float ret = 0.0;
ret = ret + mat3in[0+0*3]*mat3in[1+1*3]*mat3in[2+2*3];
ret = ret + mat3in[0+1*3]*mat3in[1+2*3]*mat3in[2+0*3];
ret = ret + mat3in[0+2*3]*mat3in[1+0*3]*mat3in[2+1*3];
ret = ret - mat3in[0+0*3]*mat3in[1+2*3]*mat3in[2+1*3];
ret = ret - mat3in[0+1*3]*mat3in[1+0*3]*mat3in[2+2*3];
ret = ret - mat3in[0+2*3]*mat3in[1+1*3]*mat3in[2+0*3];
return ret;
}
//inverts a 3x3 (9 element) matrix
__host__ __device__ void mat3f_inverse(float *mat3inout)
{
int I;
float matint[9];
float det = mat3f_det(mat3inout);
for(I=0;I<9;I++)
{
matint[I] = mat3inout[I];
}
mat3inout[0+0*3] = (matint[1+1*3]*matint[2+2*3]-matint[1+2*3]*matint[2+1*3])/det;
mat3inout[0+1*3] = -(matint[1+0*3]*matint[2+2*3]-matint[1+2*3]*matint[2+0*3])/det;
mat3inout[0+2*3] = (matint[1+0*3]*matint[2+1*3]-matint[1+1*3]*matint[2+0*3])/det;
mat3inout[1+0*3] = -(matint[0+1*3]*matint[2+2*3]-matint[0+2*3]*matint[2+1*3])/det;
mat3inout[1+1*3] = (matint[0+0*3]*matint[2+2*3]-matint[0+2*3]*matint[2+0*3])/det;
mat3inout[1+2*3] = -(matint[0+0*3]*matint[2+1*3]-matint[0+1*3]*matint[2+0*3])/det;
mat3inout[2+0*3] = (matint[0+1*3]*matint[1+2*3]-matint[0+2*3]*matint[1+1*3])/det;
mat3inout[2+1*3] = -(matint[0+0*3]*matint[1+2*3]-matint[0+2*3]*matint[1+0*3])/det;
mat3inout[2+2*3] = (matint[0+0*3]*matint[1+1*3]-matint[0+1*3]*matint[1+0*3])/det;
mat3f_transpose(mat3inout);
return;
}
__host__ __device__ cuvect3f mat3f_mult(float *mat3in, const cuvect3f &cvin)
{
int I,J;
cuvect3f ret;
for(I=0;I<3;I++)
{
ret[I] = 0.0;
for(J=0;J<3;J++)
{
ret[I] = ret[I] + mat3in[I+3*J]*cvin[J];
}
}
return ret;
}
__host__ __device__ void mat3f_mult(float *matina, float *matinb, float *matout)
{
float wrk[9];
int I,J,K;
for(I=0;I<3;I++)
{
for(J=0;J<3;J++)
{
wrk[I+3*J] = 0.0;
}
}
for(I=0;I<3;I++)
{
for(J=0;J<3;J++)
{
for(K=0;K<3;K++)
{
wrk[I+3*K] = wrk[I+3*K] + matina[I+3*J]*matinb[J+3*K];
}
}
}
for(I=0;I<3;I++)
{
for(J=0;J<3;J++)
{
matout[I+3*J] = wrk[I+3*J];
}
}
return;
}
__host__ __device__ void mat3f_hodgedual(const cuvect3f &vecin, float *matout)
{
matout[0 + 0*3] = 0.0f;
matout[1 + 0*3] = -vecin[2];
matout[2 + 0*3] = vecin[1];
matout[0 + 1*3] = vecin[2];
matout[1 + 1*3] = 0.0f;
matout[2 + 1*3] = -vecin[0];
matout[0 + 2*3] = -vecin[1];
matout[1 + 2*3] = vecin[0];
matout[2 + 2*3] = 0.0f;
return;
}
__host__ __device__ void mat3f_hodgedual(float *matin, cuvect3f &vecout)
{
vecout[0] = 0.5*(matin[1 + 2*3] - matin[2 + 1*3]);
vecout[1] = 0.5*(matin[2 + 0*3] - matin[0 + 2*3]);
vecout[2] = 0.5*(matin[0 + 1*3] - matin[1 + 0*3]);
return;
}
//returns direction cosine rotation matrix from axis and angle
__host__ __device__ void mat3f_rot_from_axisangle(cuvect3f axis, float angle, float *matout)
{
int I;
float H[9];
float Hsq[9];
float II[9];
for(I=0;I<9;I++) II[I] = 0.0;
II[0+0*3] = 1.0;
II[1+1*3] = 1.0;
II[2+2*3] = 1.0;
axis = cuvect3f_normalize(axis);
mat3f_hodgedual(axis,H);
mat3f_mult(H,H,Hsq);
for(I=0;I<9;I++)
{
matout[I] = (II[I] + Hsq[I]) + H[I]*sinf(angle) - Hsq[I]*cosf(angle);
}
return;
}
__host__ void test_cudavectf_logic1()
{
//3 dim vector and matrix functional tests on host side

View File

@ -0,0 +1,185 @@
#include <amsculib2/amsculib2.hpp>
namespace amscuda
{
//transposes a 3x3 (9 element) matrix
__host__ __device__ void mat3f_transpose(float *mat3inout)
{
int I,J;
float matint[9];
for(I=0;I<9;I++)
{
matint[I] = mat3inout[I];
}
for(I=0;I<3;I++)
{
for(J=0;J<3;J++)
{
mat3inout[I+J*3] = matint[J+I*3];
}
}
return;
}
//copies src to dest
__host__ __device__ void mat3f_copy(float *mat3f_dest, const float *mat3f_src)
{
int I;
if(mat3f_dest==NULL || mat3f_src==NULL)
return;
for(I=0;I<9;I++)
mat3f_dest[I] = mat3f_src[I];
return;
}
__host__ __device__ float mat3f_det(float *mat3in)
{
float ret = 0.0f;
ret = ret + mat3in[0+0*3]*mat3in[1+1*3]*mat3in[2+2*3];
ret = ret + mat3in[0+1*3]*mat3in[1+2*3]*mat3in[2+0*3];
ret = ret + mat3in[0+2*3]*mat3in[1+0*3]*mat3in[2+1*3];
ret = ret - mat3in[0+0*3]*mat3in[1+2*3]*mat3in[2+1*3];
ret = ret - mat3in[0+1*3]*mat3in[1+0*3]*mat3in[2+2*3];
ret = ret - mat3in[0+2*3]*mat3in[1+1*3]*mat3in[2+0*3];
return ret;
}
//inverts a 3x3 (9 element) matrix
__host__ __device__ void mat3f_inverse(float *mat3inout)
{
int I;
float matint[9];
float det = mat3f_det(mat3inout);
for(I=0;I<9;I++)
{
matint[I] = mat3inout[I];
}
mat3inout[0+0*3] = (matint[1+1*3]*matint[2+2*3]-matint[1+2*3]*matint[2+1*3])/det;
mat3inout[0+1*3] = -(matint[1+0*3]*matint[2+2*3]-matint[1+2*3]*matint[2+0*3])/det;
mat3inout[0+2*3] = (matint[1+0*3]*matint[2+1*3]-matint[1+1*3]*matint[2+0*3])/det;
mat3inout[1+0*3] = -(matint[0+1*3]*matint[2+2*3]-matint[0+2*3]*matint[2+1*3])/det;
mat3inout[1+1*3] = (matint[0+0*3]*matint[2+2*3]-matint[0+2*3]*matint[2+0*3])/det;
mat3inout[1+2*3] = -(matint[0+0*3]*matint[2+1*3]-matint[0+1*3]*matint[2+0*3])/det;
mat3inout[2+0*3] = (matint[0+1*3]*matint[1+2*3]-matint[0+2*3]*matint[1+1*3])/det;
mat3inout[2+1*3] = -(matint[0+0*3]*matint[1+2*3]-matint[0+2*3]*matint[1+0*3])/det;
mat3inout[2+2*3] = (matint[0+0*3]*matint[1+1*3]-matint[0+1*3]*matint[1+0*3])/det;
mat3f_transpose(mat3inout);
return;
}
__host__ __device__ cuvect3f mat3f_mult(float *mat3in, const cuvect3f &cvin)
{
int I,J;
cuvect3f ret;
for(I=0;I<3;I++)
{
ret[I] = 0.0f;
for(J=0;J<3;J++)
{
ret[I] = ret[I] + mat3in[I+3*J]*cvin[J];
}
}
return ret;
}
__host__ __device__ void mat3f_mult(float *matina, float *matinb, float *matout)
{
float wrk[9];
int I,J,K;
for(I=0;I<3;I++)
{
for(J=0;J<3;J++)
{
wrk[I+3*J] = 0.0f;
}
}
for(I=0;I<3;I++)
{
for(J=0;J<3;J++)
{
for(K=0;K<3;K++)
{
wrk[I+3*K] = wrk[I+3*K] + matina[I+3*J]*matinb[J+3*K];
}
}
}
for(I=0;I<3;I++)
{
for(J=0;J<3;J++)
{
matout[I+3*J] = wrk[I+3*J];
}
}
return;
}
__host__ __device__ void mat3f_hodgedual(const cuvect3f &vecin, float *matout)
{
matout[0 + 0*3] = 0.0f;
matout[1 + 0*3] = -vecin[2];
matout[2 + 0*3] = vecin[1];
matout[0 + 1*3] = vecin[2];
matout[1 + 1*3] = 0.0f;
matout[2 + 1*3] = -vecin[0];
matout[0 + 2*3] = -vecin[1];
matout[1 + 2*3] = vecin[0];
matout[2 + 2*3] = 0.0f;
return;
}
__host__ __device__ void mat3f_hodgedual(float *matin, cuvect3f &vecout)
{
vecout[0] = 0.5*(matin[1 + 2*3] - matin[2 + 1*3]);
vecout[1] = 0.5*(matin[2 + 0*3] - matin[0 + 2*3]);
vecout[2] = 0.5*(matin[0 + 1*3] - matin[1 + 0*3]);
return;
}
//returns direction cosine rotation matrix from axis and angle
__host__ __device__ void mat3f_rot_from_axisangle(cuvect3f axis, float angle, float *matout)
{
int I;
float H[9];
float Hsq[9];
float II[9];
for(I=0;I<9;I++) II[I] = 0.0f;
II[0+0*3] = 1.0;
II[1+1*3] = 1.0;
II[2+2*3] = 1.0;
axis = cuvect3f_normalize(axis);
mat3f_hodgedual(axis,H);
mat3f_mult(H,H,Hsq);
for(I=0;I<9;I++)
{
matout[I] = (II[I] + Hsq[I]) + H[I]*sinf(angle) - Hsq[I]*cosf(angle);
}
return;
}
};

View File

@ -674,30 +674,30 @@ __host__ __device__ cumat4f& cumat4f::operator*=(const cumat4f &rhs)
}
__host__ __device__ float cuvect4f_dot(cuvect4f a, cuvect4f b)
__host__ __device__ float cuvect4f_dot(cuvect4f &a, cuvect4f &b)
{
float ret = 0.0;
float ret = 0.0f;
ret = a.x*b.x + a.y*b.y + a.z*b.z + a.w*b.w;
return ret;
}
__host__ __device__ float cuvect4f_norm(cuvect4f a)
__host__ __device__ float cuvect4f_norm(cuvect4f &a)
{
float ret = 0.0;
float ret = 0.0f;
ret = ::sqrtf(cuvect4f_dot(a,a));
return ret;
}
__host__ __device__ cuvect4f cuvect4f_normalize(cuvect4f a)
__host__ __device__ cuvect4f cuvect4f_normalize(cuvect4f &a)
{
cuvect4f ret = cuvect4f(0.0f,0.0f,0.0f,0.0f);
float nrm = cuvect4f_norm(a);
if(nrm>0.0)
if(nrm>0.0f)
ret = a/nrm;
return ret;
}
__host__ __device__ cuvect4f cuvect4f_proj(cuvect4f a, cuvect4f b)
__host__ __device__ cuvect4f cuvect4f_proj(cuvect4f &a, cuvect4f &b)
{
cuvect4f ret;
cuvect4f bn = cuvect4f_normalize(b);