From fdcf090249568db90d59df09272927808945945b Mon Sep 17 00:00:00 2001
From: jan <jan@orange.mpxd.net>
Date: Mon, 4 Jul 2016 16:29:21 -0700
Subject: [PATCH] Conductor partial fixup

---
 opencl_fdfd/kernels/e2h.cl |  8 ++++----
 opencl_fdfd/kernels/h2e.cl | 12 +++---------
 opencl_fdfd/kernels/p2e.cl |  9 +++++++++
 opencl_fdfd/main.py        |  4 ++--
 opencl_fdfd/ops.py         | 15 ++++++++-------
 5 files changed, 26 insertions(+), 22 deletions(-)
 create mode 100644 opencl_fdfd/kernels/p2e.cl

diff --git a/opencl_fdfd/kernels/e2h.cl b/opencl_fdfd/kernels/e2h.cl
index 6f4f9fa..7691888 100644
--- a/opencl_fdfd/kernels/e2h.cl
+++ b/opencl_fdfd/kernels/e2h.cl
@@ -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 -%}
 {
diff --git a/opencl_fdfd/kernels/h2e.cl b/opencl_fdfd/kernels/h2e.cl
index d4fda62..b496774 100644
--- a/opencl_fdfd/kernels/h2e.cl
+++ b/opencl_fdfd/kernels/h2e.cl
@@ -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]);
diff --git a/opencl_fdfd/kernels/p2e.cl b/opencl_fdfd/kernels/p2e.cl
new file mode 100644
index 0000000..432eba9
--- /dev/null
+++ b/opencl_fdfd/kernels/p2e.cl
@@ -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]);
+}
diff --git a/opencl_fdfd/main.py b/opencl_fdfd/main.py
index f6a28d8..e41f6fa 100644
--- a/opencl_fdfd/main.py
+++ b/opencl_fdfd/main.py
@@ -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
diff --git a/opencl_fdfd/ops.py b/opencl_fdfd/ops.py
index bb686d4..8efa750 100644
--- a/opencl_fdfd/ops.py
+++ b/opencl_fdfd/ops.py
@@ -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]