22   SharedData_Hip data;  in __launch_bounds__()  local
23   data.t_id_x = threadIdx.x;  in __launch_bounds__()
24   data.t_id_y = threadIdx.y;  in __launch_bounds__()
25   data.t_id_z = threadIdx.z;  in __launch_bounds__()
26   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in __launch_bounds__()
27   data.slice  = slice + data.t_id_z * BASIS_T_1D;  in __launch_bounds__()
34   LoadMatrix<BASIS_P, BASIS_Q>(data, c_B, s_B);  in __launch_bounds__()
39 …ReadElementStrided1d<BASIS_NUM_COMP, BASIS_P>(data, elem, 1, BASIS_P * num_elem, BASIS_P, d_U, r_U…  in __launch_bounds__()
40     InterpNonTensor<BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_B, r_V);  in __launch_bounds__()
41 …WriteElementStrided1d<BASIS_NUM_COMP, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS_Q, r_V, d_…  in __launch_bounds__()
49   SharedData_Hip data;  in __launch_bounds__()  local
50   data.t_id_x = threadIdx.x;  in __launch_bounds__()
51   data.t_id_y = threadIdx.y;  in __launch_bounds__()
52   data.t_id_z = threadIdx.z;  in __launch_bounds__()
53   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in __launch_bounds__()
54   data.slice  = slice + data.t_id_z * BASIS_T_1D;  in __launch_bounds__()
61   LoadMatrix<BASIS_P, BASIS_Q>(data, c_B, s_B);  in __launch_bounds__()
66 …ReadElementStrided1d<BASIS_NUM_COMP, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS_Q, d_U, r_U…  in __launch_bounds__()
67     InterpTransposeNonTensor<BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_B, r_V);  in __launch_bounds__()
68 …WriteElementStrided1d<BASIS_NUM_COMP, BASIS_P>(data, elem, 1, BASIS_P * num_elem, BASIS_P, r_V, d_…  in __launch_bounds__()
76   SharedData_Hip data;  in __launch_bounds__()  local
77   data.t_id_x = threadIdx.x;  in __launch_bounds__()
78   data.t_id_y = threadIdx.y;  in __launch_bounds__()
79   data.t_id_z = threadIdx.z;  in __launch_bounds__()
80   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in __launch_bounds__()
81   data.slice  = slice + data.t_id_z * BASIS_T_1D;  in __launch_bounds__()
88   LoadMatrix<BASIS_P, BASIS_Q>(data, c_B, s_B);  in __launch_bounds__()
93 …ReadElementStrided1d<BASIS_NUM_COMP, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS_Q, d_U, r_U…  in __launch_bounds__()
94     InterpTransposeNonTensor<BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_B, r_V);  in __launch_bounds__()
95 …SumElementStrided1d<BASIS_NUM_COMP, BASIS_P>(data, elem, 1, BASIS_P * num_elem, BASIS_P, r_V, d_V);  in __launch_bounds__()
106   SharedData_Hip data;  in __launch_bounds__()  local
107   data.t_id_x = threadIdx.x;  in __launch_bounds__()
108   data.t_id_y = threadIdx.y;  in __launch_bounds__()
109   data.t_id_z = threadIdx.z;  in __launch_bounds__()
110   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in __launch_bounds__()
111   data.slice  = slice + data.t_id_z * BASIS_T_1D;  in __launch_bounds__()
118   LoadMatrix<BASIS_P, BASIS_Q * BASIS_DIM>(data, c_G, s_G);  in __launch_bounds__()
123 …ReadElementStrided1d<BASIS_NUM_COMP, BASIS_P>(data, elem, 1, BASIS_P * num_elem, BASIS_P, d_U, r_U…  in __launch_bounds__()
124     GradNonTensor<BASIS_NUM_COMP, BASIS_DIM, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_G, r_V);  in __launch_bounds__()
125 …WriteElementStrided1d<BASIS_NUM_COMP * BASIS_DIM, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASI…  in __launch_bounds__()
133   SharedData_Hip data;  in __launch_bounds__()  local
134   data.t_id_x = threadIdx.x;  in __launch_bounds__()
135   data.t_id_y = threadIdx.y;  in __launch_bounds__()
136   data.t_id_z = threadIdx.z;  in __launch_bounds__()
137   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in __launch_bounds__()
138   data.slice  = slice + data.t_id_z * BASIS_T_1D;  in __launch_bounds__()
145   LoadMatrix<BASIS_P, BASIS_Q * BASIS_DIM>(data, c_G, s_G);  in __launch_bounds__()
150 …ReadElementStrided1d<BASIS_NUM_COMP * BASIS_DIM, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS…  in __launch_bounds__()
151 …GradTransposeNonTensor<BASIS_NUM_COMP, BASIS_DIM, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_G, r_…  in __launch_bounds__()
152 …WriteElementStrided1d<BASIS_NUM_COMP, BASIS_P>(data, elem, 1, BASIS_P * num_elem, BASIS_P, r_V, d_…  in __launch_bounds__()
160   SharedData_Hip data;  in __launch_bounds__()  local
161   data.t_id_x = threadIdx.x;  in __launch_bounds__()
162   data.t_id_y = threadIdx.y;  in __launch_bounds__()
163   data.t_id_z = threadIdx.z;  in __launch_bounds__()
164   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in __launch_bounds__()
165   data.slice  = slice + data.t_id_z * BASIS_T_1D;  in __launch_bounds__()
172   LoadMatrix<BASIS_P, BASIS_Q * BASIS_DIM>(data, c_G, s_G);  in __launch_bounds__()
177 …ReadElementStrided1d<BASIS_NUM_COMP * BASIS_DIM, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS…  in __launch_bounds__()
178 …GradTransposeNonTensor<BASIS_NUM_COMP, BASIS_DIM, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_G, r_…  in __launch_bounds__()
179 …SumElementStrided1d<BASIS_NUM_COMP, BASIS_P>(data, elem, 1, BASIS_P * num_elem, BASIS_P, r_V, d_V);  in __launch_bounds__()
190   SharedData_Hip data;  in __launch_bounds__()  local
191   data.t_id_x = threadIdx.x;  in __launch_bounds__()
192   data.t_id_y = threadIdx.y;  in __launch_bounds__()
193   data.t_id_z = threadIdx.z;  in __launch_bounds__()
194   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in __launch_bounds__()
195   data.slice  = slice + data.t_id_z * BASIS_T_1D;  in __launch_bounds__()
200     WeightNonTensor<BASIS_P, BASIS_Q>(data, q_weight, r_W);  in __launch_bounds__()
201     WriteElementStrided1d<1, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS_Q, r_W, d_W);  in __launch_bounds__()