Lines Matching refs:data
29 CeedBasis_Hip *data; in CeedBasisApplyCore_Hip() local
32 CeedCallBackend(CeedBasisGetData(basis, &data)); in CeedBasisApplyCore_Hip()
51 …void *interp_args[] = {(void *)&num_elem, (void *)&is_transpose, &data->d_interp_1d, &d_u,… in CeedBasisApplyCore_Hip()
54 CeedCallBackend(CeedRunKernel_Hip(ceed, data->Interp, num_elem, block_size, interp_args)); in CeedBasisApplyCore_Hip()
57 … *grad_args[] = {(void *)&num_elem, (void *)&is_transpose, &data->d_interp_1d, &data->d_grad_1… in CeedBasisApplyCore_Hip()
60 CeedCallBackend(CeedRunKernel_Hip(ceed, data->Grad, num_elem, block_size, grad_args)); in CeedBasisApplyCore_Hip()
63 …CeedCheck(data->d_q_weight_1d, ceed, CEED_ERROR_BACKEND, "%s not supported; q_weights_1d not set",… in CeedBasisApplyCore_Hip()
64 void *weight_args[] = {(void *)&num_elem, (void *)&data->d_q_weight_1d, &d_v}; in CeedBasisApplyCore_Hip()
68 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Weight, num_elem, block_size_x, block_size_y, 1, … in CeedBasisApplyCore_Hip()
109 CeedBasis_Hip *data; in CeedBasisApplyAtPointsCore_Hip() local
111 CeedCallBackend(CeedBasisGetData(basis, &data)); in CeedBasisApplyAtPointsCore_Hip()
143 if (num_elem != data->num_elem_at_points) { in CeedBasisApplyAtPointsCore_Hip()
144 data->num_elem_at_points = num_elem; in CeedBasisApplyAtPointsCore_Hip()
146 if (data->d_points_per_elem) CeedCallHip(ceed, hipFree(data->d_points_per_elem)); in CeedBasisApplyAtPointsCore_Hip()
147 CeedCallHip(ceed, hipMalloc((void **)&data->d_points_per_elem, num_bytes)); in CeedBasisApplyAtPointsCore_Hip()
148 CeedCallBackend(CeedFree(&data->h_points_per_elem)); in CeedBasisApplyAtPointsCore_Hip()
149 CeedCallBackend(CeedCalloc(num_elem, &data->h_points_per_elem)); in CeedBasisApplyAtPointsCore_Hip()
151 if (memcmp(data->h_points_per_elem, num_points, num_bytes)) { in CeedBasisApplyAtPointsCore_Hip()
152 memcpy(data->h_points_per_elem, num_points, num_bytes); in CeedBasisApplyAtPointsCore_Hip()
153 …CeedCallHip(ceed, hipMemcpy(data->d_points_per_elem, num_points, num_bytes, hipMemcpyHostToDevice)… in CeedBasisApplyAtPointsCore_Hip()
158 if (data->num_points != max_num_points) { in CeedBasisApplyAtPointsCore_Hip()
162 data->num_points = max_num_points; in CeedBasisApplyAtPointsCore_Hip()
165 if (!data->d_chebyshev_interp_1d) { in CeedBasisApplyAtPointsCore_Hip()
172 CeedCallHip(ceed, hipMalloc((void **)&data->d_chebyshev_interp_1d, interp_bytes)); in CeedBasisApplyAtPointsCore_Hip()
173 …CeedCallHip(ceed, hipMemcpy(data->d_chebyshev_interp_1d, chebyshev_interp_1d, interp_bytes, hipMem… in CeedBasisApplyAtPointsCore_Hip()
181 if (data->moduleAtPoints) CeedCallHip(ceed, hipModuleUnload(data->moduleAtPoints)); in CeedBasisApplyAtPointsCore_Hip()
183 …CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->moduleAtPoints, 9, "BASIS_Q_1D",… in CeedBasisApplyAtPointsCore_Hip()
187 …CeedCallBackend(CeedGetKernel_Hip(ceed, data->moduleAtPoints, "InterpAtPoints", &data->InterpAtPoi… in CeedBasisApplyAtPointsCore_Hip()
188 …CeedCallBackend(CeedGetKernel_Hip(ceed, data->moduleAtPoints, "InterpTransposeAtPoints", &data->In… in CeedBasisApplyAtPointsCore_Hip()
189 …CeedCallBackend(CeedGetKernel_Hip(ceed, data->moduleAtPoints, "GradAtPoints", &data->GradAtPoints)… in CeedBasisApplyAtPointsCore_Hip()
190 …CeedCallBackend(CeedGetKernel_Hip(ceed, data->moduleAtPoints, "GradTransposeAtPoints", &data->Grad… in CeedBasisApplyAtPointsCore_Hip()
208 …void *interp_args[] = {(void *)&num_elem, &data->d_chebyshev_interp_1d, &data->d_points_pe… in CeedBasisApplyAtPointsCore_Hip()
211 …CeedCallBackend(CeedRunKernel_Hip(ceed, is_transpose ? data->InterpTransposeAtPoints : data->Inter… in CeedBasisApplyAtPointsCore_Hip()
215 …void *grad_args[] = {(void *)&num_elem, &data->d_chebyshev_interp_1d, &data->d_points_per_… in CeedBasisApplyAtPointsCore_Hip()
218 …CeedCallBackend(CeedRunKernel_Hip(ceed, is_transpose ? data->GradTransposeAtPoints : data->GradAtP… in CeedBasisApplyAtPointsCore_Hip()
263 CeedBasisNonTensor_Hip *data; in CeedBasisApplyNonTensorCore_Hip() local
266 CeedCallBackend(CeedBasisGetData(basis, &data)); in CeedBasisApplyNonTensorCore_Hip()
284 void *interp_args[] = {(void *)&num_elem, &data->d_interp, &d_u, &d_v}; in CeedBasisApplyNonTensorCore_Hip()
288 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->InterpTranspose, grid, block_size_x, 1, elems_per… in CeedBasisApplyNonTensorCore_Hip()
290 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Interp, grid, block_size_x, 1, elems_per_block, i… in CeedBasisApplyNonTensorCore_Hip()
294 void *grad_args[] = {(void *)&num_elem, &data->d_grad, &d_u, &d_v}; in CeedBasisApplyNonTensorCore_Hip()
298 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->DerivTranspose, grid, block_size_x, 1, elems_per_… in CeedBasisApplyNonTensorCore_Hip()
300 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Deriv, grid, block_size_x, 1, elems_per_block, gr… in CeedBasisApplyNonTensorCore_Hip()
304 void *div_args[] = {(void *)&num_elem, &data->d_div, &d_u, &d_v}; in CeedBasisApplyNonTensorCore_Hip()
308 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->DerivTranspose, grid, block_size_x, 1, elems_per_… in CeedBasisApplyNonTensorCore_Hip()
310 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Deriv, grid, block_size_x, 1, elems_per_block, di… in CeedBasisApplyNonTensorCore_Hip()
314 void *curl_args[] = {(void *)&num_elem, &data->d_curl, &d_u, &d_v}; in CeedBasisApplyNonTensorCore_Hip()
318 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->DerivTranspose, grid, block_size_x, 1, elems_per_… in CeedBasisApplyNonTensorCore_Hip()
320 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Deriv, grid, block_size_x, 1, elems_per_block, cu… in CeedBasisApplyNonTensorCore_Hip()
324 …CeedCheck(data->d_q_weight, ceed, CEED_ERROR_BACKEND, "%s not supported; q_weights not set", CeedE… in CeedBasisApplyNonTensorCore_Hip()
325 void *weight_args[] = {(void *)&num_elem, (void *)&data->d_q_weight, &d_v}; in CeedBasisApplyNonTensorCore_Hip()
327 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Weight, grid, num_qpts, 1, elems_per_block, weigh… in CeedBasisApplyNonTensorCore_Hip()
358 CeedBasis_Hip *data; in CeedBasisDestroy_Hip() local
361 CeedCallBackend(CeedBasisGetData(basis, &data)); in CeedBasisDestroy_Hip()
362 CeedCallHip(ceed, hipModuleUnload(data->module)); in CeedBasisDestroy_Hip()
363 if (data->moduleAtPoints) CeedCallHip(ceed, hipModuleUnload(data->moduleAtPoints)); in CeedBasisDestroy_Hip()
364 if (data->d_q_weight_1d) CeedCallHip(ceed, hipFree(data->d_q_weight_1d)); in CeedBasisDestroy_Hip()
365 CeedCallBackend(CeedFree(&data->h_points_per_elem)); in CeedBasisDestroy_Hip()
366 if (data->d_points_per_elem) CeedCallHip(ceed, hipFree(data->d_points_per_elem)); in CeedBasisDestroy_Hip()
367 CeedCallHip(ceed, hipFree(data->d_interp_1d)); in CeedBasisDestroy_Hip()
368 CeedCallHip(ceed, hipFree(data->d_grad_1d)); in CeedBasisDestroy_Hip()
369 CeedCallHip(ceed, hipFree(data->d_chebyshev_interp_1d)); in CeedBasisDestroy_Hip()
370 CeedCallBackend(CeedFree(&data)); in CeedBasisDestroy_Hip()
380 CeedBasisNonTensor_Hip *data; in CeedBasisDestroyNonTensor_Hip() local
383 CeedCallBackend(CeedBasisGetData(basis, &data)); in CeedBasisDestroyNonTensor_Hip()
384 CeedCallHip(ceed, hipModuleUnload(data->module)); in CeedBasisDestroyNonTensor_Hip()
385 if (data->d_q_weight) CeedCallHip(ceed, hipFree(data->d_q_weight)); in CeedBasisDestroyNonTensor_Hip()
386 CeedCallHip(ceed, hipFree(data->d_interp)); in CeedBasisDestroyNonTensor_Hip()
387 CeedCallHip(ceed, hipFree(data->d_grad)); in CeedBasisDestroyNonTensor_Hip()
388 CeedCallHip(ceed, hipFree(data->d_div)); in CeedBasisDestroyNonTensor_Hip()
389 CeedCallHip(ceed, hipFree(data->d_curl)); in CeedBasisDestroyNonTensor_Hip()
390 CeedCallBackend(CeedFree(&data)); in CeedBasisDestroyNonTensor_Hip()
404 CeedBasis_Hip *data; in CeedBasisCreateTensorH1_Hip() local
407 CeedCallBackend(CeedCalloc(1, &data)); in CeedBasisCreateTensorH1_Hip()
411 CeedCallHip(ceed, hipMalloc((void **)&data->d_q_weight_1d, q_bytes)); in CeedBasisCreateTensorH1_Hip()
412 CeedCallHip(ceed, hipMemcpy(data->d_q_weight_1d, q_weight_1d, q_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateTensorH1_Hip()
414 CeedCallHip(ceed, hipMalloc((void **)&data->d_interp_1d, interp_bytes)); in CeedBasisCreateTensorH1_Hip()
415 CeedCallHip(ceed, hipMemcpy(data->d_interp_1d, interp_1d, interp_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateTensorH1_Hip()
416 CeedCallHip(ceed, hipMalloc((void **)&data->d_grad_1d, interp_bytes)); in CeedBasisCreateTensorH1_Hip()
417 CeedCallHip(ceed, hipMemcpy(data->d_grad_1d, grad_1d, interp_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateTensorH1_Hip()
423 …CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->module, 7, "BASIS_Q_1D", Q_1d, "… in CeedBasisCreateTensorH1_Hip()
426 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Interp", &data->Interp)); in CeedBasisCreateTensorH1_Hip()
427 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Grad", &data->Grad)); in CeedBasisCreateTensorH1_Hip()
428 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Weight", &data->Weight)); in CeedBasisCreateTensorH1_Hip()
430 CeedCallBackend(CeedBasisSetData(basis, data)); in CeedBasisCreateTensorH1_Hip()
450 CeedBasisNonTensor_Hip *data; in CeedBasisCreateH1_Hip() local
453 CeedCallBackend(CeedCalloc(1, &data)); in CeedBasisCreateH1_Hip()
459 CeedCallHip(ceed, hipMalloc((void **)&data->d_q_weight, q_bytes)); in CeedBasisCreateH1_Hip()
460 CeedCallHip(ceed, hipMemcpy(data->d_q_weight, q_weight, q_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateH1_Hip()
465 CeedCallHip(ceed, hipMalloc((void **)&data->d_interp, interp_bytes)); in CeedBasisCreateH1_Hip()
466 CeedCallHip(ceed, hipMemcpy(data->d_interp, interp, interp_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateH1_Hip()
471 CeedCallHip(ceed, hipMalloc((void **)&data->d_grad, grad_bytes)); in CeedBasisCreateH1_Hip()
472 CeedCallHip(ceed, hipMemcpy(data->d_grad, grad, grad_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateH1_Hip()
479 …CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->module, 5, "BASIS_Q", num_qpts, … in CeedBasisCreateH1_Hip()
481 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Interp", &data->Interp)); in CeedBasisCreateH1_Hip()
482 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "InterpTranspose", &data->InterpTranspose)); in CeedBasisCreateH1_Hip()
483 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Deriv", &data->Deriv)); in CeedBasisCreateH1_Hip()
484 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "DerivTranspose", &data->DerivTranspose)); in CeedBasisCreateH1_Hip()
485 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Weight", &data->Weight)); in CeedBasisCreateH1_Hip()
487 CeedCallBackend(CeedBasisSetData(basis, data)); in CeedBasisCreateH1_Hip()
505 CeedBasisNonTensor_Hip *data; in CeedBasisCreateHdiv_Hip() local
508 CeedCallBackend(CeedCalloc(1, &data)); in CeedBasisCreateHdiv_Hip()
514 CeedCallHip(ceed, hipMalloc((void **)&data->d_q_weight, q_bytes)); in CeedBasisCreateHdiv_Hip()
515 CeedCallHip(ceed, hipMemcpy(data->d_q_weight, q_weight, q_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateHdiv_Hip()
520 CeedCallHip(ceed, hipMalloc((void **)&data->d_interp, interp_bytes)); in CeedBasisCreateHdiv_Hip()
521 CeedCallHip(ceed, hipMemcpy(data->d_interp, interp, interp_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateHdiv_Hip()
526 CeedCallHip(ceed, hipMalloc((void **)&data->d_div, div_bytes)); in CeedBasisCreateHdiv_Hip()
527 CeedCallHip(ceed, hipMemcpy(data->d_div, div, div_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateHdiv_Hip()
534 …CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->module, 5, "BASIS_Q", num_qpts, … in CeedBasisCreateHdiv_Hip()
536 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Interp", &data->Interp)); in CeedBasisCreateHdiv_Hip()
537 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "InterpTranspose", &data->InterpTranspose)); in CeedBasisCreateHdiv_Hip()
538 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Deriv", &data->Deriv)); in CeedBasisCreateHdiv_Hip()
539 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "DerivTranspose", &data->DerivTranspose)); in CeedBasisCreateHdiv_Hip()
540 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Weight", &data->Weight)); in CeedBasisCreateHdiv_Hip()
542 CeedCallBackend(CeedBasisSetData(basis, data)); in CeedBasisCreateHdiv_Hip()
560 CeedBasisNonTensor_Hip *data; in CeedBasisCreateHcurl_Hip() local
563 CeedCallBackend(CeedCalloc(1, &data)); in CeedBasisCreateHcurl_Hip()
569 CeedCallHip(ceed, hipMalloc((void **)&data->d_q_weight, q_bytes)); in CeedBasisCreateHcurl_Hip()
570 CeedCallHip(ceed, hipMemcpy(data->d_q_weight, q_weight, q_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateHcurl_Hip()
575 CeedCallHip(ceed, hipMalloc((void **)&data->d_interp, interp_bytes)); in CeedBasisCreateHcurl_Hip()
576 CeedCallHip(ceed, hipMemcpy(data->d_interp, interp, interp_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateHcurl_Hip()
581 CeedCallHip(ceed, hipMalloc((void **)&data->d_curl, curl_bytes)); in CeedBasisCreateHcurl_Hip()
582 CeedCallHip(ceed, hipMemcpy(data->d_curl, curl, curl_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateHcurl_Hip()
589 …CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->module, 5, "BASIS_Q", num_qpts, … in CeedBasisCreateHcurl_Hip()
591 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Interp", &data->Interp)); in CeedBasisCreateHcurl_Hip()
592 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "InterpTranspose", &data->InterpTranspose)); in CeedBasisCreateHcurl_Hip()
593 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Deriv", &data->Deriv)); in CeedBasisCreateHcurl_Hip()
594 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "DerivTranspose", &data->DerivTranspose)); in CeedBasisCreateHcurl_Hip()
595 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Weight", &data->Weight)); in CeedBasisCreateHcurl_Hip()
597 CeedCallBackend(CeedBasisSetData(basis, data)); in CeedBasisCreateHcurl_Hip()