Lines Matching refs:data
98 CeedBasis_Hip_shared *data; in CeedBasisApplyTensorCore_Hip_shared() local
102 CeedCallBackend(CeedBasisGetData(basis, &data)); in CeedBasisApplyTensorCore_Hip_shared()
119 CeedInt block_size = data->block_sizes[0]; in CeedBasisApplyTensorCore_Hip_shared()
121 …CeedCheck(data->d_interp_1d, ceed, CEED_ERROR_BACKEND, "%s not supported; interp_1d not set", Ceed… in CeedBasisApplyTensorCore_Hip_shared()
125 void *interp_args[] = {(void *)&num_elem, &data->d_interp_1d, &d_u, &d_v}; in CeedBasisApplyTensorCore_Hip_shared()
134 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, apply_add ? data->InterpTransposeAdd : data->Inte… in CeedBasisApplyTensorCore_Hip_shared()
137 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, data->Interp, NULL, grid, thread_1d, 1, elems_per… in CeedBasisApplyTensorCore_Hip_shared()
146 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, apply_add ? data->InterpTransposeAdd : data->Inte… in CeedBasisApplyTensorCore_Hip_shared()
149 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, data->Interp, NULL, grid, thread_1d, thread_1d, e… in CeedBasisApplyTensorCore_Hip_shared()
157 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, apply_add ? data->InterpTransposeAdd : data->Inte… in CeedBasisApplyTensorCore_Hip_shared()
160 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, data->Interp, NULL, grid, thread_1d, thread_1d, e… in CeedBasisApplyTensorCore_Hip_shared()
166 CeedInt block_size = data->block_sizes[1]; in CeedBasisApplyTensorCore_Hip_shared()
168 …CeedCheck(data->d_grad_1d, ceed, CEED_ERROR_BACKEND, "%s not supported; grad_1d not set", CeedEval… in CeedBasisApplyTensorCore_Hip_shared()
172 CeedScalar *d_grad_1d = data->d_grad_1d; in CeedBasisApplyTensorCore_Hip_shared()
174 if (data->d_collo_grad_1d) { in CeedBasisApplyTensorCore_Hip_shared()
175 d_grad_1d = data->d_collo_grad_1d; in CeedBasisApplyTensorCore_Hip_shared()
177 void *grad_args[] = {(void *)&num_elem, &data->d_interp_1d, &d_grad_1d, &d_u, &d_v}; in CeedBasisApplyTensorCore_Hip_shared()
186 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, apply_add ? data->GradTransposeAdd : data->GradTr… in CeedBasisApplyTensorCore_Hip_shared()
189 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, data->Grad, NULL, grid, thread_1d, 1, elems_per_b… in CeedBasisApplyTensorCore_Hip_shared()
198 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, apply_add ? data->GradTransposeAdd : data->GradTr… in CeedBasisApplyTensorCore_Hip_shared()
201 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, data->Grad, NULL, grid, thread_1d, thread_1d, ele… in CeedBasisApplyTensorCore_Hip_shared()
209 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, apply_add ? data->GradTransposeAdd : data->GradTr… in CeedBasisApplyTensorCore_Hip_shared()
212 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, data->Grad, NULL, grid, thread_1d, thread_1d, ele… in CeedBasisApplyTensorCore_Hip_shared()
218 CeedInt block_size = data->block_sizes[2]; in CeedBasisApplyTensorCore_Hip_shared()
220 …CeedCheck(data->d_q_weight_1d, ceed, CEED_ERROR_BACKEND, "%s not supported; q_weights_1d not set",… in CeedBasisApplyTensorCore_Hip_shared()
222 void *weight_args[] = {(void *)&num_elem, (void *)&data->d_q_weight_1d, &d_v}; in CeedBasisApplyTensorCore_Hip_shared()
229 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Weight, grid_size, Q_1d, elems_per_block, 1, weig… in CeedBasisApplyTensorCore_Hip_shared()
235 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Weight, grid_size, Q_1d, Q_1d, elems_per_block, w… in CeedBasisApplyTensorCore_Hip_shared()
241 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Weight, grid_size, Q_1d, Q_1d, elems_per_block, w… in CeedBasisApplyTensorCore_Hip_shared()
283 CeedBasis_Hip_shared *data; in CeedBasisApplyAtPointsCore_Hip_shared() local
285 CeedCallBackend(CeedBasisGetData(basis, &data)); in CeedBasisApplyAtPointsCore_Hip_shared()
317 if (num_elem != data->num_elem_at_points) { in CeedBasisApplyAtPointsCore_Hip_shared()
318 data->num_elem_at_points = num_elem; in CeedBasisApplyAtPointsCore_Hip_shared()
320 if (data->d_points_per_elem) CeedCallHip(ceed, hipFree(data->d_points_per_elem)); in CeedBasisApplyAtPointsCore_Hip_shared()
321 CeedCallHip(ceed, hipMalloc((void **)&data->d_points_per_elem, num_bytes)); in CeedBasisApplyAtPointsCore_Hip_shared()
322 CeedCallBackend(CeedFree(&data->h_points_per_elem)); in CeedBasisApplyAtPointsCore_Hip_shared()
323 CeedCallBackend(CeedCalloc(num_elem, &data->h_points_per_elem)); in CeedBasisApplyAtPointsCore_Hip_shared()
325 if (memcmp(data->h_points_per_elem, num_points, num_bytes)) { in CeedBasisApplyAtPointsCore_Hip_shared()
326 memcpy(data->h_points_per_elem, num_points, num_bytes); in CeedBasisApplyAtPointsCore_Hip_shared()
327 …CeedCallHip(ceed, hipMemcpy(data->d_points_per_elem, num_points, num_bytes, hipMemcpyHostToDevice)… in CeedBasisApplyAtPointsCore_Hip_shared()
332 if (data->num_points != max_num_points) { in CeedBasisApplyAtPointsCore_Hip_shared()
336 data->num_points = max_num_points; in CeedBasisApplyAtPointsCore_Hip_shared()
339 if (!data->d_chebyshev_interp_1d) { in CeedBasisApplyAtPointsCore_Hip_shared()
346 CeedCallHip(ceed, hipMalloc((void **)&data->d_chebyshev_interp_1d, interp_bytes)); in CeedBasisApplyAtPointsCore_Hip_shared()
347 …CeedCallHip(ceed, hipMemcpy(data->d_chebyshev_interp_1d, chebyshev_interp_1d, interp_bytes, hipMem… in CeedBasisApplyAtPointsCore_Hip_shared()
355 if (data->moduleAtPoints) CeedCallHip(ceed, hipModuleUnload(data->moduleAtPoints)); in CeedBasisApplyAtPointsCore_Hip_shared()
357 …CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->moduleAtPoints, 9, "BASIS_Q_1D",… in CeedBasisApplyAtPointsCore_Hip_shared()
360 data->block_sizes[0])); in CeedBasisApplyAtPointsCore_Hip_shared()
361 …CeedCallBackend(CeedGetKernel_Hip(ceed, data->moduleAtPoints, "InterpAtPoints", &data->InterpAtPoi… in CeedBasisApplyAtPointsCore_Hip_shared()
362 …CeedCallBackend(CeedGetKernel_Hip(ceed, data->moduleAtPoints, "InterpTransposeAtPoints", &data->In… in CeedBasisApplyAtPointsCore_Hip_shared()
363 …CeedCallBackend(CeedGetKernel_Hip(ceed, data->moduleAtPoints, "InterpTransposeAddAtPoints", &data-… in CeedBasisApplyAtPointsCore_Hip_shared()
364 …CeedCallBackend(CeedGetKernel_Hip(ceed, data->moduleAtPoints, "GradAtPoints", &data->GradAtPoints)… in CeedBasisApplyAtPointsCore_Hip_shared()
365 …CeedCallBackend(CeedGetKernel_Hip(ceed, data->moduleAtPoints, "GradTransposeAtPoints", &data->Grad… in CeedBasisApplyAtPointsCore_Hip_shared()
366 …CeedCallBackend(CeedGetKernel_Hip(ceed, data->moduleAtPoints, "GradTransposeAddAtPoints", &data->G… in CeedBasisApplyAtPointsCore_Hip_shared()
383 CeedInt block_size = data->block_sizes[0]; in CeedBasisApplyAtPointsCore_Hip_shared()
388 …void *interp_args[] = {(void *)&num_elem, &data->d_chebyshev_interp_1d, &data->d_points_per_elem… in CeedBasisApplyAtPointsCore_Hip_shared()
397 …CallBackend(CeedRunKernelDimShared_Hip(ceed, apply_add ? data->InterpTransposeAddAtPoints : data->… in CeedBasisApplyAtPointsCore_Hip_shared()
400 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, data->InterpAtPoints, NULL, grid, thread_1d, 1, e… in CeedBasisApplyAtPointsCore_Hip_shared()
409 …CallBackend(CeedRunKernelDimShared_Hip(ceed, apply_add ? data->InterpTransposeAddAtPoints : data->… in CeedBasisApplyAtPointsCore_Hip_shared()
412 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, data->InterpAtPoints, NULL, grid, thread_1d, thre… in CeedBasisApplyAtPointsCore_Hip_shared()
421 …CallBackend(CeedRunKernelDimShared_Hip(ceed, apply_add ? data->InterpTransposeAddAtPoints : data->… in CeedBasisApplyAtPointsCore_Hip_shared()
424 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, data->InterpAtPoints, NULL, grid, thread_1d, thre… in CeedBasisApplyAtPointsCore_Hip_shared()
431 CeedInt block_size = data->block_sizes[0]; in CeedBasisApplyAtPointsCore_Hip_shared()
436 …void *grad_args[] = {(void *)&num_elem, &data->d_chebyshev_interp_1d, &data->d_points_per_elem, … in CeedBasisApplyAtPointsCore_Hip_shared()
445 …dCallBackend(CeedRunKernelDimShared_Hip(ceed, apply_add ? data->GradTransposeAddAtPoints : data->G… in CeedBasisApplyAtPointsCore_Hip_shared()
448 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, data->GradAtPoints, NULL, grid, thread_1d, 1, ele… in CeedBasisApplyAtPointsCore_Hip_shared()
457 …dCallBackend(CeedRunKernelDimShared_Hip(ceed, apply_add ? data->GradTransposeAddAtPoints : data->G… in CeedBasisApplyAtPointsCore_Hip_shared()
460 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, data->GradAtPoints, NULL, grid, thread_1d, thread… in CeedBasisApplyAtPointsCore_Hip_shared()
469 …dCallBackend(CeedRunKernelDimShared_Hip(ceed, apply_add ? data->GradTransposeAddAtPoints : data->G… in CeedBasisApplyAtPointsCore_Hip_shared()
472 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, data->GradAtPoints, NULL, grid, thread_1d, thread… in CeedBasisApplyAtPointsCore_Hip_shared()
518 CeedBasis_Hip_shared *data; in CeedBasisApplyNonTensorCore_Hip_shared() local
522 CeedCallBackend(CeedBasisGetData(basis, &data)); in CeedBasisApplyNonTensorCore_Hip_shared()
540 …CeedCheck(data->d_interp_1d, ceed, CEED_ERROR_BACKEND, "%s not supported; interp not set", CeedEva… in CeedBasisApplyNonTensorCore_Hip_shared()
544 void *interp_args[] = {(void *)&num_elem, &data->d_interp_1d, &d_u, &d_v}; in CeedBasisApplyNonTensorCore_Hip_shared()
553 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, apply_add ? data->InterpTransposeAdd : data->Inte… in CeedBasisApplyNonTensorCore_Hip_shared()
556 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, data->Interp, NULL, grid, thread, 1, elems_per_bl… in CeedBasisApplyNonTensorCore_Hip_shared()
563 …CeedCheck(data->d_grad_1d, ceed, CEED_ERROR_BACKEND, "%s not supported; grad not set", CeedEvalMod… in CeedBasisApplyNonTensorCore_Hip_shared()
567 void *grad_args[] = {(void *)&num_elem, &data->d_grad_1d, &d_u, &d_v}; in CeedBasisApplyNonTensorCore_Hip_shared()
576 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, apply_add ? data->GradTransposeAdd : data->GradTr… in CeedBasisApplyNonTensorCore_Hip_shared()
579 …CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, data->Grad, NULL, grid, thread, 1, elems_per_bloc… in CeedBasisApplyNonTensorCore_Hip_shared()
586 …CeedCheck(data->d_q_weight_1d, ceed, CEED_ERROR_BACKEND, "%s not supported; q_weights not set", Ce… in CeedBasisApplyNonTensorCore_Hip_shared()
590 void *weight_args[] = {(void *)&num_elem, (void *)&data->d_q_weight_1d, &d_v}; in CeedBasisApplyNonTensorCore_Hip_shared()
597 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Weight, grid_size, thread, elems_per_block, 1, we… in CeedBasisApplyNonTensorCore_Hip_shared()
634 CeedBasis_Hip_shared *data; in CeedBasisDestroy_Hip_shared() local
637 CeedCallBackend(CeedBasisGetData(basis, &data)); in CeedBasisDestroy_Hip_shared()
638 CeedCallHip(ceed, hipModuleUnload(data->module)); in CeedBasisDestroy_Hip_shared()
639 if (data->moduleAtPoints) CeedCallHip(ceed, hipModuleUnload(data->moduleAtPoints)); in CeedBasisDestroy_Hip_shared()
640 if (data->d_q_weight_1d) CeedCallHip(ceed, hipFree(data->d_q_weight_1d)); in CeedBasisDestroy_Hip_shared()
641 CeedCallBackend(CeedFree(&data->h_points_per_elem)); in CeedBasisDestroy_Hip_shared()
642 if (data->d_points_per_elem) CeedCallHip(ceed, hipFree(data->d_points_per_elem)); in CeedBasisDestroy_Hip_shared()
643 CeedCallHip(ceed, hipFree(data->d_interp_1d)); in CeedBasisDestroy_Hip_shared()
644 CeedCallHip(ceed, hipFree(data->d_grad_1d)); in CeedBasisDestroy_Hip_shared()
645 CeedCallHip(ceed, hipFree(data->d_collo_grad_1d)); in CeedBasisDestroy_Hip_shared()
646 CeedCallHip(ceed, hipFree(data->d_chebyshev_interp_1d)); in CeedBasisDestroy_Hip_shared()
647 CeedCallBackend(CeedFree(&data)); in CeedBasisDestroy_Hip_shared()
661 CeedBasis_Hip_shared *data; in CeedBasisCreateTensorH1_Hip_shared() local
664 CeedCallBackend(CeedCalloc(1, &data)); in CeedBasisCreateTensorH1_Hip_shared()
668 CeedCallHip(ceed, hipMalloc((void **)&data->d_q_weight_1d, q_bytes)); in CeedBasisCreateTensorH1_Hip_shared()
669 CeedCallHip(ceed, hipMemcpy(data->d_q_weight_1d, q_weight_1d, q_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateTensorH1_Hip_shared()
671 CeedCallHip(ceed, hipMalloc((void **)&data->d_interp_1d, interp_bytes)); in CeedBasisCreateTensorH1_Hip_shared()
672 CeedCallHip(ceed, hipMemcpy(data->d_interp_1d, interp_1d, interp_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateTensorH1_Hip_shared()
673 CeedCallHip(ceed, hipMalloc((void **)&data->d_grad_1d, interp_bytes)); in CeedBasisCreateTensorH1_Hip_shared()
674 CeedCallHip(ceed, hipMemcpy(data->d_grad_1d, grad_1d, interp_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateTensorH1_Hip_shared()
677 data->d_collo_grad_1d = NULL; in CeedBasisCreateTensorH1_Hip_shared()
685 CeedCallHip(ceed, hipMalloc((void **)&data->d_collo_grad_1d, q_bytes * Q_1d)); in CeedBasisCreateTensorH1_Hip_shared()
686 …CeedCallHip(ceed, hipMemcpy(data->d_collo_grad_1d, collo_grad_1d, q_bytes * Q_1d, hipMemcpyHostToD… in CeedBasisCreateTensorH1_Hip_shared()
692 CeedCallBackend(ComputeBasisThreadBlockSizes(dim, P_1d, Q_1d, num_comp, data->block_sizes)); in CeedBasisCreateTensorH1_Hip_shared()
698 …CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->module, 11, "BASIS_Q_1D", Q_1d, … in CeedBasisCreateTensorH1_Hip_shared()
700 …"BASIS_NUM_QPTS", CeedIntPow(Q_1d, dim), "BASIS_INTERP_BLOCK_SIZE", data->block_sizes[0], "BASIS_G… in CeedBasisCreateTensorH1_Hip_shared()
701 …data->block_sizes[1], "BASIS_WEIGHT_BLOCK_SIZE", data->block_sizes[2], "BASIS_HAS_COLLOCATED_GRAD", in CeedBasisCreateTensorH1_Hip_shared()
704 …CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, is_collocated ? "InterpCollocated" : "Interp… in CeedBasisCreateTensorH1_Hip_shared()
705 …llBackend(CeedGetKernel_Hip(ceed, data->module, is_collocated ? "InterpCollocatedTranspose" : "Int… in CeedBasisCreateTensorH1_Hip_shared()
706 …CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, is_collocated ? "InterpCollocatedTransposeAd… in CeedBasisCreateTensorH1_Hip_shared()
707 &data->InterpTransposeAdd)); in CeedBasisCreateTensorH1_Hip_shared()
708 …CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, is_collocated ? "GradCollocated" : "Grad", &… in CeedBasisCreateTensorH1_Hip_shared()
709 …CallBackend(CeedGetKernel_Hip(ceed, data->module, is_collocated ? "GradCollocatedTranspose" : "Gra… in CeedBasisCreateTensorH1_Hip_shared()
710 …lBackend(CeedGetKernel_Hip(ceed, data->module, is_collocated ? "GradCollocatedTransposeAdd" : "Gra… in CeedBasisCreateTensorH1_Hip_shared()
711 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Weight", &data->Weight)); in CeedBasisCreateTensorH1_Hip_shared()
713 CeedCallBackend(CeedBasisSetData(basis, data)); in CeedBasisCreateTensorH1_Hip_shared()
733 CeedBasis_Hip_shared *data; in CeedBasisCreateH1_Hip_shared() local
750 CeedCallBackend(CeedCalloc(1, &data)); in CeedBasisCreateH1_Hip_shared()
756 CeedCallHip(ceed, hipMalloc((void **)&data->d_q_weight_1d, q_bytes)); in CeedBasisCreateH1_Hip_shared()
757 CeedCallHip(ceed, hipMemcpy(data->d_q_weight_1d, q_weight, q_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateH1_Hip_shared()
762 CeedCallHip(ceed, hipMalloc((void **)&data->d_interp_1d, interp_bytes)); in CeedBasisCreateH1_Hip_shared()
763 CeedCallHip(ceed, hipMemcpy(data->d_interp_1d, interp, interp_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateH1_Hip_shared()
768 CeedCallHip(ceed, hipMalloc((void **)&data->d_grad_1d, grad_bytes)); in CeedBasisCreateH1_Hip_shared()
769 CeedCallHip(ceed, hipMemcpy(data->d_grad_1d, grad, grad_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateH1_Hip_shared()
776 …CeedCallBackend(ComputeBasisThreadBlockSizes(dim, num_nodes, num_qpts, num_comp, data->block_sizes… in CeedBasisCreateH1_Hip_shared()
777 …CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->module, 6, "BASIS_Q", num_qpts, … in CeedBasisCreateH1_Hip_shared()
779 data->block_sizes[0])); in CeedBasisCreateH1_Hip_shared()
780 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Interp", &data->Interp)); in CeedBasisCreateH1_Hip_shared()
781 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "InterpTranspose", &data->InterpTranspose)); in CeedBasisCreateH1_Hip_shared()
782 …CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "InterpTransposeAdd", &data->InterpTranspose… in CeedBasisCreateH1_Hip_shared()
783 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Grad", &data->Grad)); in CeedBasisCreateH1_Hip_shared()
784 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "GradTranspose", &data->GradTranspose)); in CeedBasisCreateH1_Hip_shared()
785 …CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "GradTransposeAdd", &data->GradTransposeAdd)… in CeedBasisCreateH1_Hip_shared()
786 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Weight", &data->Weight)); in CeedBasisCreateH1_Hip_shared()
788 CeedCallBackend(CeedBasisSetData(basis, data)); in CeedBasisCreateH1_Hip_shared()