Index: sasmodels/kernel_iq.c
===================================================================
--- sasmodels/kernel_iq.c	(revision 705307783b90bea1a66cdf1ff38974433eb673cb)
+++ sasmodels/kernel_iq.c	(revision e44432d7890080c5b9a440696903718a52c60eba)
@@ -26,5 +26,5 @@
 //  MAGNETIC_PARS : a comma-separated list of indices to the sld
 //      parameters in the parameter table.
-//  CALL_VOLUME(table) : call the form volume function
+//  CALL_VOLUME(form, shell, table) : assign form and shell values
 //  CALL_EFFECTIVE_RADIUS(type, table) : call the R_eff function
 //  CALL_IQ(q, table) : call the Iq function for 1D calcs.
@@ -275,5 +275,4 @@
 
 // ==================== KERNEL CODE ========================
-#define COMPUTE_F1_F2 defined(CALL_FQ)
 kernel
 void KERNEL_NAME(
@@ -345,25 +344,29 @@
   // version must loop over all q.
   #ifdef USE_OPENCL
-    #if COMPUTE_F1_F2
+    #if defined(CALL_FQ)
       double weight_norm = (pd_start == 0 ? 0.0 : result[2*nq]);
-      double weighted_volume = (pd_start == 0 ? 0.0 : result[2*nq+1]);
-      double weighted_radius = (pd_start == 0 ? 0.0 : result[2*nq+2]);
+      double weighted_form = (pd_start == 0 ? 0.0 : result[2*nq+1]);
+      double weighted_shell = (pd_start == 0 ? 0.0 : result[2*nq+2]);
+      double weighted_radius = (pd_start == 0 ? 0.0 : result[2*nq+3]);
       double this_F2 = (pd_start == 0 ? 0.0 : result[2*q_index+0]);
       double this_F1 = (pd_start == 0 ? 0.0 : result[2*q_index+1]);
     #else
       double weight_norm = (pd_start == 0 ? 0.0 : result[nq]);
-      double weighted_volume = (pd_start == 0 ? 0.0 : result[nq+1]);
-      double weighted_radius = (pd_start == 0 ? 0.0 : result[nq+2]);
+      double weighted_form = (pd_start == 0 ? 0.0 : result[nq+1]);
+      double weighted_shell = (pd_start == 0 ? 0.0 : result[nq+2]);
+      double weighted_radius = (pd_start == 0 ? 0.0 : result[nq+3]);
       double this_result = (pd_start == 0 ? 0.0 : result[q_index]);
     #endif
   #else // !USE_OPENCL
-    #if COMPUTE_F1_F2
+    #if defined(CALL_FQ)
       double weight_norm = (pd_start == 0 ? 0.0 : result[2*nq]);
-      double weighted_volume = (pd_start == 0 ? 0.0 : result[2*nq+1]);
-      double weighted_radius = (pd_start == 0 ? 0.0 : result[2*nq+2]);
+      double weighted_form = (pd_start == 0 ? 0.0 : result[2*nq+1]);
+      double weighted_shell = (pd_start == 0 ? 0.0 : result[2*nq+2]);
+      double weighted_radius = (pd_start == 0 ? 0.0 : result[2*nq+3]);
     #else
       double weight_norm = (pd_start == 0 ? 0.0 : result[nq]);
-      double weighted_volume = (pd_start == 0 ? 0.0 : result[nq+1]);
-      double weighted_radius = (pd_start == 0 ? 0.0 : result[nq+2]);
+      double weighted_form = (pd_start == 0 ? 0.0 : result[nq+1]);
+      double weighted_shell = (pd_start == 0 ? 0.0 : result[nq+2]);
+      double weighted_radius = (pd_start == 0 ? 0.0 : result[nq+3]);
     #endif
     if (pd_start == 0) {
@@ -371,5 +374,5 @@
       #pragma omp parallel for
       #endif
-      #if COMPUTE_F1_F2
+      #if defined(CALL_FQ)
           // 2*nq for F^2,F pairs
           for (int q_index=0; q_index < 2*nq; q_index++) result[q_index] = 0.0;
@@ -378,5 +381,5 @@
       #endif
     }
-    //if (q_index==0) printf("start %d %g %g\n", pd_start, weighted_volume, result[0]);
+    //if (q_index==0) printf("start %d %g %g\n", pd_start, weighted_shell, result[0]);
 #endif // !USE_OPENCL
 
@@ -693,8 +696,9 @@
     // Note: weight==0 must always be excluded
     if (weight > cutoff) {
-      weighted_volume += weight * CALL_VOLUME(local_values.table);
-      #if COMPUTE_F1_F2
+      double form, shell;
+      CALL_VOLUME(form, shell, local_values.table);
       weight_norm += weight;
-      #endif
+      weighted_form += weight * form;
+      weighted_shell += weight * shell;
       if (effective_radius_type != 0) {
         weighted_radius += weight * CALL_EFFECTIVE_RADIUS(effective_radius_type, local_values.table);
@@ -747,5 +751,5 @@
           }
         #else  // !MAGNETIC
-          #if COMPUTE_F1_F2
+          #if defined(CALL_FQ)
             CALL_KERNEL(); // sets F1 and F2 by reference
           #else
@@ -756,5 +760,5 @@
 
         #ifdef USE_OPENCL
-          #if COMPUTE_F1_F2
+          #if defined(CALL_FQ)
             this_F2 += weight * F2;
             this_F1 += weight * F1;
@@ -763,5 +767,5 @@
           #endif
         #else // !USE_OPENCL
-          #if COMPUTE_F1_F2
+          #if defined(CALL_FQ)
             result[2*q_index+0] += weight * F2;
             result[2*q_index+1] += weight * F1;
@@ -793,11 +797,12 @@
 // Remember the current result and the updated norm.
 #ifdef USE_OPENCL
-  #if COMPUTE_F1_F2
+  #if defined(CALL_FQ)
     result[2*q_index+0] = this_F2;
     result[2*q_index+1] = this_F1;
     if (q_index == 0) {
       result[2*nq+0] = weight_norm;
-      result[2*nq+1] = weighted_volume;
-      result[2*nq+2] = weighted_radius;
+      result[2*nq+1] = weighted_form;
+      result[2*nq+3] = weighted_shell;
+      result[2*nq+3] = weighted_radius;
     }
   #else
@@ -805,25 +810,27 @@
     if (q_index == 0) {
       result[nq+0] = weight_norm;
-      result[nq+1] = weighted_volume;
-      result[nq+2] = weighted_radius;
+      result[nq+1] = weighted_form;
+      result[nq+2] = weighted_shell;
+      result[nq+3] = weighted_radius;
     }
   #endif
 
-//if (q_index == 0) printf("res: %g/%g\n", result[0], weigthed_volume);
+//if (q_index == 0) printf("res: %g/%g\n", result[0], weighted_shell);
 #else // !USE_OPENCL
-  #if COMPUTE_F1_F2
+  #if defined(CALL_FQ)
     result[2*nq] = weight_norm;
-    result[2*nq+1] = weighted_volume;
-    result[2*nq+2] = weighted_radius;
+    result[2*nq+1] = weighted_form;
+    result[2*nq+2] = weighted_shell;
+    result[2*nq+3] = weighted_radius;
   #else
     result[nq] = weight_norm;
-    result[nq+1] = weighted_volume;
-    result[nq+2] = weighted_radius;
+    result[nq+1] = weighted_form;
+    result[nq+2] = weighted_shell;
+    result[nq+3] = weighted_radius;
   #endif
-//printf("res: %g/%g\n", result[0], weighted_volume);
+//printf("res: %g/%g\n", result[0], weighted_shell);
 #endif // !USE_OPENCL
 
 // ** clear the macros in preparation for the next kernel **
-#undef COMPUTE_F1_F2
 #undef PD_INIT
 #undef PD_OPEN
