This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
--------------------------------------------------------------------------- | |
KERNEL: resample_by_mat | |
--------------------------------------------------------------------------- | |
ARGUMENTS: | |
ary: type: np:dtype('float64'), shape: (nelements_vec, n_from_nodes), dim_tags: (N1:stride:n_from_nodes, N0:stride:1), offset: <class 'loopy.kernel.data.auto'> aspace: global | |
from_element_indices: type: np:dtype('int64'), shape: (nelements), dim_tags: (N0:stride:1), offset: <class 'loopy.kernel.data.auto'> aspace: global | |
n_from_nodes: ValueArg, type: np:dtype('int64') | |
n_to_nodes: ValueArg, type: np:dtype('int64') | |
nelements: ValueArg, type: np:dtype('int64') | |
nelements_vec: ValueArg, type: np:dtype('int32') |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
from kanren import var, eq, run, facts, Relation | |
from pytools.tag import Tag, tag_dataclass, UniqueTag | |
from typing import FrozenSet | |
from dataclasses import dataclass | |
from unification import unifiable | |
@unifiable | |
@tag_dataclass | |
class AxisPropagatedFromInput(UniqueTag): |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
import loopy as lp | |
import numpy as np | |
_pt_kernel_knl = lp.make_kernel([ | |
"{ [_pt_sum_r0_1, _pt_sum_r0_17] : 0<=_pt_sum_r0_1<=2 and 0<=_pt_sum_r0_17<=2 }", | |
"{ [_pt_sum_r0_0] : 0 <= _pt_sum_r0_0 <= 34 }", | |
"{ [iel_cse_1, idof_cse_1] : 0<=iel_cse_1<=187199 and 0<=idof_cse_1<=14 }", | |
"{ [iel_cse_0, idof_cse_0] : 0<=iel_cse_0<=187199 and 0<=idof_cse_0<=14 }", | |
"{ [iel_cse_2, idof_cse_2] : 0<=iel_cse_2<=4799 and 0<=idof_cse_2<=14 }", |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
for idof_face_mass, iel_face_mass | |
acc_face_f = 0.0 {id=face_insn_face_f_init} | |
acc_face_f_0 = 0.0 {id=face_insn_0_face_f_0_init} | |
acc_face_f_1 = 0.0 {id=face_insn_1_face_f_1_init} | |
acc_face_f_2 = 0.0 {id=face_insn_2_face_f_2_init} | |
for face_f | |
acc_face_jdof = 0.0 {id=face_insn_face_f_update_face_jdof_init} | |
acc_face_jdof_0 = 0.0 {id=face_insn_0_face_f_0_update_face_jdof_0_init} | |
acc_face_jdof_1 = 0.0 {id=face_insn_1_face_f_1_update_face_jdof_1_init} | |
acc_face_jdof_2 = 0.0 {id=face_insn_2_face_f_2_update_face_jdof_2_init} |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
#define lid(N) ((int) get_local_id(N)) | |
#define gid(N) ((int) get_group_id(N)) | |
#if __OPENCL_C_VERSION__ < 120 | |
#pragma OPENCL EXTENSION cl_khr_fp64: enable | |
#endif | |
#define LOOPY_CALL_WITH_INTEGER_TYPES(MACRO_NAME) \ | |
MACRO_NAME(int8, char) \ | |
MACRO_NAME(int16, short) \ | |
MACRO_NAME(int32, int) \ | |
MACRO_NAME(int64, long) |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
--------------------------------------------------------------------------- | |
KERNEL: _pt_kernel | |
--------------------------------------------------------------------------- | |
ARGUMENTS: | |
_actx_in_0: type: np:dtype('float64'), shape: () aspace: global | |
_actx_in_1_0_0: type: np:dtype('float64'), shape: (48000, 35), dim_tags: (N1:stride:35, N0:stride:1) aspace: global | |
_actx_in_1_1_0: type: np:dtype('float64'), shape: (48000, 35), dim_tags: (N1:stride:35, N0:stride:1) aspace: global | |
_actx_in_1_2_0: type: np:dtype('float64'), shape: (48000, 35), dim_tags: (N1:stride:35, N0:stride:1) aspace: global | |
_actx_in_1_3_0: type: np:dtype('float64'), shape: (48000, 35), dim_tags: (N1:stride:35, N0:stride:1) aspace: global | |
_pt_in: type: np:dtype('float64'), shape: (48000, 35), dim_tags: (N1:stride:35, N0:stride:1) aspace: global |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
--------------------------------------------------------------------------- | |
KERNEL: face_mass | |
--------------------------------------------------------------------------- | |
ARGUMENTS: | |
jac_surf: type: np:dtype('float64'), shape: (nfaces, nelements, nface_nodes), dim_tags: (N2:stride:nface_nodes*nelements, N1:stride:nface_nodes, N0:stride:1), offset: <class 'loopy.kernel.data.auto'> aspace: global | |
mat: type: np:dtype('float64'), shape: (nvol_nodes, nfaces, nface_nodes), dim_tags: (N2:stride:nface_nodes*nfaces, N1:stride:nface_nodes, N0:stride:1), offset: <class 'loopy.kernel.data.auto'> aspace: global | |
nelements: ValueArg, type: np:dtype('int64') | |
nface_nodes: ValueArg, type: np:dtype('int64') | |
nfaces: ValueArg, type: np:dtype('int64') | |
nvol_nodes: ValueArg, type: np:dtype('int64') |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
import numpy as np | |
import pyopencl as cl | |
from grudge.discretization import DiscretizationCollection | |
import grudge.op as op | |
from meshmode.mesh.generation import generate_box_mesh | |
from meshmode.array_context import PyOpenCLArrayContext | |
from arraycontext import freeze, thaw, to_numpy | |
from grudge.dof_desc import DTAG_BOUNDARY | |
from meshmode.discretization.connection import FACE_RESTR_INTERIOR |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
diff --git a/pytato/analysis.py b/pytato/analysis.py | |
index 3bd2839..93147b1 100644 | |
--- a/pytato/analysis.py | |
+++ b/pytato/analysis.py | |
@@ -24,13 +24,19 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN | |
THE SOFTWARE. | |
""" | |
-from typing import Mapping, Dict, Union | |
+from typing import Mapping, Dict, Union, FrozenSet |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
#define lid(N) ((int) get_local_id(N)) | |
#define gid(N) ((int) get_group_id(N)) | |
#if __OPENCL_C_VERSION__ < 120 | |
#pragma OPENCL EXTENSION cl_khr_fp64: enable | |
#endif | |
static void face_mass(__global double const *__restrict__ jac_surf, __global double const *__restrict__ mat, int const nelements, int const nface_nodes, int const nfaces, int const nvol_nodes, __global double *__restrict__ result, __global double const *__restrict__ vec); | |
static void resample_by_picking(__global double const *__restrict__ ary, int const nelements_vec, int const n_from_nodes, __global long const *__restrict__ from_element_indices, int const n_to_nodes, int const nelements, __global int const *__restrict__ pick_list, __global double *__restrict__ result); | |
static void resample_by_picking_0(__global double const *__restrict__ ary, int const nelements_vec, int const n_from_nodes, __global int const *__restrict__ from_element_indices, int const n_to_nodes, int const nelements, __global int const *__restrict__ pick_list, __global double *__ |