| /libCEED/backends/sycl-shared/ |
| H A D | ceed-sycl-shared.hpp | 22 sycl::kernel *interp_kernel;
23 sycl::kernel *interp_transpose_kernel;
24 sycl::kernel *grad_kernel;
25 sycl::kernel *grad_transpose_kernel;
26 sycl::kernel *weight_kernel;
|
| H A D | ceed-sycl-shared-basis.sycl.cpp | 64 …sycl::kernel *interp_kernel = (t_mode == CEED_TRANSPOSE) ? impl->interp_transpose_kernel : impl->i… in CeedBasisApplyTensor_Sycl_shared()
85 …sycl::kernel *grad_kernel = (t_mode == CEED_TRANSPOSE) ? impl->grad_transpose_kernel : impl->g… in CeedBasisApplyTensor_Sycl_shared()
|
| /libCEED/backends/cuda/ |
| H A D | ceed-cuda-compile.h | 18 CEED_INTERN int CeedGetKernel_Cuda(Ceed ceed, CUmodule module, const char *name, CUfunction *kernel…
20 CEED_INTERN int CeedRunKernel_Cuda(Ceed ceed, CUfunction kernel, int grid_size, int block_size, voi…
22 CEED_INTERN int CeedRunKernelAutoblockCuda(Ceed ceed, CUfunction kernel, size_t points, void **args…
24 CEED_INTERN int CeedRunKernelDim_Cuda(Ceed ceed, CUfunction kernel, int grid_size, int block_size_x…
26 CEED_INTERN int CeedRunKernelDimShared_Cuda(Ceed ceed, CUfunction kernel, CUstream stream, int grid…
28 CEED_INTERN int CeedTryRunKernelDimShared_Cuda(Ceed ceed, CUfunction kernel, CUstream stream, int g…
|
| H A D | ceed-cuda-compile.cpp | 433 int CeedGetKernel_Cuda(Ceed ceed, CUmodule module, const char *name, CUfunction *kernel) { in CeedGetKernel_Cuda() argument
434 CeedCallCuda(ceed, cuModuleGetFunction(kernel, module, name)); in CeedGetKernel_Cuda()
443 int CeedRunKernelAutoblockCuda(Ceed ceed, CUfunction kernel, size_t points, void **args) { in CeedRunKernelAutoblockCuda() argument
446 …CeedCallCuda(ceed, cuOccupancyMaxPotentialBlockSize(&min_grid_size, &max_block_size, kernel, NULL,… in CeedRunKernelAutoblockCuda()
447 …CeedCallBackend(CeedRunKernel_Cuda(ceed, kernel, CeedDivUpInt(points, max_block_size), max_block_s… in CeedRunKernelAutoblockCuda()
454 int CeedRunKernel_Cuda(Ceed ceed, CUfunction kernel, const int grid_size, const int block_size, voi… in CeedRunKernel_Cuda() argument
455 …CeedCallBackend(CeedRunKernelDimShared_Cuda(ceed, kernel, NULL, grid_size, block_size, 1, 1, 0, ar… in CeedRunKernel_Cuda()
462 int CeedRunKernelDim_Cuda(Ceed ceed, CUfunction kernel, const int grid_size, const int block_size_x… in CeedRunKernelDim_Cuda() argument
464 …CeedCallBackend(CeedRunKernelDimShared_Cuda(ceed, kernel, NULL, grid_size, block_size_x, block_siz… in CeedRunKernelDim_Cuda()
471 static int CeedRunKernelDimSharedCore_Cuda(Ceed ceed, CUfunction kernel, CUstream stream, const int… in CeedRunKernelDimSharedCore_Cuda() argument
[all …]
|
| /libCEED/backends/hip/ |
| H A D | ceed-hip-compile.h | 18 …TERN int CeedGetKernel_Hip(Ceed ceed, hipModule_t module, const char *name, hipFunction_t *kernel);
20 CEED_INTERN int CeedRunKernel_Hip(Ceed ceed, hipFunction_t kernel, int grid_size, int block_size, v…
22 CEED_INTERN int CeedRunKernelDim_Hip(Ceed ceed, hipFunction_t kernel, int grid_size, int block_size…
25 CEED_INTERN int CeedRunKernelDimShared_Hip(Ceed ceed, hipFunction_t kernel, hipStream_t stream, int…
27 CEED_INTERN int CeedTryRunKernelDimShared_Hip(Ceed ceed, hipFunction_t kernel, hipStream_t stream, …
|
| H A D | ceed-hip-compile.cpp | 204 int CeedGetKernel_Hip(Ceed ceed, hipModule_t module, const char *name, hipFunction_t *kernel) { in CeedGetKernel_Hip() argument
205 CeedCallHip(ceed, hipModuleGetFunction(kernel, module, name)); in CeedGetKernel_Hip()
212 int CeedRunKernel_Hip(Ceed ceed, hipFunction_t kernel, const int grid_size, const int block_size, v… in CeedRunKernel_Hip() argument
213 …CeedCallHip(ceed, hipModuleLaunchKernel(kernel, grid_size, 1, 1, block_size, 1, 1, 0, NULL, args, … in CeedRunKernel_Hip()
220 int CeedRunKernelDim_Hip(Ceed ceed, hipFunction_t kernel, const int grid_size, const int block_size… in CeedRunKernelDim_Hip() argument
222 …CeedCallHip(ceed, hipModuleLaunchKernel(kernel, grid_size, 1, 1, block_size_x, block_size_y, block… in CeedRunKernelDim_Hip()
229 static int CeedRunKernelDimSharedCore_Hip(Ceed ceed, hipFunction_t kernel, hipStream_t stream, cons… in CeedRunKernelDimSharedCore_Hip() argument
232 …hipError_t result = hipModuleLaunchKernel(kernel, grid_size, 1, 1, block_size_x, block_size_y, blo… in CeedRunKernelDimSharedCore_Hip()
253 int CeedRunKernelDimShared_Hip(Ceed ceed, hipFunction_t kernel, hipStream_t stream, const int grid_… in CeedRunKernelDimShared_Hip() argument
257 …CeedCallBackend(CeedRunKernelDimSharedCore_Hip(ceed, kernel, stream, grid_size, block_size_x, bloc… in CeedRunKernelDimShared_Hip()
[all …]
|
| /libCEED/backends/xsmm/ |
| H A D | ceed-xsmm-tensor.c | 29 …const libxsmm_gemmfunction kernel = libxsmm_dispatch_gemm(gemm_shape, (libxsmm_bitfield)(flags), (… in CeedTensorContractApply_Xsmm() local
32 …CeedCheck(kernel, CeedTensorContractReturnCeed(contract), CEED_ERROR_BACKEND, "LIBXSMM kernel fail… in CeedTensorContractApply_Xsmm()
38 kernel(&gemm_param); in CeedTensorContractApply_Xsmm()
49 …const libxsmm_gemmfunction kernel = libxsmm_dispatch_gemm(gemm_shape, (libxsmm_bitfield)(flags), (… in CeedTensorContractApply_Xsmm() local
52 …CeedCheck(kernel, CeedTensorContractReturnCeed(contract), CEED_ERROR_BACKEND, "LIBXSMM kernel fail… in CeedTensorContractApply_Xsmm()
59 kernel(&gemm_param); in CeedTensorContractApply_Xsmm()
|
| /libCEED/backends/sycl/ |
| H A D | ceed-sycl-compile.hpp | 19 …ceed, const SyclModule_t *sycl_module, const std::string &kernel_name, sycl::kernel **sycl_kernel);
21 CEED_INTERN int CeedRunKernelDimSharedSycl(Ceed ceed, sycl::kernel *kernel, const int grid_size, co…
|
| H A D | ceed-sycl-compile.sycl.cpp | 143 …eed, const SyclModule_t *sycl_module, const std::string &kernel_name, sycl::kernel **sycl_kernel) { in CeedGetKernel_Sycl()
160 …*sycl_kernel = new sycl::kernel(sycl::make_kernel<sycl::backend::ext_oneapi_level_zero>({*sycl_mod… in CeedGetKernel_Sycl()
169 int CeedRunKernelDimSharedSycl(Ceed ceed, sycl::kernel *kernel, const int grid_size, const int bloc… in CeedRunKernelDimSharedSycl() argument
185 cgh.parallel_for(kernel_range, *kernel); in CeedRunKernelDimSharedSycl()
|
| /libCEED/include/ceed/jit-source/sycl/ |
| H A D | sycl-shared-basis-tensor.h | 22 kernel void Interp(const CeedInt num_elem, global const CeedScalar *restrict d_interp_1d, global co… in Interp()
53 kernel void InterpTranspose(const CeedInt num_elem, global const CeedScalar *restrict d_interp_1d, … in InterpTranspose()
90 kernel void Grad(const CeedInt num_elem, global const CeedScalar *restrict d_interp_1d, global cons… in Grad()
125 kernel void GradTranspose(const CeedInt num_elem, global const CeedScalar *restrict d_interp_1d, gl… in GradTranspose()
163 kernel void Weight(const CeedInt num_elem, global const CeedScalar *restrict q_weight_1d, global Ce… in Weight()
|
| /libCEED/backends/sycl-gen/ |
| H A D | ceed-sycl-gen.hpp | 21 sycl::kernel *op;
|
| /libCEED/backends/sycl-ref/ |
| H A D | ceed-sycl-ref.hpp | 73 sycl::kernel *QFunction;
|
| /libCEED/doc/sphinx/source/ |
| H A D | libCEEDdev.md | 73 …kernel launches, following the libCEED operator decomposition, where first {ref}`CeedElemRestricti…
74 …y to all points across all elements in order to maximize the amount of work each kernel launch has.
83 …kernel to apply the action of the {ref}`CeedOperator`, significantly improving performance by elim…
84 This kernel is compiled at runtime via NVRTC, HIPRTC, or OpenCL RTC.
|
| H A D | releasenotes.md | 168 …ionGetKernelName`; refactored {c:func}`CeedQFunctionGetSourcePath` to exclude function kernel name.
186 - Added {c:func}`CeedPathConcatenate` to facilitate loading kernel source files with a path relativ…
196 - Put GPU JiTed kernel source code into separate files.
|
| H A D | libCEEDapi.md | 158 … $\bm{B}$, $\bm{\mathcal{E}}$ and $\bm{P}$, starting from its point-wise kernel $\bm{D}$, a "matve…
|
| /libCEED/examples/ |
| H A D | README.md | 47 …ctor** and include parallel scatter, element scatter, element evaluation kernel, element gather, a…
|
| /libCEED/examples/fluids/ |
| H A D | index.md | 437 …filter width scaling tensor (also a rank 2 SPD tensor), and $\beta$ is a kernel scaling factor on …
452 …s common to denote a filter width dimensioned relative to the radial distance of the filter kernel.
512 #### Filter kernel scaling, β
513 … a certain physical filter width, the actual width of the implied filter kernel is quite larger th…
515 To match the "size" of a normal kernel to our differential kernel, we attempt to have them match se…
|
| H A D | README.md | 653 - Scaling to make differential kernel size equivalent to other filter kernels
|