pcluster
new_cluster(uint size, uint * idx, uint sons, uint dim)
{
pcluster t;
uint i;
t = (pcluster) allocmem((size_t) sizeof(cluster));
t->type = 0;
t->size = size;
t->dim = dim;
t->idx = idx;
t->bmin = allocreal(dim);
t->bmax = allocreal(dim);
t->sons = sons;
if (sons > 0) {
t->son = (pcluster *) allocmem((size_t) sons * sizeof(pcluster));
for (i = 0; i < sons; i++)
t->son[i] = 0;
}
else
t->son = 0;
return t;
}
real
max_rel_inner_error(pcbem3d bem, helmholtz_data * hdata, pcavector x,
boundary_func3d rhs)
{
uint nx, nz, npoints;
real(*xdata)[3];
field *ydata;
uint i, j;
real error, maxerror;
real eta_bw =
REAL_SQRT(ABSSQR(bem->kvec[0]) + ABSSQR(bem->kvec[1]) +
ABSSQR(bem->kvec[2]));
nx = 20;
nz = 20;
npoints = nx * nz;
xdata = (real(*)[3]) allocreal(3 * npoints);
npoints = 0;
for (j = 0; j < nz; ++j) {
for (i = 0; i < nx; ++i) {
xdata[npoints][0] = -1.0 + (2.0 / (nx - 1)) * i;
xdata[npoints][1] = 0.0;
xdata[npoints][2] = -1.0 + (2.0 / (nz - 1)) * j;
if (REAL_SQR(xdata[npoints][0]) + REAL_SQR(xdata[npoints][2]) < 1) {
npoints++;
}
}
}
ydata = allocfield(npoints);
#pragma omp parallel for
for (j = 0; j < npoints; ++j) {
ydata[j] = eval_brakhage_werner_c(bem, x, eta_bw, xdata[j]);
}
j = 0;
maxerror =
ABS(ydata[j] -
rhs(xdata[j], NULL, hdata)) / ABS(rhs(xdata[j], NULL, hdata));
for (j = 1; j < npoints; ++j) {
error =
ABS(ydata[j] -
rhs(xdata[j], NULL, hdata)) / ABS(rhs(xdata[j], NULL, hdata));
maxerror = error > maxerror ? error : maxerror;
}
freemem(ydata);
freemem(xdata);
return maxerror;
}
pleveldir
remove_unused_direction(pdblock b, pdcluster t, pleveldir lold)
{
uint i, j;
uint entrys, dirs;
preal dirmem;
pleveldir ldir;
if (t->directions == 0) {
return lold;
}
ldir = new_leveldir(getdepth_dcluster(t), t->dim);
uint depth = getdepth_dblock(b);
uint **idx = (uint **) allocmem(sizeof(uint *) * (depth + 1));
uint **used = (uint **) allocmem(sizeof(uint *) * (depth + 1));
uint **tmp_dirson = (uint **) allocmem(sizeof(uint *) * (depth));
uint *fill = (uint *) allocmem(sizeof(uint) * (depth + 1));
/* Copy Information that won't change */
for (i = 0; i < ldir->depth + 1; i++) {
ldir->maxdiam[i] = lold->maxdiam[i];
ldir->splits[i] = lold->splits[i];
}
/* Set up used array */
for (i = 0; i < depth; i++) {
used[i] = (uint *) allocmem(sizeof(uint) * lold->directions[i]);
tmp_dirson[i] = (uint *) allocmem(sizeof(uint) * lold->directions[i]);
fill[i] = 0;
for (j = 0; j < lold->directions[i]; j++) {
used[i][j] = 0;
}
}
used[depth] = (uint *) allocmem(sizeof(uint) * lold->directions[depth]);
fill[depth] = 0;
for (j = 0; j < lold->directions[depth]; j++) {
used[depth][j] = 0;
}
/* Scan directional block tree for used directions */
scan_tree(b, used, tmp_dirson, fill, 0);
freemem(fill);
/* Check if all needed directions are inside and copy important directions */
check_directions(used, tmp_dirson, lold);
/* Set up memory for new leveldir object and array for new numeration */
dirs = 0;
for (i = 0; i < depth + 1; i++) {
entrys = 0;
for (j = 0; j < lold->directions[i]; j++) { /* Count new directions */
if (used[i][j] > 0) {
entrys++;
}
}
idx[i] = (uint *) allocmem(sizeof(uint) * lold->directions[i]);
ldir->directions[i] = entrys;
ldir->dir[i] = (preal *) allocmem(sizeof(preal) * entrys);
dirs += entrys;
}
ldir->dirmem = dirmem = allocreal(dirs * ldir->dim);
for (i = 0; i < depth + 1; i++) {
for (j = 0; j < ldir->directions[i]; j++) {
ldir->dir[i][j] = dirmem;
dirmem += ldir->dim;
}
}
/* Fill new leveldir object */
copy_direction(used, idx, ldir, lold);
/* Last part is changing directional cluster and block */
/* Since we only have one directional cluster for both rows and columns we only need to call it once */
change_dcluster(t, used, idx, tmp_dirson, ldir, 0);
for (i = 0; i < depth; i++) {
freemem(tmp_dirson[i]);
freemem(used[i]);
}
freemem(used[depth]);
free(tmp_dirson);
free(used);
/* Now the block */
change_dblock(b, idx, 0);
for (i = 0; i < depth + 1; i++) {
freemem(idx[i]);
}
freemem(idx);
del_leveldir(lold);
return ldir;
}
static void
test_h2matrix_system(const char *apprxtype, pcamatrix Vfull,
pcamatrix KMfull, pblock block, pbem3d bem_slp,
ph2matrix V, pbem3d bem_dlp, ph2matrix KM, bool linear,
bool exterior, real low, real high)
{
struct _eval_A eval;
helmholtz_data hdata;
pavector x, b;
real errorV, errorKM, error_solve, eps_solve;
uint steps;
boundary_func3d rhs = (boundary_func3d) rhs_dirichlet_point_helmholtzbem3d;
eps_solve = 1.0e-12;
steps = 1000;
printf("Testing: %s H2matrix %s\n"
"====================================\n\n",
(exterior == true ? "exterior" : "interior"), apprxtype);
assemble_bem3d_h2matrix_row_clusterbasis(bem_slp, V->rb);
assemble_bem3d_h2matrix_col_clusterbasis(bem_slp, V->cb);
SCHEDULE_OPENCL(0, 1, assemble_bem3d_h2matrix, bem_slp, block, V);
assemble_bem3d_h2matrix_row_clusterbasis(bem_dlp, KM->rb);
assemble_bem3d_h2matrix_col_clusterbasis(bem_dlp, KM->cb);
SCHEDULE_OPENCL(0, 1, assemble_bem3d_h2matrix, bem_dlp, block, KM);
eval.V = V;
eval.Vtype = H2MATRIX;
eval.KM = KM;
eval.KMtype = H2MATRIX;
eval.eta =
REAL_SQRT(ABSSQR(bem_slp->kvec[0]) + ABSSQR(bem_slp->kvec[1]) +
ABSSQR(bem_slp->kvec[2]));
hdata.kvec = bem_slp->kvec;
hdata.source = allocreal(3);
if (exterior) {
hdata.source[0] = 0.0, hdata.source[1] = 0.0, hdata.source[2] = 0.2;
}
else {
hdata.source[0] = 0.0, hdata.source[1] = 0.0, hdata.source[2] = 5.0;
}
errorV = norm2diff_amatrix_h2matrix(V, Vfull) / norm2_amatrix(Vfull);
printf("rel. error V : %.5e\n", errorV);
errorKM = norm2diff_amatrix_h2matrix(KM, KMfull) / norm2_amatrix(KMfull);
printf("rel. error K%c0.5*M : %.5e\n", (exterior == true ? '-' : '+'),
errorKM);
x = new_avector(Vfull->rows);
b = new_avector(KMfull->cols);
printf("Solving Dirichlet problem:\n");
integrate_bem3d_const_avector(bem_dlp, rhs, b, (void *) &hdata);
solve_gmres_bem3d(HMATRIX, &eval, b, x, eps_solve, steps);
error_solve = max_rel_outer_error(bem_slp, &hdata, x, rhs);
printf("max. rel. error : %.5e %s\n", error_solve,
(IS_IN_RANGE(low, error_solve, high) ? " okay" : "NOT okay"));
if (!IS_IN_RANGE(low, error_solve, high))
problems++;
printf("\n");
del_avector(x);
del_avector(b);
freemem(hdata.source);
}
static void
init_helmholtzbem3d_opencl(pcbem3d bem)
{
uint num_kernels;
const char *kernel_names[] = {
"assemble_slp_cc_list_0", "assemble_slp_cc_list_1",
"assemble_slp_cc_list_2", "assemble_slp_cc_list_3",
"assemble_dlp_cc_list_0", "assemble_dlp_cc_list_1",
"assemble_dlp_cc_list_2", "assemble_dlp_cc_list_3"
};
cl_uint num_devices = ocl_system.num_devices;
cl_uint num_queues = ocl_system.queues_per_device;
cl_uint nthreads = num_devices * num_queues;
cl_int res;
cl_mem_flags mem_rflags, mem_wflags;
uint i, j;
real *gr_x;
uint *gr_t;
real *gr_n;
real *q_xw, *x, *w;
uint q, v, t;
num_kernels = sizeof(kernel_names) / sizeof(kernel_names[0]);
mem_rflags = CL_MEM_READ_ONLY;
mem_wflags = CL_MEM_WRITE_ONLY;
if (ocl_bem3d.num_kernels == 0) {
/****************************************************
* Setup all necessary kernels
****************************************************/
setup_kernels(helmholtzbem3d_ocl_src, num_kernels, kernel_names,
&ocl_bem3d.kernels);
ocl_bem3d.num_kernels = num_kernels;
/****************************************************
* Create buffers for matrix chunks
****************************************************/
ocl_bem3d.mem_N = (cl_mem *) allocmem(nthreads * sizeof(cl_mem));
ocl_bem3d.mem_ridx = (cl_mem *) allocmem(nthreads * sizeof(cl_mem));
ocl_bem3d.mem_cidx = (cl_mem *) allocmem(nthreads * sizeof(cl_mem));
for (i = 0; i < num_devices; ++i) {
for (j = 0; j < num_queues; ++j) {
ocl_bem3d.mem_N[j + i * num_queues] =
clCreateBuffer(ocl_system.contexts[i], mem_wflags,
ocl_system.max_package_size, NULL, &res);
CL_CHECK(res)
ocl_bem3d.mem_ridx[j + i * num_queues] =
clCreateBuffer(ocl_system.contexts[i], mem_rflags,
ocl_system.max_package_size / sizeof(field) *
sizeof(uint), NULL, &res);
CL_CHECK(res)
ocl_bem3d.mem_cidx[j + i * num_queues] =
clCreateBuffer(ocl_system.contexts[i], mem_rflags,
ocl_system.max_package_size / sizeof(field) *
sizeof(uint), NULL, &res);
CL_CHECK(res)
}
}
/****************************************************
* Create buffer for non singular quadrature rules
* and copy the contents
****************************************************/
q = bem->sq->q;
x = allocreal(q);
w = allocreal(q);
q_xw = allocreal(2 * q);
assemble_gauss(q, x, w);
for (i = 0; i < q; ++i) {
q_xw[2 * i] = 0.5 + 0.5 * x[i];
q_xw[2 * i + 1] = 0.5 * w[i];
}
ocl_bem3d.mem_q_xw = (cl_mem *) allocmem(num_devices * sizeof(cl_mem));
for (i = 0; i < num_devices; ++i) {
ocl_bem3d.mem_q_xw[i] =
clCreateBuffer(ocl_system.contexts[i], mem_rflags,
2 * q * sizeof(real), NULL, &res);
CL_CHECK(res)
res = clEnqueueWriteBuffer(ocl_system.queues[i * num_queues],
ocl_bem3d.mem_q_xw[i], CL_TRUE, 0,
2 * q * sizeof(real), q_xw, 0, NULL, NULL);
CL_CHECK(res);
}
ocl_bem3d.nq = 2 * q;
freemem(x);
freemem(w);
freemem(q_xw);
/****************************************************
* Create buffer for singular quadrature rules
* and copy the contents
****************************************************/
q = bem->sq->q2;
x = allocreal(q);
w = allocreal(q);
q_xw = allocreal(2 * q);
assemble_gauss(q, x, w);
for (i = 0; i < q; ++i) {
q_xw[2 * i] = 0.5 + 0.5 * x[i];
q_xw[2 * i + 1] = 0.5 * w[i];
}
ocl_bem3d.mem_q2_xw = (cl_mem *) allocmem(num_devices * sizeof(cl_mem));
for (i = 0; i < num_devices; ++i) {
ocl_bem3d.mem_q2_xw[i] = clCreateBuffer(ocl_system.contexts[i],
mem_rflags,
2 * q * sizeof(real), NULL,
&res);
CL_CHECK(res)
res = clEnqueueWriteBuffer(ocl_system.queues[i * num_queues],
ocl_bem3d.mem_q2_xw[i], CL_TRUE, 0,
2 * q * sizeof(real), q_xw, 0, NULL, NULL);
CL_CHECK(res);
}
ocl_bem3d.nq2 = 2 * q;
freemem(x);
freemem(w);
freemem(q_xw);
/****************************************************
* Create buffers for geometry data and copy the contents
****************************************************/
v = bem->gr->vertices;
t = bem->gr->triangles;
gr_x = allocreal(3 * v);
gr_t = allocuint(3 * t);
gr_n = allocreal(3 * t);
for (i = 0; i < v; ++i) {
gr_x[3 * i + 0] = bem->gr->x[i][0];
gr_x[3 * i + 1] = bem->gr->x[i][1];
gr_x[3 * i + 2] = bem->gr->x[i][2];
}
for (i = 0; i < t; ++i) {
gr_t[i + 0 * t] = bem->gr->t[i][0];
gr_t[i + 1 * t] = bem->gr->t[i][1];
gr_t[i + 2 * t] = bem->gr->t[i][2];
}
for (i = 0; i < t; ++i) {
gr_n[3 * i + 0] = bem->gr->n[i][0];
gr_n[3 * i + 1] = bem->gr->n[i][1];
gr_n[3 * i + 2] = bem->gr->n[i][2];
}
ocl_bem3d.mem_gr_t = (cl_mem *) allocmem(num_devices * sizeof(cl_mem));
ocl_bem3d.mem_gr_x = (cl_mem *) allocmem(num_devices * sizeof(cl_mem));
for (i = 0; i < num_devices; ++i) {
ocl_bem3d.mem_gr_x[i] =
clCreateBuffer(ocl_system.contexts[i], mem_rflags,
3 * v * sizeof(real), NULL, &res);
CL_CHECK(res)
ocl_bem3d.mem_gr_t[i] =
clCreateBuffer(ocl_system.contexts[i], mem_rflags,
3 * t * sizeof(uint), NULL, &res);
CL_CHECK(res)
res = clEnqueueWriteBuffer(ocl_system.queues[i * num_queues],
ocl_bem3d.mem_gr_x[i], CL_TRUE, 0,
3 * v * sizeof(real), gr_x, 0, NULL, NULL);
CL_CHECK(res);
res = clEnqueueWriteBuffer(ocl_system.queues[i * num_queues],
ocl_bem3d.mem_gr_t[i], CL_TRUE, 0,
3 * t * sizeof(uint), gr_t, 0, NULL, NULL);
CL_CHECK(res);
}
ocl_bem3d.triangles = t;
freemem(gr_x);
freemem(gr_t);
freemem(gr_n);
omp_init_lock(&nf_lock);
omp_init_lock(&nf_dist_lock);
omp_init_lock(&nf_vert_lock);
omp_init_lock(&nf_edge_lock);
omp_init_lock(&nf_iden_lock);
}
pcluster
build_pca_cluster(pclustergeometry cf, uint size, uint * idx, uint clf)
{
const uint dim = cf->dim;
pamatrix C, Q;
pavector v;
prealavector lambda;
real *x, *y;
real w;
uint i, j, k, size0, size1;
pcluster t;
assert(size > 0);
size0 = 0;
size1 = 0;
if (size > clf) {
x = allocreal(dim);
y = allocreal(dim);
/* determine weight of current cluster */
w = 0.0;
for (i = 0; i < size; ++i) {
w += cf->w[idx[i]];
}
w = 1.0 / w;
for (j = 0; j < dim; ++j) {
x[j] = 0.0;
}
/* determine center of mass */
for (i = 0; i < size; ++i) {
for (j = 0; j < dim; ++j) {
x[j] += cf->w[idx[i]] * cf->x[idx[i]][j];
}
}
for (j = 0; j < dim; ++j) {
x[j] *= w;
}
C = new_zero_amatrix(dim, dim);
Q = new_zero_amatrix(dim, dim);
lambda = new_realavector(dim);
/* setup covariance matrix */
for (i = 0; i < size; ++i) {
for (j = 0; j < dim; ++j) {
y[j] = cf->x[idx[i]][j] - x[j];
}
for (j = 0; j < dim; ++j) {
for (k = 0; k < dim; ++k) {
C->a[j + k * C->ld] += cf->w[idx[i]] * y[j] * y[k];
}
}
}
/* get eigenvalues and eigenvectors of covariance matrix */
eig_amatrix(C, lambda, Q);
/* get eigenvector from largest eigenvalue */
v = new_avector(0);
init_column_avector(v, Q, dim - 1);
/* separate cluster with v as separation-plane */
for (i = 0; i < size; ++i) {
/* x_i - X */
for (j = 0; j < dim; ++j) {
y[j] = cf->x[idx[i]][j] - x[j];
}
/* <y,v> */
w = 0.0;
for (j = 0; j < dim; ++j) {
w += y[j] * v->v[j];
}
if (w >= 0.0) {
j = idx[i];
idx[i] = idx[size0];
idx[size0] = j;
size0++;
}
else {
size1++;
}
}
assert(size0 + size1 == size);
del_amatrix(Q);
del_amatrix(C);
del_realavector(lambda);
del_avector(v);
freemem(x);
freemem(y);
/* recursion */
if (size0 > 0) {
if (size1 > 0) {
t = new_cluster(size, idx, 2, cf->dim);
t->son[0] = build_pca_cluster(cf, size0, idx, clf);
t->son[1] = build_pca_cluster(cf, size1, idx + size0, clf);
update_bbox_cluster(t);
}
else {
t = new_cluster(size, idx, 1, cf->dim);
t->son[0] = build_pca_cluster(cf, size0, idx, clf);
update_bbox_cluster(t);
}
}
else {
assert(size1 > 0);
t = new_cluster(size, idx, 1, cf->dim);
t->son[0] = build_pca_cluster(cf, size1, idx, clf);
update_bbox_cluster(t);
}
}
else {
t = new_cluster(size, idx, 0, cf->dim);
update_support_bbox_cluster(cf, t);
}
update_cluster(t);
return t;
}
pbem3d
new_dlp_helmholtz_bem3d(field * kvec, pcsurface3d gr, uint q_regular,
uint q_singular, basisfunctionbem3d basis_neumann,
basisfunctionbem3d basis_dirichlet, field alpha)
{
pkernelbem3d kernels;
pbem3d bem;
bem = new_bem3d(gr);
kernels = bem->kernels;
bem->sq = build_singquad2d(gr, q_regular, q_singular);
if (basis_neumann == BASIS_LINEAR_BEM3D || basis_dirichlet
== BASIS_LINEAR_BEM3D) {
setup_vertex_to_triangle_map_bem3d(bem);
}
bem->basis_neumann = basis_neumann;
bem->basis_dirichlet = basis_dirichlet;
bem->kernel_const = KERNEL_CONST_HELMHOLTZBEM3D;
bem->kvec = kvec;
bem->k = NORM3(kvec[0], kvec[1], kvec[2]);
bem->alpha = alpha;
kernels->fundamental = fill_kernel_helmholtzbem3d;
kernels->dny_fundamental = fill_dny_kernel_helmholtzbem3d;
kernels->dnx_dny_fundamental = fill_dnx_dny_kernel_helmholtzbem3d;
if (basis_neumann == BASIS_CONSTANT_BEM3D) {
bem->N_neumann = gr->triangles;
}
else {
assert(basis_neumann == BASIS_LINEAR_BEM3D);
bem->N_neumann = gr->vertices;
}
if (basis_dirichlet == BASIS_CONSTANT_BEM3D) {
bem->N_dirichlet = gr->triangles;
}
else {
assert(basis_dirichlet == BASIS_LINEAR_BEM3D);
bem->N_dirichlet = gr->vertices;
}
if (basis_neumann == BASIS_CONSTANT_BEM3D && basis_dirichlet
== BASIS_CONSTANT_BEM3D) {
bem->nearfield = fill_dlp_cc_near_helmholtzbem3d;
bem->nearfield_far = fill_dlp_cc_far_helmholtzbem3d;
kernels->lagrange_row = assemble_bem3d_lagrange_const_amatrix;
kernels->lagrange_col = assemble_bem3d_dn_lagrange_const_amatrix;
kernels->fundamental_row = fill_kernel_c_helmholtzbem3d;
kernels->fundamental_col = fill_kernel_c_helmholtzbem3d;
kernels->dnz_fundamental_row = fill_dnz_kernel_c_helmholtzbem3d;
kernels->dnz_fundamental_col = fill_dnz_kernel_c_helmholtzbem3d;
kernels->kernel_row = fill_kernel_c_helmholtzbem3d;
kernels->kernel_col = fill_dcol_kernel_col_c_helmholtzbem3d;
kernels->dnz_kernel_row = fill_dnz_kernel_c_helmholtzbem3d;
kernels->dnz_kernel_col = fill_dnzdcol_kernel_c_helmholtzbem3d;
}
else if (basis_neumann == BASIS_LINEAR_BEM3D
&& basis_dirichlet == BASIS_CONSTANT_BEM3D) {
bem->mass = allocreal(3);
/* TODO MASS-MATRIX */
}
else if (basis_neumann == BASIS_CONSTANT_BEM3D
&& basis_dirichlet == BASIS_LINEAR_BEM3D) {
weight_basisfunc_cl_singquad2d(bem->sq->x_id, bem->sq->y_id,
bem->sq->w_id, bem->sq->n_id);
weight_basisfunc_cl_singquad2d(bem->sq->x_edge, bem->sq->y_edge,
bem->sq->w_edge, bem->sq->n_edge);
weight_basisfunc_cl_singquad2d(bem->sq->x_vert, bem->sq->y_vert,
bem->sq->w_vert, bem->sq->n_vert);
weight_basisfunc_cl_singquad2d(bem->sq->x_dist, bem->sq->y_dist,
bem->sq->w_dist, bem->sq->n_dist);
weight_basisfunc_l_singquad2d(bem->sq->x_single, bem->sq->y_single,
bem->sq->w_single, bem->sq->n_single);
bem->nearfield = fill_dlp_cl_near_helmholtzbem3d;
bem->nearfield_far = fill_dlp_cl_far_helmholtzbem3d;
kernels->lagrange_row = assemble_bem3d_lagrange_const_amatrix;
kernels->lagrange_col = assemble_bem3d_dn_lagrange_linear_amatrix;
kernels->fundamental_row = fill_kernel_c_helmholtzbem3d;
kernels->fundamental_col = fill_kernel_l_helmholtzbem3d;
kernels->dnz_fundamental_row = fill_dnz_kernel_c_helmholtzbem3d;
kernels->dnz_fundamental_col = fill_dnz_kernel_l_helmholtzbem3d;
kernels->kernel_row = fill_kernel_c_helmholtzbem3d;
kernels->kernel_col = fill_dcol_kernel_col_l_helmholtzbem3d;
kernels->dnz_kernel_row = fill_dnz_kernel_c_helmholtzbem3d;
kernels->dnz_kernel_col = fill_dnzdcol_kernel_l_helmholtzbem3d;
bem->mass = allocreal(3);
bem->mass[0] = 1.0 / 6.0;
bem->mass[1] = 1.0 / 6.0;
bem->mass[2] = 1.0 / 6.0;
}
else {
assert(basis_neumann == BASIS_LINEAR_BEM3D && basis_dirichlet
== BASIS_LINEAR_BEM3D);
weight_basisfunc_ll_singquad2d(bem->sq->x_id, bem->sq->y_id,
bem->sq->w_id, bem->sq->n_id);
weight_basisfunc_ll_singquad2d(bem->sq->x_edge, bem->sq->y_edge,
bem->sq->w_edge, bem->sq->n_edge);
weight_basisfunc_ll_singquad2d(bem->sq->x_vert, bem->sq->y_vert,
bem->sq->w_vert, bem->sq->n_vert);
weight_basisfunc_ll_singquad2d(bem->sq->x_dist, bem->sq->y_dist,
bem->sq->w_dist, bem->sq->n_dist);
weight_basisfunc_l_singquad2d(bem->sq->x_single, bem->sq->y_single,
bem->sq->w_single, bem->sq->n_single);
bem->nearfield = fill_dlp_ll_near_helmholtzbem3d;
bem->nearfield_far = fill_dlp_ll_far_helmholtzbem3d;
kernels->lagrange_row = assemble_bem3d_lagrange_linear_amatrix;
kernels->lagrange_col = assemble_bem3d_dn_lagrange_linear_amatrix;
kernels->fundamental_row = fill_kernel_l_helmholtzbem3d;
kernels->fundamental_col = fill_kernel_l_helmholtzbem3d;
kernels->dnz_fundamental_row = fill_dnz_kernel_l_helmholtzbem3d;
kernels->dnz_fundamental_col = fill_dnz_kernel_l_helmholtzbem3d;
kernels->kernel_row = fill_kernel_l_helmholtzbem3d;
kernels->kernel_col = fill_dcol_kernel_col_l_helmholtzbem3d;
kernels->dnz_kernel_row = NULL;
kernels->dnz_kernel_col = fill_dnzdcol_kernel_l_helmholtzbem3d;
bem->mass = allocreal(9);
bem->mass[0] = 1.0 / 12.0;
bem->mass[1] = 1.0 / 24.0;
bem->mass[2] = 1.0 / 24.0;
bem->mass[3] = 1.0 / 24.0;
bem->mass[4] = 1.0 / 12.0;
bem->mass[5] = 1.0 / 24.0;
bem->mass[6] = 1.0 / 24.0;
bem->mass[7] = 1.0 / 24.0;
bem->mass[8] = 1.0 / 12.0;
}
return bem;
}
