| /libCEED/backends/hip-gen/ |
| H A D | ceed-hip-gen-operator.c | 152 …CeedInt block_sizes[3] = {data->thread_1d, ((!is_tensor || data->dim == 1) ? 1 : data->thread_1d),… in CeedOperatorApplyAddCore_Hip_gen() local
155 …lBackend(BlockGridCalculate_Hip_gen(data->dim, num_elem, data->max_P_1d, data->Q_1d, block_sizes)); in CeedOperatorApplyAddCore_Hip_gen()
160 block_sizes[2] = elems_per_block; in CeedOperatorApplyAddCore_Hip_gen()
163 …CeedInt grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num… in CeedOperatorApplyAddCore_Hip_gen()
164 CeedInt sharedMem = block_sizes[2] * data->thread_1d * sizeof(CeedScalar); in CeedOperatorApplyAddCore_Hip_gen()
166 …ryRunKernelDimShared_Hip(ceed, data->op, stream, grid, block_sizes[0], block_sizes[1], block_sizes… in CeedOperatorApplyAddCore_Hip_gen()
169 …CeedInt grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num… in CeedOperatorApplyAddCore_Hip_gen()
170 CeedInt sharedMem = block_sizes[2] * data->thread_1d * data->thread_1d * sizeof(CeedScalar); in CeedOperatorApplyAddCore_Hip_gen()
172 …ryRunKernelDimShared_Hip(ceed, data->op, stream, grid, block_sizes[0], block_sizes[1], block_sizes… in CeedOperatorApplyAddCore_Hip_gen()
175 …CeedInt grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num… in CeedOperatorApplyAddCore_Hip_gen()
[all …]
|
| H A D | ceed-hip-gen-operator-build.h | 9 …Calculate_Hip_gen(CeedInt dim, CeedInt num_elem, CeedInt P_1d, CeedInt Q_1d, CeedInt *block_sizes);
|
| H A D | ceed-hip-gen-operator-build.cpp | 34 …eedInt dim, const CeedInt num_elem, const CeedInt P_1d, const CeedInt Q_1d, CeedInt *block_sizes) { in BlockGridCalculate_Hip_gen() argument
40 block_sizes[0] = thread_1d; in BlockGridCalculate_Hip_gen()
41 block_sizes[1] = 1; in BlockGridCalculate_Hip_gen()
42 block_sizes[2] = elems_per_block; in BlockGridCalculate_Hip_gen()
46 block_sizes[0] = thread_1d; in BlockGridCalculate_Hip_gen()
47 block_sizes[1] = thread_1d; in BlockGridCalculate_Hip_gen()
48 block_sizes[2] = elems_per_block; in BlockGridCalculate_Hip_gen()
52 block_sizes[0] = thread_1d; in BlockGridCalculate_Hip_gen()
53 block_sizes[1] = thread_1d; in BlockGridCalculate_Hip_gen()
54 block_sizes[2] = elems_per_block; in BlockGridCalculate_Hip_gen()
[all …]
|
| /libCEED/backends/sycl-gen/ |
| H A D | ceed-sycl-gen-operator.sycl.cpp | 126 CeedInt block_sizes[3], grid = 0; in CeedOperatorApplyAdd_Sycl_gen() local
128 CeedCallBackend(BlockGridCalculate_Sycl_gen(dim, P_1d, Q_1d, block_sizes)); in CeedOperatorApplyAdd_Sycl_gen()
130 …grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num_elem) ? 1 : … in CeedOperatorApplyAdd_Sycl_gen()
133 …grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num_elem) ? 1 : … in CeedOperatorApplyAdd_Sycl_gen()
136 …grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num_elem) ? 1 : … in CeedOperatorApplyAdd_Sycl_gen()
140 sycl::range<3> local_range(block_sizes[2], block_sizes[1], block_sizes[0]); in CeedOperatorApplyAdd_Sycl_gen()
141 sycl::range<3> global_range(grid * block_sizes[2], block_sizes[1], block_sizes[0]); in CeedOperatorApplyAdd_Sycl_gen()
|
| H A D | ceed-sycl-gen-operator-build.sycl.cpp | 30 …lculate_Sycl_gen(const CeedInt dim, const CeedInt P_1d, const CeedInt Q_1d, CeedInt *block_sizes) { in BlockGridCalculate_Sycl_gen() argument
37 block_sizes[0] = thread1d; in BlockGridCalculate_Sycl_gen()
38 block_sizes[1] = 1; in BlockGridCalculate_Sycl_gen()
39 block_sizes[2] = elems_per_block; in BlockGridCalculate_Sycl_gen()
43 block_sizes[0] = thread1d; in BlockGridCalculate_Sycl_gen()
44 block_sizes[1] = thread1d; in BlockGridCalculate_Sycl_gen()
45 block_sizes[2] = elems_per_block; in BlockGridCalculate_Sycl_gen()
49 block_sizes[0] = thread1d; in BlockGridCalculate_Sycl_gen()
50 block_sizes[1] = thread1d; in BlockGridCalculate_Sycl_gen()
51 block_sizes[2] = elems_per_block; in BlockGridCalculate_Sycl_gen()
[all …]
|
| H A D | ceed-sycl-gen-operator-build.hpp | 9 …alculate_Sycl_gen(const CeedInt dim, const CeedInt P_1d, const CeedInt Q_1d, CeedInt *block_sizes);
|
| /libCEED/backends/hip-shared/ |
| H A D | ceed-hip-shared-basis.c | 39 …eedInt dim, const CeedInt P_1d, const CeedInt Q_1d, const CeedInt num_comp, CeedInt *block_sizes) { in ComputeBasisThreadBlockSizes() argument
49 block_sizes[0] = 256; in ComputeBasisThreadBlockSizes()
52 block_sizes[1] = 256; in ComputeBasisThreadBlockSizes()
55 block_sizes[2] = 256; in ComputeBasisThreadBlockSizes()
61 block_sizes[0] = CeedIntMax(256, ComputeBlockSizeFromRequirement(required)); in ComputeBasisThreadBlockSizes()
64 block_sizes[1] = CeedIntMax(256, ComputeBlockSizeFromRequirement(required)); in ComputeBasisThreadBlockSizes()
68 block_sizes[2] = CeedIntMax(256, ComputeBlockSizeFromRequirement(required)); in ComputeBasisThreadBlockSizes()
75 block_sizes[0] = CeedIntMax(256, ComputeBlockSizeFromRequirement(required)); in ComputeBasisThreadBlockSizes()
78 block_sizes[1] = CeedIntMax(256, ComputeBlockSizeFromRequirement(required)); in ComputeBasisThreadBlockSizes()
82 block_sizes[2] = CeedIntMax(256, ComputeBlockSizeFromRequirement(required)); in ComputeBasisThreadBlockSizes()
[all …]
|
| H A D | ceed-hip-shared.h | 30 CeedInt block_sizes[3]; // interp, grad, weight thread block sizes member
|
| /libCEED/examples/fluids/src/ |
| H A D | mat-ceed.c | 551 const PetscInt *block_sizes; in MatCeedCopy() local
553 PetscCall(MatGetVariableBlockSizes(mat_ceed, &num_blocks, &block_sizes)); in MatCeedCopy()
554 …f (num_blocks) PetscCall(MatSetVariableBlockSizes(mat_other, num_blocks, (PetscInt *)block_sizes)); in MatCeedCopy()
|