shuren/opencl_fdfd
1
0
Fork 0
forked from jan/opencl_fdfd
Code Issues Pull Requests Releases Wiki Activity

Conductor partial fixup

Browse Source
This commit is contained in:
jan 2016-07-04 16:29:21 -07:00
parent 985fca76f4
commit fdcf090249
5 changed files with 26 additions and 22 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

8
opencl_fdfd/kernels/e2h.cl
View File

@ -35,7 +35,7 @@ if ( z == sz - 1 ) {
//Update H components; set them to 0 if PMC is enabled there.
// Also divide by mu only if requested.
{% if pmc -%}
if (pmc[XX + i]) {
if (pmc[XX + i] != 0) {
Hx[i] = cdouble_new(0.0, 0.0);
} else
{%- endif -%}
@ -51,7 +51,7 @@ if (pmc[XX + i]) {
}
{% if pmc -%}
if (pmc[YY + i]) {
if (pmc[YY + i] != 0) {
Hy[i] = cdouble_new(0.0, 0.0);
} else
{%- endif -%}
@ -67,8 +67,8 @@ if (pmc[YY + i]) {
}
{% if pmc -%}
if (pmc[XX + i]) {
Hx[i] = cdouble_new(0.0, 0.0);
if (pmc[ZZ + i] != 0) {
Hz[i] = cdouble_new(0.0, 0.0);
} else
{%- endif -%}
{

12
opencl_fdfd/kernels/h2e.cl
View File

@ -34,9 +34,7 @@ if ( z == 0 ) {
//Update E components; set them to 0 if PEC is enabled there.
{% if pec -%}
if (pec[XX + i]) {
Ex[i] = cdouble_new(0.0, 0.0);
} else
if (pec[XX + i] == 0)
{%- endif -%}
{
cdouble_t tEx = cdouble_mul(Ex[i], oeps[XX + i]);
@ -47,9 +45,7 @@ if (pec[XX + i]) {
}
{% if pec -%}
if (pec[YY + i]) {
Ey[i] = cdouble_new(0.0, 0.0);
} else
if (pec[YY + i] == 0)
{%- endif -%}
{
cdouble_t tEy = cdouble_mul(Ey[i], oeps[YY + i]);
@ -60,9 +56,7 @@ if (pec[YY + i]) {
}
{% if pec -%}
if (pec[ZZ + i]) {
Ez[i] = cdouble_new(0.0, 0.0);
} else
if (pec[ZZ + i] == 0)
{%- endif -%}
{
cdouble_t tEz = cdouble_mul(Ez[i], oeps[ZZ + i]);

9
opencl_fdfd/kernels/p2e.cl Normal file
View File

@ -0,0 +1,9 @@
{%- if pec -%}
if (pec[i] != 0) {
E[i] = cdouble_new(0.0, 0.0);
} else
{%- endif -%}
{
E[i] = cdouble_mul(Pr[i], p[i]);
}

4
opencl_fdfd/main.py
View File

@ -91,12 +91,12 @@ def cg_solver(omega, dxes, J, epsilon, mu=None, pec=None, pmc=None, adjoint=Fals
if pec is None:
gpec = load_field(numpy.array([]), dtype=numpy.int8)
else:
gpec = load_field(pec, dtype=numpy.int8)
gpec = load_field(pec.astype(bool), dtype=numpy.int8)
if pmc is None:
gpmc = load_field(numpy.array([]), dtype=numpy.int8)
else:
gpmc = load_field(pmc, dtype=numpy.int8)
gpmc = load_field(pmc.astype(bool), dtype=numpy.int8)
'''
Generate OpenCL kernels

15
opencl_fdfd/ops.py
View File

@ -101,15 +101,18 @@ def create_a(context, shape, mu=False, pec=False, pmc=False):
header = shape_source(shape) + dixyz_source + xyz_source
vec_h = vec_source + E_ptrs + H_ptrs
p2e_source = 'E[i] = cdouble_mul(Pr[i], p[i]);'
pec_arg = ['char *pec']
pmc_arg = ['char *pmc']
des = [ctype + ' *inv_de' + a for a in 'xyz']
dhs = [ctype + ' *inv_dh' + a for a in 'xyz']
p2e_source = jinja_env.get_template('p2e.cl').render(pec=pec)
P2E_kernel = ElementwiseKernel(context,
name='P2E',
preamble=preamble,
operation=p2e_source,
arguments=', '.join(ptrs('E', 'p', 'Pr')))
arguments=', '.join(ptrs('E', 'p', 'Pr') + pec_arg))
pmc_arg = ['char *pmc']
des = [ctype + ' *inv_de' + a for a in 'xyz']
e2h_source = jinja_env.get_template('e2h.cl').render(mu=mu,
pmc=pmc,
dixyz_source=header,
@ -120,8 +123,6 @@ def create_a(context, shape, mu=False, pec=False, pmc=False):
operation=e2h_source,
arguments=', '.join(ptrs('E', 'H', 'inv_mu') + pmc_arg + des))
pec_arg = ['char *pec']
dhs = [ctype + ' *inv_dh' + a for a in 'xyz']
h2e_source = jinja_env.get_template('h2e.cl').render(pmc=pec,
dixyz_source=header,
vec_source=vec_h)
@ -132,7 +133,7 @@ def create_a(context, shape, mu=False, pec=False, pmc=False):
arguments=', '.join(ptrs('E', 'H', 'oeps', 'Pl') + pec_arg + dhs))
def spmv(E, H, p, idxes, oeps, inv_mu, pec, pmc, Pl, Pr, e):
e2 = P2E_kernel(E, p, Pr, wait_for=e)
e2 = P2E_kernel(E, p, Pr, pec, wait_for=e)
e2 = E2H_kernel(E, H, inv_mu, pmc, *idxes[0], wait_for=[e2])
e2 = H2E_kernel(E, H, oeps, Pl, pec, *idxes[1], wait_for=[e2])
return [e2]