48namespace internal_collidesweepsingleprecisionleesedwardscuda_collidesweepsingleprecisionleesedwardscuda {
49static FUNC_PREFIX __launch_bounds__(256) void collidesweepsingleprecisionleesedwardscuda_collidesweepsingleprecisionleesedwardscuda(
float *
RESTRICT const _data_force,
float *
RESTRICT _data_pdfs, int64_t const
_size_force_0, int64_t const
_size_force_1, int64_t const
_size_force_2, int64_t const
_stride_force_0, int64_t const
_stride_force_1, int64_t const
_stride_force_2, int64_t const
_stride_force_3, int64_t const
_stride_pdfs_0, int64_t const
_stride_pdfs_1, int64_t const
_stride_pdfs_2, int64_t const
_stride_pdfs_3,
float grid_size,
float omega_shear,
float v_s) {
51 const int64_t ctr_0 = blockDim.x * blockIdx.x + threadIdx.x;
52 const int64_t ctr_1 = blockDim.y * blockIdx.y + threadIdx.y;
53 const int64_t ctr_2 = blockDim.z * blockIdx.z + threadIdx.z;
76 const float xi_3 = xi_41;
77 const float xi_4 = xi_36;
78 const float xi_5 = xi_28;
79 const float xi_6 = xi_45;
80 const float xi_7 = xi_46;
81 const float xi_8 = xi_40;
82 const float xi_9 = xi_35;
83 const float xi_10 = xi_42;
84 const float xi_11 = xi_33;
85 const float xi_12 = xi_34;
86 const float xi_13 = xi_32;
87 const float xi_14 = xi_38;
88 const float xi_15 = xi_39;
89 const float xi_16 = xi_44;
90 const float xi_17 = xi_25;
91 const float xi_18 = xi_30;
92 const float xi_19 = xi_31;
93 const float xi_20 = xi_26;
94 const float xi_21 = xi_27;
95 const float xi_22 = xi_29;
96 const float xi_23 = xi_37;
97 const float xi_24 = xi_43;
98 const float xi_0 = ((1.0f) / (
omega_shear * -0.25f + 2.0f));
99 const float rr_0 = xi_0 * (
omega_shear * -2.0f + 4.0f);
100 const float vel0Term = xi_11 + xi_12 + xi_14 + xi_18 + xi_24;
101 const float vel1Term = xi_20 + xi_21 + xi_7 + xi_8;
102 const float vel2Term = xi_15 + xi_22 + xi_23;
103 const float rho = vel0Term + vel1Term + vel2Term + xi_13 + xi_16 + xi_17 + xi_19 + xi_3 + xi_4 + xi_6;
104 const float xi_1 = ((1.0f) / (rho));
105 const float u_0 = xi_1 * xi_10 * 0.5f + xi_1 * (vel0Term - xi_13 - xi_15 - xi_3 - xi_6 - xi_7);
106 const float u_1 = xi_1 * xi_9 * 0.5f + xi_1 * (vel1Term + xi_12 - xi_13 - xi_19 - xi_23 - xi_24 - xi_4);
107 const float u_2 = xi_1 * xi_5 * 0.5f + xi_1 * (vel2Term + xi_11 - xi_16 - xi_18 - xi_20 + xi_21 - xi_3 - xi_4);
108 const float forceTerm_0 =
omega_shear * u_0 * xi_10 * 0.5f +
omega_shear * u_1 * xi_9 * 0.5f +
omega_shear * u_2 * xi_5 * 0.5f - u_0 * xi_10 - u_1 * xi_9 - u_2 * xi_5;
109 const float forceTerm_1 =
omega_shear * u_0 * xi_10 * 0.083333333333333329f +
omega_shear * u_1 * xi_9 * -0.16666666666666666f +
omega_shear * u_2 * xi_5 * 0.083333333333333329f + rr_0 * xi_9 * -0.083333333333333329f + u_0 * xi_10 * -0.16666666666666666f + u_1 * xi_9 * 0.33333333333333331f + u_2 * xi_5 * -0.16666666666666666f + xi_9 * 0.16666666666666666f;
110 const float forceTerm_2 =
omega_shear * u_0 * xi_10 * 0.083333333333333329f +
omega_shear * u_1 * xi_9 * -0.16666666666666666f +
omega_shear * u_2 * xi_5 * 0.083333333333333329f + rr_0 * xi_9 * 0.083333333333333329f + u_0 * xi_10 * -0.16666666666666666f + u_1 * xi_9 * 0.33333333333333331f + u_2 * xi_5 * -0.16666666666666666f + xi_9 * -0.16666666666666666f;
111 const float forceTerm_3 =
omega_shear * u_0 * xi_10 * -0.16666666666666666f +
omega_shear * u_1 * xi_9 * 0.083333333333333329f +
omega_shear * u_2 * xi_5 * 0.083333333333333329f + rr_0 * xi_10 * 0.083333333333333329f + u_0 * xi_10 * 0.33333333333333331f + u_1 * xi_9 * -0.16666666666666666f + u_2 * xi_5 * -0.16666666666666666f + xi_10 * -0.16666666666666666f;
112 const float forceTerm_4 =
omega_shear * u_0 * xi_10 * -0.16666666666666666f +
omega_shear * u_1 * xi_9 * 0.083333333333333329f +
omega_shear * u_2 * xi_5 * 0.083333333333333329f + rr_0 * xi_10 * -0.083333333333333329f + u_0 * xi_10 * 0.33333333333333331f + u_1 * xi_9 * -0.16666666666666666f + u_2 * xi_5 * -0.16666666666666666f + xi_10 * 0.16666666666666666f;
113 const float forceTerm_5 =
omega_shear * u_0 * xi_10 * 0.083333333333333329f +
omega_shear * u_1 * xi_9 * 0.083333333333333329f +
omega_shear * u_2 * xi_5 * -0.16666666666666666f + rr_0 * xi_5 * -0.083333333333333329f + u_0 * xi_10 * -0.16666666666666666f + u_1 * xi_9 * -0.16666666666666666f + u_2 * xi_5 * 0.33333333333333331f + xi_5 * 0.16666666666666666f;
114 const float forceTerm_6 =
omega_shear * u_0 * xi_10 * 0.083333333333333329f +
omega_shear * u_1 * xi_9 * 0.083333333333333329f +
omega_shear * u_2 * xi_5 * -0.16666666666666666f + rr_0 * xi_5 * 0.083333333333333329f + u_0 * xi_10 * -0.16666666666666666f + u_1 * xi_9 * -0.16666666666666666f + u_2 * xi_5 * 0.33333333333333331f + xi_5 * -0.16666666666666666f;
115 const float forceTerm_7 =
omega_shear * u_0 * xi_10 * -0.083333333333333329f +
omega_shear * u_0 * xi_9 * 0.125f +
omega_shear * u_1 * xi_10 * 0.125f +
omega_shear * u_1 * xi_9 * -0.083333333333333329f +
omega_shear * u_2 * xi_5 * 0.041666666666666664f + rr_0 * xi_10 * 0.041666666666666664f + rr_0 * xi_9 * -0.041666666666666664f + u_0 * xi_10 * 0.16666666666666666f + u_0 * xi_9 * -0.25f + u_1 * xi_10 * -0.25f + u_1 * xi_9 * 0.16666666666666666f + u_2 * xi_5 * -0.083333333333333329f + xi_10 * -0.083333333333333329f + xi_9 * 0.083333333333333329f;
116 const float forceTerm_8 =
omega_shear * u_0 * xi_10 * -0.083333333333333329f +
omega_shear * u_0 * xi_9 * -0.125f +
omega_shear * u_1 * xi_10 * -0.125f +
omega_shear * u_1 * xi_9 * -0.083333333333333329f +
omega_shear * u_2 * xi_5 * 0.041666666666666664f + rr_0 * xi_10 * -0.041666666666666664f + rr_0 * xi_9 * -0.041666666666666664f + u_0 * xi_10 * 0.16666666666666666f + u_0 * xi_9 * 0.25f + u_1 * xi_10 * 0.25f + u_1 * xi_9 * 0.16666666666666666f + u_2 * xi_5 * -0.083333333333333329f + xi_10 * 0.083333333333333329f + xi_9 * 0.083333333333333329f;
117 const float forceTerm_9 =
omega_shear * u_0 * xi_10 * -0.083333333333333329f +
omega_shear * u_0 * xi_9 * -0.125f +
omega_shear * u_1 * xi_10 * -0.125f +
omega_shear * u_1 * xi_9 * -0.083333333333333329f +
omega_shear * u_2 * xi_5 * 0.041666666666666664f + rr_0 * xi_10 * 0.041666666666666664f + rr_0 * xi_9 * 0.041666666666666664f + u_0 * xi_10 * 0.16666666666666666f + u_0 * xi_9 * 0.25f + u_1 * xi_10 * 0.25f + u_1 * xi_9 * 0.16666666666666666f + u_2 * xi_5 * -0.083333333333333329f + xi_10 * -0.083333333333333329f + xi_9 * -0.083333333333333329f;
118 const float forceTerm_10 =
omega_shear * u_0 * xi_10 * -0.083333333333333329f +
omega_shear * u_0 * xi_9 * 0.125f +
omega_shear * u_1 * xi_10 * 0.125f +
omega_shear * u_1 * xi_9 * -0.083333333333333329f +
omega_shear * u_2 * xi_5 * 0.041666666666666664f + rr_0 * xi_10 * -0.041666666666666664f + rr_0 * xi_9 * 0.041666666666666664f + u_0 * xi_10 * 0.16666666666666666f + u_0 * xi_9 * -0.25f + u_1 * xi_10 * -0.25f + u_1 * xi_9 * 0.16666666666666666f + u_2 * xi_5 * -0.083333333333333329f + xi_10 * 0.083333333333333329f + xi_9 * -0.083333333333333329f;
119 const float forceTerm_11 =
omega_shear * u_0 * xi_10 * 0.041666666666666664f +
omega_shear * u_1 * xi_5 * -0.125f +
omega_shear * u_1 * xi_9 * -0.083333333333333329f +
omega_shear * u_2 * xi_5 * -0.083333333333333329f +
omega_shear * u_2 * xi_9 * -0.125f + rr_0 * xi_5 * -0.041666666666666664f + rr_0 * xi_9 * -0.041666666666666664f + u_0 * xi_10 * -0.083333333333333329f + u_1 * xi_5 * 0.25f + u_1 * xi_9 * 0.16666666666666666f + u_2 * xi_5 * 0.16666666666666666f + u_2 * xi_9 * 0.25f + xi_5 * 0.083333333333333329f + xi_9 * 0.083333333333333329f;
120 const float forceTerm_12 =
omega_shear * u_0 * xi_10 * 0.041666666666666664f +
omega_shear * u_1 * xi_5 * 0.125f +
omega_shear * u_1 * xi_9 * -0.083333333333333329f +
omega_shear * u_2 * xi_5 * -0.083333333333333329f +
omega_shear * u_2 * xi_9 * 0.125f + rr_0 * xi_5 * -0.041666666666666664f + rr_0 * xi_9 * 0.041666666666666664f + u_0 * xi_10 * -0.083333333333333329f + u_1 * xi_5 * -0.25f + u_1 * xi_9 * 0.16666666666666666f + u_2 * xi_5 * 0.16666666666666666f + u_2 * xi_9 * -0.25f + xi_5 * 0.083333333333333329f + xi_9 * -0.083333333333333329f;
121 const float forceTerm_13 =
omega_shear * u_0 * xi_10 * -0.083333333333333329f +
omega_shear * u_0 * xi_5 * 0.125f +
omega_shear * u_1 * xi_9 * 0.041666666666666664f +
omega_shear * u_2 * xi_10 * 0.125f +
omega_shear * u_2 * xi_5 * -0.083333333333333329f + rr_0 * xi_10 * 0.041666666666666664f + rr_0 * xi_5 * -0.041666666666666664f + u_0 * xi_10 * 0.16666666666666666f + u_0 * xi_5 * -0.25f + u_1 * xi_9 * -0.083333333333333329f + u_2 * xi_10 * -0.25f + u_2 * xi_5 * 0.16666666666666666f + xi_10 * -0.083333333333333329f + xi_5 * 0.083333333333333329f;
122 const float forceTerm_14 =
omega_shear * u_0 * xi_10 * -0.083333333333333329f +
omega_shear * u_0 * xi_5 * -0.125f +
omega_shear * u_1 * xi_9 * 0.041666666666666664f +
omega_shear * u_2 * xi_10 * -0.125f +
omega_shear * u_2 * xi_5 * -0.083333333333333329f + rr_0 * xi_10 * -0.041666666666666664f + rr_0 * xi_5 * -0.041666666666666664f + u_0 * xi_10 * 0.16666666666666666f + u_0 * xi_5 * 0.25f + u_1 * xi_9 * -0.083333333333333329f + u_2 * xi_10 * 0.25f + u_2 * xi_5 * 0.16666666666666666f + xi_10 * 0.083333333333333329f + xi_5 * 0.083333333333333329f;
123 const float forceTerm_15 =
omega_shear * u_0 * xi_10 * 0.041666666666666664f +
omega_shear * u_1 * xi_5 * 0.125f +
omega_shear * u_1 * xi_9 * -0.083333333333333329f +
omega_shear * u_2 * xi_5 * -0.083333333333333329f +
omega_shear * u_2 * xi_9 * 0.125f + rr_0 * xi_5 * 0.041666666666666664f + rr_0 * xi_9 * -0.041666666666666664f + u_0 * xi_10 * -0.083333333333333329f + u_1 * xi_5 * -0.25f + u_1 * xi_9 * 0.16666666666666666f + u_2 * xi_5 * 0.16666666666666666f + u_2 * xi_9 * -0.25f + xi_5 * -0.083333333333333329f + xi_9 * 0.083333333333333329f;
124 const float forceTerm_16 =
omega_shear * u_0 * xi_10 * 0.041666666666666664f +
omega_shear * u_1 * xi_5 * -0.125f +
omega_shear * u_1 * xi_9 * -0.083333333333333329f +
omega_shear * u_2 * xi_5 * -0.083333333333333329f +
omega_shear * u_2 * xi_9 * -0.125f + rr_0 * xi_5 * 0.041666666666666664f + rr_0 * xi_9 * 0.041666666666666664f + u_0 * xi_10 * -0.083333333333333329f + u_1 * xi_5 * 0.25f + u_1 * xi_9 * 0.16666666666666666f + u_2 * xi_5 * 0.16666666666666666f + u_2 * xi_9 * 0.25f + xi_5 * -0.083333333333333329f + xi_9 * -0.083333333333333329f;
125 const float forceTerm_17 =
omega_shear * u_0 * xi_10 * -0.083333333333333329f +
omega_shear * u_0 * xi_5 * -0.125f +
omega_shear * u_1 * xi_9 * 0.041666666666666664f +
omega_shear * u_2 * xi_10 * -0.125f +
omega_shear * u_2 * xi_5 * -0.083333333333333329f + rr_0 * xi_10 * 0.041666666666666664f + rr_0 * xi_5 * 0.041666666666666664f + u_0 * xi_10 * 0.16666666666666666f + u_0 * xi_5 * 0.25f + u_1 * xi_9 * -0.083333333333333329f + u_2 * xi_10 * 0.25f + u_2 * xi_5 * 0.16666666666666666f + xi_10 * -0.083333333333333329f + xi_5 * -0.083333333333333329f;
126 const float forceTerm_18 =
omega_shear * u_0 * xi_10 * -0.083333333333333329f +
omega_shear * u_0 * xi_5 * 0.125f +
omega_shear * u_1 * xi_9 * 0.041666666666666664f +
omega_shear * u_2 * xi_10 * 0.125f +
omega_shear * u_2 * xi_5 * -0.083333333333333329f + rr_0 * xi_10 * -0.041666666666666664f + rr_0 * xi_5 * 0.041666666666666664f + u_0 * xi_10 * 0.16666666666666666f + u_0 * xi_5 * -0.25f + u_1 * xi_9 * -0.083333333333333329f + u_2 * xi_10 * -0.25f + u_2 * xi_5 * 0.16666666666666666f + xi_10 * 0.083333333333333329f + xi_5 * -0.083333333333333329f;
127 const float u0Mu1 = u_0 - u_1;
128 const float u0Pu1 = u_0 + u_1;
129 const float u1Pu2 = u_1 + u_2;
130 const float u1Mu2 = u_1 - u_2;
131 const float u0Mu2 = u_0 - u_2;
132 const float u0Pu2 = u_0 + u_2;
133 const float f_eq_common = rho - rho * u_0 * u_0 - rho * u_1 * u_1 - rho * u_2 * u_2;
135 _data_pdfs[
_stride_pdfs_0 * ctr_0 +
_stride_pdfs_1 * ctr_1 +
_stride_pdfs_2 * ctr_2 +
_stride_pdfs_3] = forceTerm_1 +
omega_shear * (f_eq_common * 0.16666666666666666f + rho * (-0.1111111111111111f + 0.33333333333333331f * (u_1 * u_1)) + xi_19 * -0.5f + xi_8 * -0.5f) + rr_0 * (rho * u_1 * 0.16666666666666666f + xi_19 * 0.5f + xi_8 * -0.5f) + xi_8 + ((-1.0f <= -
grid_size + ((float)(ctr_1))) ? (rho * v_s * (u_0 * 2.0f + v_s) * 0.16666666666666666f) : (0.0f));
136 _data_pdfs[
_stride_pdfs_0 * ctr_0 +
_stride_pdfs_1 * ctr_1 +
_stride_pdfs_2 * ctr_2 + 2 *
_stride_pdfs_3] = forceTerm_2 +
omega_shear * (f_eq_common * 0.16666666666666666f + rho * (-0.1111111111111111f + 0.33333333333333331f * (u_1 * u_1)) + xi_19 * -0.5f + xi_8 * -0.5f) + rr_0 * (rho * u_1 * -0.16666666666666666f + xi_19 * -0.5f + xi_8 * 0.5f) + xi_19 + ((0.0f >= ((float)(ctr_1))) ? (rho * v_s * (u_0 * -2.0f + v_s) * 0.16666666666666666f) : (0.0f));
141 _data_pdfs[
_stride_pdfs_0 * ctr_0 +
_stride_pdfs_1 * ctr_1 +
_stride_pdfs_2 * ctr_2 + 7 *
_stride_pdfs_3] = forceTerm_7 +
omega_shear * (f_eq_common * 0.041666666666666664f + rho * (-0.013888888888888888f + 0.041666666666666664f * (u_2 * u_2) + 0.125f * (u0Mu1 * u0Mu1)) + xi_24 * -0.5f + xi_7 * -0.5f) + rr_0 * (rho * u0Mu1 * -0.083333333333333329f + xi_24 * 0.5f + xi_7 * -0.5f) + xi_7 + ((-1.0f <= -
grid_size + ((
float)(ctr_1))) ? (rho * v_s * (u_0 * -2.0f + u_1 * 3.0f - v_s + 1.0f) * 0.083333333333333329f) : (0.0f));
142 _data_pdfs[
_stride_pdfs_0 * ctr_0 +
_stride_pdfs_1 * ctr_1 +
_stride_pdfs_2 * ctr_2 + 8 *
_stride_pdfs_3] = forceTerm_8 +
omega_shear * (f_eq_common * 0.041666666666666664f + rho * (-0.013888888888888888f + 0.041666666666666664f * (u_2 * u_2) + 0.125f * (u0Pu1 * u0Pu1)) + xi_12 * -0.5f + xi_13 * -0.5f) + rr_0 * (rho * u0Pu1 * 0.083333333333333329f + xi_12 * -0.5f + xi_13 * 0.5f) + xi_12 + ((-1.0f <= -
grid_size + ((
float)(ctr_1))) ? (rho * v_s * (u_0 * 2.0f + u_1 * 3.0f + v_s + 1.0f) * -0.083333333333333329f) : (0.0f));
143 _data_pdfs[
_stride_pdfs_0 * ctr_0 +
_stride_pdfs_1 * ctr_1 +
_stride_pdfs_2 * ctr_2 + 9 *
_stride_pdfs_3] = forceTerm_9 +
omega_shear * (f_eq_common * 0.041666666666666664f + rho * (-0.013888888888888888f + 0.041666666666666664f * (u_2 * u_2) + 0.125f * (u0Pu1 * u0Pu1)) + xi_12 * -0.5f + xi_13 * -0.5f) + rr_0 * (rho * u0Pu1 * -0.083333333333333329f + xi_12 * 0.5f + xi_13 * -0.5f) + xi_13 + ((0.0f >= ((
float)(ctr_1))) ? (rho * v_s * (u_0 * 2.0f + u_1 * 3.0f - v_s - 1.0f) * 0.083333333333333329f) : (0.0f));
144 _data_pdfs[
_stride_pdfs_0 * ctr_0 +
_stride_pdfs_1 * ctr_1 +
_stride_pdfs_2 * ctr_2 + 10 *
_stride_pdfs_3] = forceTerm_10 +
omega_shear * (f_eq_common * 0.041666666666666664f + rho * (-0.013888888888888888f + 0.041666666666666664f * (u_2 * u_2) + 0.125f * (u0Mu1 * u0Mu1)) + xi_24 * -0.5f + xi_7 * -0.5f) + rr_0 * (rho * u0Mu1 * 0.083333333333333329f + xi_24 * -0.5f + xi_7 * 0.5f) + xi_24 + ((0.0f >= ((
float)(ctr_1))) ? (rho * v_s * (u_0 * 2.0f + u_1 * -3.0f - v_s + 1.0f) * 0.083333333333333329f) : (0.0f));
145 _data_pdfs[
_stride_pdfs_0 * ctr_0 +
_stride_pdfs_1 * ctr_1 +
_stride_pdfs_2 * ctr_2 + 11 *
_stride_pdfs_3] = forceTerm_11 +
omega_shear * (f_eq_common * 0.041666666666666664f + rho * (-0.013888888888888888f + 0.041666666666666664f * (u_0 * u_0) + 0.125f * (u1Pu2 * u1Pu2)) + xi_21 * -0.5f + xi_4 * -0.5f) + rr_0 * (rho * u1Pu2 * 0.083333333333333329f + xi_21 * -0.5f + xi_4 * 0.5f) + xi_21;
146 _data_pdfs[
_stride_pdfs_0 * ctr_0 +
_stride_pdfs_1 * ctr_1 +
_stride_pdfs_2 * ctr_2 + 12 *
_stride_pdfs_3] = forceTerm_12 +
omega_shear * (f_eq_common * 0.041666666666666664f + rho * (-0.013888888888888888f + 0.041666666666666664f * (u_0 * u_0) + 0.125f * (u1Mu2 * u1Mu2)) + xi_20 * -0.5f + xi_23 * -0.5f) + rr_0 * (rho * u1Mu2 * -0.083333333333333329f + xi_20 * 0.5f + xi_23 * -0.5f) + xi_23;
147 _data_pdfs[
_stride_pdfs_0 * ctr_0 +
_stride_pdfs_1 * ctr_1 +
_stride_pdfs_2 * ctr_2 + 13 *
_stride_pdfs_3] = forceTerm_13 +
omega_shear * (f_eq_common * 0.041666666666666664f + rho * (-0.013888888888888888f + 0.041666666666666664f * (u_1 * u_1) + 0.125f * (u0Mu2 * u0Mu2)) + xi_15 * -0.5f + xi_18 * -0.5f) + rr_0 * (rho * u0Mu2 * -0.083333333333333329f + xi_15 * -0.5f + xi_18 * 0.5f) + xi_15;
148 _data_pdfs[
_stride_pdfs_0 * ctr_0 +
_stride_pdfs_1 * ctr_1 +
_stride_pdfs_2 * ctr_2 + 14 *
_stride_pdfs_3] = forceTerm_14 +
omega_shear * (f_eq_common * 0.041666666666666664f + rho * (-0.013888888888888888f + 0.041666666666666664f * (u_1 * u_1) + 0.125f * (u0Pu2 * u0Pu2)) + xi_11 * -0.5f + xi_3 * -0.5f) + rr_0 * (rho * u0Pu2 * 0.083333333333333329f + xi_11 * -0.5f + xi_3 * 0.5f) + xi_11;
149 _data_pdfs[
_stride_pdfs_0 * ctr_0 +
_stride_pdfs_1 * ctr_1 +
_stride_pdfs_2 * ctr_2 + 15 *
_stride_pdfs_3] = forceTerm_15 +
omega_shear * (f_eq_common * 0.041666666666666664f + rho * (-0.013888888888888888f + 0.041666666666666664f * (u_0 * u_0) + 0.125f * (u1Mu2 * u1Mu2)) + xi_20 * -0.5f + xi_23 * -0.5f) + rr_0 * (rho * u1Mu2 * 0.083333333333333329f + xi_20 * -0.5f + xi_23 * 0.5f) + xi_20;
150 _data_pdfs[
_stride_pdfs_0 * ctr_0 +
_stride_pdfs_1 * ctr_1 +
_stride_pdfs_2 * ctr_2 + 16 *
_stride_pdfs_3] = forceTerm_16 +
omega_shear * (f_eq_common * 0.041666666666666664f + rho * (-0.013888888888888888f + 0.041666666666666664f * (u_0 * u_0) + 0.125f * (u1Pu2 * u1Pu2)) + xi_21 * -0.5f + xi_4 * -0.5f) + rr_0 * (rho * u1Pu2 * -0.083333333333333329f + xi_21 * 0.5f + xi_4 * -0.5f) + xi_4;
151 _data_pdfs[
_stride_pdfs_0 * ctr_0 +
_stride_pdfs_1 * ctr_1 +
_stride_pdfs_2 * ctr_2 + 17 *
_stride_pdfs_3] = forceTerm_17 +
omega_shear * (f_eq_common * 0.041666666666666664f + rho * (-0.013888888888888888f + 0.041666666666666664f * (u_1 * u_1) + 0.125f * (u0Pu2 * u0Pu2)) + xi_11 * -0.5f + xi_3 * -0.5f) + rr_0 * (rho * u0Pu2 * -0.083333333333333329f + xi_11 * 0.5f + xi_3 * -0.5f) + xi_3;
152 _data_pdfs[
_stride_pdfs_0 * ctr_0 +
_stride_pdfs_1 * ctr_1 +
_stride_pdfs_2 * ctr_2 + 18 *
_stride_pdfs_3] = forceTerm_18 +
omega_shear * (f_eq_common * 0.041666666666666664f + rho * (-0.013888888888888888f + 0.041666666666666664f * (u_1 * u_1) + 0.125f * (u0Mu2 * u0Mu2)) + xi_15 * -0.5f + xi_18 * -0.5f) + rr_0 * (rho * u0Mu2 * 0.083333333333333329f + xi_15 * 0.5f + xi_18 * -0.5f) + xi_18;
159 auto pdfs =
block->getData<gpu::GPUField<float>>(
pdfsID);
160 auto force =
block->getData<gpu::GPUField<float>>(
forceID);
162 auto &v_s = this->
v_s_;
165 WALBERLA_ASSERT_GREATER_EQUAL(0, -int_c(force->nrOfGhostLayers()))
166 float *
RESTRICT const _data_force = force->dataAt(0, 0, 0, 0);
167 WALBERLA_ASSERT_EQUAL(force->layout(), field::fzyx)
168 WALBERLA_ASSERT_GREATER_EQUAL(0, -int_c(pdfs->nrOfGhostLayers()))
169 float *
RESTRICT _data_pdfs = pdfs->dataAt(0, 0, 0, 0);
170 WALBERLA_ASSERT_EQUAL(pdfs->layout(), field::fzyx)
171 WALBERLA_ASSERT_GREATER_EQUAL(force->xSizeWithGhostLayer(), int64_t(int64_c(force->xSize()) + 0))
172 const int64_t _size_force_0 = int64_t(int64_c(force->xSize()) + 0);
173 WALBERLA_ASSERT_EQUAL(force->layout(), field::fzyx)
174 WALBERLA_ASSERT_GREATER_EQUAL(force->ySizeWithGhostLayer(), int64_t(int64_c(force->ySize()) + 0))
175 const int64_t _size_force_1 = int64_t(int64_c(force->ySize()) + 0);
176 WALBERLA_ASSERT_EQUAL(force->layout(), field::fzyx)
177 WALBERLA_ASSERT_GREATER_EQUAL(force->zSizeWithGhostLayer(), int64_t(int64_c(force->zSize()) + 0))
178 const int64_t _size_force_2 = int64_t(int64_c(force->zSize()) + 0);
179 WALBERLA_ASSERT_EQUAL(force->layout(), field::fzyx)
180 const int64_t _stride_force_0 = int64_t(force->xStride());
181 const int64_t _stride_force_1 = int64_t(force->yStride());
182 const int64_t _stride_force_2 = int64_t(force->zStride());
183 const int64_t _stride_force_3 = int64_t(1 * int64_t(force->fStride()));
184 const int64_t _stride_pdfs_0 = int64_t(pdfs->xStride());
185 const int64_t _stride_pdfs_1 = int64_t(pdfs->yStride());
186 const int64_t _stride_pdfs_2 = int64_t(pdfs->zStride());
187 const int64_t _stride_pdfs_3 = int64_t(1 * int64_t(pdfs->fStride()));
188 dim3 _block(uint32_c(((128 < _size_force_0) ? 128 : _size_force_0)), uint32_c(((1024 < ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))) ? 1024 : ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))))), uint32_c(((64 < ((_size_force_2 < ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))))) ? _size_force_2 : ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))))))) ? 64 : ((_size_force_2 < ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))))) ? _size_force_2 : ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))))))));
189 dim3 _grid(uint32_c(((_size_force_0) % (((128 < _size_force_0) ? 128 : _size_force_0)) == 0 ? (int64_t)(_size_force_0) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)) : ((int64_t)(_size_force_0) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))) + 1)), uint32_c(((_size_force_1) % (((1024 < ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))) ? 1024 : ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))))) == 0 ? (int64_t)(_size_force_1) / (int64_t)(((1024 < ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))) ? 1024 : ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))))) : ((int64_t)(_size_force_1) / (int64_t)(((1024 < ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))) ? 1024 : ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))))) + 1)), uint32_c(((_size_force_2) % (((64 < ((_size_force_2 < ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))))) ? _size_force_2 : ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))))))) ? 64 : ((_size_force_2 < ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))))) ? _size_force_2 : ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))))))) == 0 ? (int64_t)(_size_force_2) / (int64_t)(((64 < ((_size_force_2 < ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))))) ? _size_force_2 : ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))))))) ? 64 : ((_size_force_2 < ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))))) ? _size_force_2 : ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))))))) : ((int64_t)(_size_force_2) / (int64_t)(((64 < ((_size_force_2 < ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))))) ? _size_force_2 : ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))))))) ? 64 : ((_size_force_2 < ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))))) ? _size_force_2 : ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))))))))) + 1)));
190 internal_collidesweepsingleprecisionleesedwardscuda_collidesweepsingleprecisionleesedwardscuda::collidesweepsingleprecisionleesedwardscuda_collidesweepsingleprecisionleesedwardscuda<<<_grid, _block, 0, stream>>>(_data_force, _data_pdfs, _size_force_0, _size_force_1, _size_force_2, _stride_force_0, _stride_force_1, _stride_force_2, _stride_force_3, _stride_pdfs_0, _stride_pdfs_1, _stride_pdfs_2, _stride_pdfs_3, grid_size, omega_shear, v_s);
195 CellInterval ci = globalCellInterval;
196 CellInterval blockBB = blocks->getBlockCellBB(*
block);
197 blockBB.expand(ghostLayers);
198 ci.intersect(blockBB);
199 blocks->transformGlobalToBlockLocalCellInterval(ci, *
block);
203 auto pdfs =
block->getData<gpu::GPUField<float>>(
pdfsID);
204 auto force =
block->getData<gpu::GPUField<float>>(
forceID);
206 auto &v_s = this->
v_s_;
209 WALBERLA_ASSERT_GREATER_EQUAL(ci.xMin(), -int_c(force->nrOfGhostLayers()))
210 WALBERLA_ASSERT_GREATER_EQUAL(ci.yMin(), -int_c(force->nrOfGhostLayers()))
211 WALBERLA_ASSERT_GREATER_EQUAL(ci.zMin(), -int_c(force->nrOfGhostLayers()))
212 float *
RESTRICT const _data_force = force->dataAt(ci.xMin(), ci.yMin(), ci.zMin(), 0);
213 WALBERLA_ASSERT_EQUAL(force->layout(), field::fzyx)
214 WALBERLA_ASSERT_GREATER_EQUAL(ci.xMin(), -int_c(pdfs->nrOfGhostLayers()))
215 WALBERLA_ASSERT_GREATER_EQUAL(ci.yMin(), -int_c(pdfs->nrOfGhostLayers()))
216 WALBERLA_ASSERT_GREATER_EQUAL(ci.zMin(), -int_c(pdfs->nrOfGhostLayers()))
217 float *
RESTRICT _data_pdfs = pdfs->dataAt(ci.xMin(), ci.yMin(), ci.zMin(), 0);
218 WALBERLA_ASSERT_EQUAL(pdfs->layout(), field::fzyx)
219 WALBERLA_ASSERT_GREATER_EQUAL(force->xSizeWithGhostLayer(), int64_t(int64_c(ci.xSize()) + 0))
220 const int64_t _size_force_0 = int64_t(int64_c(ci.xSize()) + 0);
221 WALBERLA_ASSERT_EQUAL(force->layout(), field::fzyx)
222 WALBERLA_ASSERT_GREATER_EQUAL(force->ySizeWithGhostLayer(), int64_t(int64_c(ci.ySize()) + 0))
223 const int64_t _size_force_1 = int64_t(int64_c(ci.ySize()) + 0);
224 WALBERLA_ASSERT_EQUAL(force->layout(), field::fzyx)
225 WALBERLA_ASSERT_GREATER_EQUAL(force->zSizeWithGhostLayer(), int64_t(int64_c(ci.zSize()) + 0))
226 const int64_t _size_force_2 = int64_t(int64_c(ci.zSize()) + 0);
227 WALBERLA_ASSERT_EQUAL(force->layout(), field::fzyx)
228 const int64_t _stride_force_0 = int64_t(force->xStride());
229 const int64_t _stride_force_1 = int64_t(force->yStride());
230 const int64_t _stride_force_2 = int64_t(force->zStride());
231 const int64_t _stride_force_3 = int64_t(1 * int64_t(force->fStride()));
232 const int64_t _stride_pdfs_0 = int64_t(pdfs->xStride());
233 const int64_t _stride_pdfs_1 = int64_t(pdfs->yStride());
234 const int64_t _stride_pdfs_2 = int64_t(pdfs->zStride());
235 const int64_t _stride_pdfs_3 = int64_t(1 * int64_t(pdfs->fStride()));
236 dim3 _block(uint32_c(((128 < _size_force_0) ? 128 : _size_force_0)), uint32_c(((1024 < ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))) ? 1024 : ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))))), uint32_c(((64 < ((_size_force_2 < ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))))) ? _size_force_2 : ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))))))) ? 64 : ((_size_force_2 < ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))))) ? _size_force_2 : ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))))))));
237 dim3 _grid(uint32_c(((_size_force_0) % (((128 < _size_force_0) ? 128 : _size_force_0)) == 0 ? (int64_t)(_size_force_0) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)) : ((int64_t)(_size_force_0) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))) + 1)), uint32_c(((_size_force_1) % (((1024 < ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))) ? 1024 : ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))))) == 0 ? (int64_t)(_size_force_1) / (int64_t)(((1024 < ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))) ? 1024 : ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))))) : ((int64_t)(_size_force_1) / (int64_t)(((1024 < ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))) ? 1024 : ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))))) + 1)), uint32_c(((_size_force_2) % (((64 < ((_size_force_2 < ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))))) ? _size_force_2 : ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))))))) ? 64 : ((_size_force_2 < ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))))) ? _size_force_2 : ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))))))) == 0 ? (int64_t)(_size_force_2) / (int64_t)(((64 < ((_size_force_2 < ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))))) ? _size_force_2 : ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))))))) ? 64 : ((_size_force_2 < ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))))) ? _size_force_2 : ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))))))) : ((int64_t)(_size_force_2) / (int64_t)(((64 < ((_size_force_2 < ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))))) ? _size_force_2 : ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))))))) ? 64 : ((_size_force_2 < ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0))))))) ? _size_force_2 : ((int64_t)(256) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0) * ((_size_force_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))) ? _size_force_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_force_0) ? 128 : _size_force_0)))))))))) + 1)));
238 internal_collidesweepsingleprecisionleesedwardscuda_collidesweepsingleprecisionleesedwardscuda::collidesweepsingleprecisionleesedwardscuda_collidesweepsingleprecisionleesedwardscuda<<<_grid, _block, 0, stream>>>(_data_force, _data_pdfs, _size_force_0, _size_force_1, _size_force_2, _stride_force_0, _stride_force_1, _stride_force_2, _stride_force_3, _stride_pdfs_0, _stride_pdfs_1, _stride_pdfs_2, _stride_pdfs_3, grid_size, omega_shear, v_s);