88   if (impl->asmb) {  in CeedOperatorDestroy_Hip()
92     CeedCallHip(ceed, hipModuleUnload(impl->asmb->module));  in CeedOperatorDestroy_Hip()
93     CeedCallHip(ceed, hipFree(impl->asmb->d_B_in));  in CeedOperatorDestroy_Hip()
94     CeedCallHip(ceed, hipFree(impl->asmb->d_B_out));  in CeedOperatorDestroy_Hip()
97   CeedCallBackend(CeedFree(&impl->asmb));  in CeedOperatorDestroy_Hip()
1601   CeedCallBackend(CeedCalloc(1, &impl->asmb));  in CeedOperatorAssembleSingleSetup_Hip()
1602   CeedOperatorAssemble_Hip *asmb = impl->asmb;  in CeedOperatorAssembleSingleSetup_Hip()  local
1603   asmb->elems_per_block          = 1;  in CeedOperatorAssembleSingleSetup_Hip()
1604   asmb->block_size_x             = elem_size_in;  in CeedOperatorAssembleSingleSetup_Hip()
1605   asmb->block_size_y             = elem_size_out;  in CeedOperatorAssembleSingleSetup_Hip()
1608 …bool fallback = asmb->block_size_x * asmb->block_size_y * asmb->elems_per_block > hip_data->device…  in CeedOperatorAssembleSingleSetup_Hip()
1612     asmb->block_size_y = 1;  in CeedOperatorAssembleSingleSetup_Hip()
1620 …CeedCallBackend(CeedCompile_Hip(ceed, assembly_kernel_source, &asmb->module, 10, "NUM_EVAL_MODES_I…  in CeedOperatorAssembleSingleSetup_Hip()
1623 …asmb->block_size_x * asmb->block_size_y * asmb->elems_per_block, "BLOCK_SIZE_Y", asmb->block_size_…  in CeedOperatorAssembleSingleSetup_Hip()
1625   CeedCallBackend(CeedGetKernel_Hip(ceed, asmb->module, "LinearAssemble", &asmb->LinearAssemble));  in CeedOperatorAssembleSingleSetup_Hip()
1643     CeedCallHip(ceed, hipMalloc((void **)&asmb->d_B_in, in_bytes));  in CeedOperatorAssembleSingleSetup_Hip()
1655 …CeedCallHip(ceed, hipMemcpy(&asmb->d_B_in[i * elem_size_in * num_qpts_in], h_B_in, elem_size_in * …  in CeedOperatorAssembleSingleSetup_Hip()
1678     CeedCallHip(ceed, hipMalloc((void **)&asmb->d_B_out, out_bytes));  in CeedOperatorAssembleSingleSetup_Hip()
1690 …CeedCallHip(ceed, hipMemcpy(&asmb->d_B_out[i * elem_size_out * num_qpts_out], h_B_out, elem_size_o…  in CeedOperatorAssembleSingleSetup_Hip()
1739   if (!impl->asmb) CeedCallBackend(CeedOperatorAssembleSingleSetup_Hip(op, use_ceedsize_idx));  in CeedOperatorAssembleSingle_Hip()
1740   CeedOperatorAssemble_Hip *asmb = impl->asmb;  in CeedOperatorAssembleSingle_Hip()  local
1742   assert(asmb != NULL);  in CeedOperatorAssembleSingle_Hip()
1779 …t ? elem_size_in * elem_size_out : 0) + (curl_orients_in ? elem_size_in * asmb->block_size_y : 0))…  in CeedOperatorAssembleSingle_Hip()
1781   CeedInt grid   = CeedDivUpInt(num_elem_in, asmb->elems_per_block);  in CeedOperatorAssembleSingle_Hip()
1782 …void   *args[] = {(void *)&num_elem_in, &asmb->d_B_in,     &asmb->d_B_out,      &orients_in,  &cur…  in CeedOperatorAssembleSingle_Hip()
1785 …end(CeedRunKernelDimShared_Hip(ceed, asmb->LinearAssemble, NULL, grid, asmb->block_size_x, asmb->b…  in CeedOperatorAssembleSingle_Hip()