89   if (impl->asmb) {  in CeedOperatorDestroy_Cuda()
93     CeedCallCuda(ceed, cuModuleUnload(impl->asmb->module));  in CeedOperatorDestroy_Cuda()
94     CeedCallCuda(ceed, cudaFree(impl->asmb->d_B_in));  in CeedOperatorDestroy_Cuda()
95     CeedCallCuda(ceed, cudaFree(impl->asmb->d_B_out));  in CeedOperatorDestroy_Cuda()
98   CeedCallBackend(CeedFree(&impl->asmb));  in CeedOperatorDestroy_Cuda()
1604   CeedCallBackend(CeedCalloc(1, &impl->asmb));  in CeedOperatorAssembleSingleSetup_Cuda()
1605   CeedOperatorAssemble_Cuda *asmb = impl->asmb;  in CeedOperatorAssembleSingleSetup_Cuda()  local
1606   asmb->elems_per_block           = 1;  in CeedOperatorAssembleSingleSetup_Cuda()
1607   asmb->block_size_x              = elem_size_in;  in CeedOperatorAssembleSingleSetup_Cuda()
1608   asmb->block_size_y              = elem_size_out;  in CeedOperatorAssembleSingleSetup_Cuda()
1611 …bool fallback = asmb->block_size_x * asmb->block_size_y * asmb->elems_per_block > cuda_data->devic…  in CeedOperatorAssembleSingleSetup_Cuda()
1615     asmb->block_size_y = 1;  in CeedOperatorAssembleSingleSetup_Cuda()
1623 …CeedCallBackend(CeedCompile_Cuda(ceed, assembly_kernel_source, &asmb->module, 10, "NUM_EVAL_MODES_…  in CeedOperatorAssembleSingleSetup_Cuda()
1626 …asmb->block_size_x * asmb->block_size_y * asmb->elems_per_block, "BLOCK_SIZE_Y", asmb->block_size_…  in CeedOperatorAssembleSingleSetup_Cuda()
1628   CeedCallBackend(CeedGetKernel_Cuda(ceed, asmb->module, "LinearAssemble", &asmb->LinearAssemble));  in CeedOperatorAssembleSingleSetup_Cuda()
1646     CeedCallCuda(ceed, cudaMalloc((void **)&asmb->d_B_in, in_bytes));  in CeedOperatorAssembleSingleSetup_Cuda()
1658 …CeedCallCuda(ceed, cudaMemcpy(&asmb->d_B_in[i * elem_size_in * num_qpts_in], h_B_in, elem_size_in …  in CeedOperatorAssembleSingleSetup_Cuda()
1681     CeedCallCuda(ceed, cudaMalloc((void **)&asmb->d_B_out, out_bytes));  in CeedOperatorAssembleSingleSetup_Cuda()
1693 …CeedCallCuda(ceed, cudaMemcpy(&asmb->d_B_out[i * elem_size_out * num_qpts_out], h_B_out, elem_size…  in CeedOperatorAssembleSingleSetup_Cuda()
1742   if (!impl->asmb) CeedCallBackend(CeedOperatorAssembleSingleSetup_Cuda(op, use_ceedsize_idx));  in CeedOperatorAssembleSingle_Cuda()
1743   CeedOperatorAssemble_Cuda *asmb = impl->asmb;  in CeedOperatorAssembleSingle_Cuda()  local
1745   assert(asmb != NULL);  in CeedOperatorAssembleSingle_Cuda()
1782 …t ? elem_size_in * elem_size_out : 0) + (curl_orients_in ? elem_size_in * asmb->block_size_y : 0))…  in CeedOperatorAssembleSingle_Cuda()
1784   CeedInt grid   = CeedDivUpInt(num_elem_in, asmb->elems_per_block);  in CeedOperatorAssembleSingle_Cuda()
1785 …void   *args[] = {(void *)&num_elem_in, &asmb->d_B_in,     &asmb->d_B_out,      &orients_in,  &cur…  in CeedOperatorAssembleSingle_Cuda()
1788 …nd(CeedRunKernelDimShared_Cuda(ceed, asmb->LinearAssemble, NULL, grid, asmb->block_size_x, asmb->b…  in CeedOperatorAssembleSingle_Cuda()