21   SharedData_Cuda data;  in Interp()  local
22   data.t_id_x = threadIdx.x;  in Interp()
23   data.t_id_y = threadIdx.y;  in Interp()
24   data.t_id_z = threadIdx.z;  in Interp()
25   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in Interp()
26   data.slice  = slice + data.t_id_z * BASIS_T_1D;  in Interp()
33   LoadMatrix<BASIS_P, BASIS_Q>(data, c_B, s_B);  in Interp()
38 …ReadElementStrided1d<BASIS_NUM_COMP, BASIS_P>(data, elem, 1, BASIS_P * num_elem, BASIS_P, d_U, r_U…  in Interp()
39     InterpNonTensor<BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_B, r_V);  in Interp()
40 …WriteElementStrided1d<BASIS_NUM_COMP, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS_Q, r_V, d_…  in Interp()
48   SharedData_Cuda data;  in InterpTranspose()  local
49   data.t_id_x = threadIdx.x;  in InterpTranspose()
50   data.t_id_y = threadIdx.y;  in InterpTranspose()
51   data.t_id_z = threadIdx.z;  in InterpTranspose()
52   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in InterpTranspose()
53   data.slice  = slice + data.t_id_z * BASIS_T_1D;  in InterpTranspose()
60   LoadMatrix<BASIS_P, BASIS_Q>(data, c_B, s_B);  in InterpTranspose()
65 …ReadElementStrided1d<BASIS_NUM_COMP, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS_Q, d_U, r_U…  in InterpTranspose()
66     InterpTransposeNonTensor<BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_B, r_V);  in InterpTranspose()
67 …WriteElementStrided1d<BASIS_NUM_COMP, BASIS_P>(data, elem, 1, BASIS_P * num_elem, BASIS_P, r_V, d_…  in InterpTranspose()
75   SharedData_Cuda data;  in InterpTransposeAdd()  local
76   data.t_id_x = threadIdx.x;  in InterpTransposeAdd()
77   data.t_id_y = threadIdx.y;  in InterpTransposeAdd()
78   data.t_id_z = threadIdx.z;  in InterpTransposeAdd()
79   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in InterpTransposeAdd()
80   data.slice  = slice + data.t_id_z * BASIS_T_1D;  in InterpTransposeAdd()
87   LoadMatrix<BASIS_P, BASIS_Q>(data, c_B, s_B);  in InterpTransposeAdd()
92 …ReadElementStrided1d<BASIS_NUM_COMP, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS_Q, d_U, r_U…  in InterpTransposeAdd()
93     InterpTransposeNonTensor<BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_B, r_V);  in InterpTransposeAdd()
94 …SumElementStrided1d<BASIS_NUM_COMP, BASIS_P>(data, elem, 1, BASIS_P * num_elem, BASIS_P, r_V, d_V);  in InterpTransposeAdd()
104   SharedData_Cuda data;  in Grad()  local
105   data.t_id_x = threadIdx.x;  in Grad()
106   data.t_id_y = threadIdx.y;  in Grad()
107   data.t_id_z = threadIdx.z;  in Grad()
108   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in Grad()
109   data.slice  = slice + data.t_id_z * BASIS_T_1D;  in Grad()
116   LoadMatrix<BASIS_P, BASIS_Q * BASIS_DIM>(data, c_G, s_G);  in Grad()
121 …ReadElementStrided1d<BASIS_NUM_COMP, BASIS_P>(data, elem, 1, BASIS_P * num_elem, BASIS_P, d_U, r_U…  in Grad()
122     GradNonTensor<BASIS_NUM_COMP, BASIS_DIM, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_G, r_V);  in Grad()
123 …WriteElementStrided1d<BASIS_NUM_COMP * BASIS_DIM, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASI…  in Grad()
131   SharedData_Cuda data;  in GradTranspose()  local
132   data.t_id_x = threadIdx.x;  in GradTranspose()
133   data.t_id_y = threadIdx.y;  in GradTranspose()
134   data.t_id_z = threadIdx.z;  in GradTranspose()
135   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in GradTranspose()
136   data.slice  = slice + data.t_id_z * BASIS_T_1D;  in GradTranspose()
143   LoadMatrix<BASIS_P, BASIS_Q * BASIS_DIM>(data, c_G, s_G);  in GradTranspose()
148 …ReadElementStrided1d<BASIS_NUM_COMP * BASIS_DIM, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS…  in GradTranspose()
149 …GradTransposeNonTensor<BASIS_NUM_COMP, BASIS_DIM, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_G, r_…  in GradTranspose()
150 …WriteElementStrided1d<BASIS_NUM_COMP, BASIS_P>(data, elem, 1, BASIS_P * num_elem, BASIS_P, r_V, d_…  in GradTranspose()
158   SharedData_Cuda data;  in GradTransposeAdd()  local
159   data.t_id_x = threadIdx.x;  in GradTransposeAdd()
160   data.t_id_y = threadIdx.y;  in GradTransposeAdd()
161   data.t_id_z = threadIdx.z;  in GradTransposeAdd()
162   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in GradTransposeAdd()
163   data.slice  = slice + data.t_id_z * BASIS_T_1D;  in GradTransposeAdd()
170   LoadMatrix<BASIS_P, BASIS_Q * BASIS_DIM>(data, c_G, s_G);  in GradTransposeAdd()
175 …ReadElementStrided1d<BASIS_NUM_COMP * BASIS_DIM, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS…  in GradTransposeAdd()
176 …GradTransposeNonTensor<BASIS_NUM_COMP, BASIS_DIM, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_G, r_…  in GradTransposeAdd()
177 …SumElementStrided1d<BASIS_NUM_COMP, BASIS_P>(data, elem, 1, BASIS_P * num_elem, BASIS_P, r_V, d_V);  in GradTransposeAdd()
187   SharedData_Cuda data;  in Weight()  local
188   data.t_id_x = threadIdx.x;  in Weight()
189   data.t_id_y = threadIdx.y;  in Weight()
190   data.t_id_z = threadIdx.z;  in Weight()
191   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in Weight()
192   data.slice  = slice + data.t_id_z * BASIS_T_1D;  in Weight()
197     WeightNonTensor<BASIS_P, BASIS_Q>(data, q_weight, r_W);  in Weight()
198     WriteElementStrided1d<1, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS_Q, r_W, d_W);  in Weight()