Changeset 74e9b5f in sasmodels for sasmodels/kernel_iq.c

Timestamp:

Oct 12, 2018 10:52:48 PM (6 years ago)

Author:

pkienzle

Branches:

master, core_shell_microgels, magnetic_model, ticket-1257-vesicle-product, ticket_1156, ticket_1265_superball, ticket_822_more_unit_tests

Children:

4de14584

Parents:

b0de252

Message:

autotag functions as device functions for cuda. Refs #1076.

File:

• sasmodels/kernel_iq.c (modified) (11 diffs)

sasmodels/kernel_iq.c

@@ -67 +67 @@
 
 // Return value restricted between low and high
-__device__
 static double clip(double value, double low, double high)
 {
@@ -80 +79 @@
 //     du * (m_sigma_y + 1j*m_sigma_z);
 // weights for spin crosssections: dd du real, ud real, uu, du imag, ud imag
-__device__
 static void set_spin_weights(double in_spin, double out_spin, double weight[6])
 {
@@ -105 +103 @@
 
 // Compute the magnetic sld
-__device__
 static double mag_sld(
   const unsigned int xs, // 0=dd, 1=du.real, 2=ud.real, 3=uu, 4=du.imag, 5=ud.imag
@@ -154 +151 @@
 // jitter angles (dtheta, dphi).  This matrix can be applied to all of the
 // (qx, qy) points in the image to produce R*[qx,qy]' = [qa,qc]'
-__device__
 static void
 qac_rotation(
@@ -188 +184 @@
 // Apply the rotation matrix returned from qac_rotation to the point (qx,qy),
 // returning R*[qx,qy]' = [qa,qc]'
-__device__
 static void
 qac_apply(
@@ -216 +211 @@
 // jitter angles (dtheta, dphi, dpsi).  This matrix can be applied to all of the
 // (qx, qy) points in the image to produce R*[qx,qy]' = [qa,qb,qc]'
-__device__
 static void
 qabc_rotation(
@@ -263 +257 @@
 // Apply the rotation matrix returned from qabc_rotation to the point (qx,qy),
 // returning R*[qx,qy]' = [qa,qb,qc]'
-__device__
 static void
 qabc_apply(
@@ -285 +278 @@
     const int32_t pd_start,     // where we are in the dispersity loop
     const int32_t pd_stop,      // where we are stopping in the dispersity loop
-    global_par const ProblemDetails *details,
-    global_par const double *values,
-    global_par const double *q, // nq q values, with padding to boundary
-    global_par double *result,  // nq+1 return values, again with padding
+    pglobal const ProblemDetails *details,
+    pglobal const double *values,
+    pglobal const double *q, // nq q values, with padding to boundary
+    pglobal double *result,  // nq+1 return values, again with padding
     const double cutoff     // cutoff in the dispersity weight product
     )
@@ -386 +379 @@
   const int n4 = pd_length[4];
   const int p4 = pd_par[4];
-  global_var const double *v4 = pd_value + pd_offset[4];
-  global_var const double *w4 = pd_weight + pd_offset[4];
+  pglobal const double *v4 = pd_value + pd_offset[4];
+  pglobal const double *w4 = pd_weight + pd_offset[4];
   int i4 = (pd_start/pd_stride[4])%n4;  // position in level 4 at pd_start
 
@@ -573 +566 @@
   const int n##_LOOP = details->pd_length[_LOOP]; \
   const int p##_LOOP = details->pd_par[_LOOP]; \
-  global_var const double *v##_LOOP = pd_value + details->pd_offset[_LOOP]; \
-  global_var const double *w##_LOOP = pd_weight + details->pd_offset[_LOOP]; \
+  pglobal const double *v##_LOOP = pd_value + details->pd_offset[_LOOP]; \
+  pglobal const double *w##_LOOP = pd_weight + details->pd_offset[_LOOP]; \
   int i##_LOOP = (pd_start/details->pd_stride[_LOOP])%n##_LOOP;
 
@@ -598 +591 @@
 // Pointers to the start of the dispersity and weight vectors, if needed.
 #if MAX_PD>0
-  global_var const double *pd_value = values + NUM_VALUES;
-  global_var const double *pd_weight = pd_value + details->num_weights;
+  pglobal const double *pd_value = values + NUM_VALUES;
+  pglobal const double *pd_weight = pd_value + details->num_weights;
 #endif