Skip to content

Instantly share code, notes, and snippets.

__copyright__ = "Copyright (C) 2023 Kaushik Kulkarni"
__license__ = """
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
#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
__kernel void __attribute__ ((reqd_work_group_size(8, 23, 1))) loopy_kernel(__global double const *__restrict__ A, __global double const *__restrict__ B, __global double *__restrict__ C)
{
__local double A_fetch[8 * 11];
__local double B_fetch[11 * 23];
#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
__kernel void __attribute__ ((reqd_work_group_size(9, 10, 1))) loopy_kernel(__global double const *__restrict__ u_ary, __global double const *__restrict__ J_ary, __global double const *__restrict__ D_ary, int const Nel, __global double *__restrict__ out)
{
__local double D_fetch[3 * 9 * 9];
double J_0[3 * 3];
@kaushikcfd
kaushikcfd / simple_grad_0.c
Last active October 13, 2022 23:41
2D DG-FEM Grad Operator over 16 ranks
void rhs(double const *__restrict__ _pt_data, double const *__restrict__ _pt_data_0, double const *__restrict__ _pt_part_ph_id, double const *__restrict__ _pt_data_1, double const *__restrict__ _pt_data_2, char const *__restrict__ from_el_present, double const *__restrict__ _pt_part_ph_id_0, double const *__restrict__ normal_2_b_all, int const *__restrict__ from_el_indices, int const *__restrict__ _pt_data_3, char const *__restrict__ from_el_present_0, double const *__restrict__ _pt_part_ph_id_1, double const *__restrict__ normal_2_b_face_restr_interior, int const *__restrict__ from_el_indices_0, int const *__restrict__ _pt_data_4, char const *__restrict__ from_el_present_1, double const *__restrict__ _pt_part_ph_id_2, double const *__restrict__ _pt_dist_id, int const *__restrict__ from_el_indices_1, int const *__restrict__ _pt_data_5, double const *__restrict__ normal_2_b_BTAG_PARTITION, int const *__restrict__ from_el_indices_2, int const *__restrict__ _pt_data_6, char const *__restrict__ from_el_present_2,
from pymbolic import parse
from pymbolic.mapper import CachedIdentityMapper
code = ("(-1)*((cse_577[_pt_data_48[((iface_ensm15*1075540 + iel_ensm15*10 + idof_ensm15) % 4302160) // 10, 0],"
"_pt_data_49[(iface_ensm15*1075540 + iel_ensm15*10 + idof_ensm15) % 10]]"
" if _pt_data_48[((iface_ensm15*1075540 + iel_ensm15*10 + idof_ensm15) % 4302160) // 10, 0] != -1 else 0)"
" + (cse_577[_pt_data_46[((iface_ensm15*1075540 + iel_ensm15*10 + idof_ensm15) % 4302160) // 10, 0],"
" _pt_data_47[(iface_ensm15*1075540 + iel_ensm15*10 + idof_ensm15) % 10]]"
" if _pt_data_46[((iface_ensm15*1075540 + iel_ensm15*10 + idof_ensm15) % 4302160) // 10, 0] != -1 else 0)"
" + (cse_577[_pt_data_7[((iface_ensm15*1075540 + iel_ensm15*10 + idof_ensm15) % 4302160) // 10, 0],"
#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
inline double powf64(double x, double y) {
return pow(x, y);
}
__kernel void __attribute__ ((reqd_work_group_size(3, 32, 1))) my_rhs(__global double const *__restrict__ _actx_in_1_mass_0, __global double const *__restrict__ _actx_in_1_momentum_2_0, __global double const *__restrict__ nodes0_3d_0, __global double const *__restrict__ nodes1_3d_0, __global int const *__restrict__ from_el_indices_1, __global int const *__restrict__ dof_pick_lists_1, __global char const *__restrict__ dof_pick_list_indices_1, __global double const *__restrict__ normal_1_b_outflow, __global double const *__restrict__ _actx_in_1_momentum_0_0, __global double const *__restrict__ _actx_in_1_momentum_1_0, __global double const *__restrict__ _actx_in_1_energy_0, __global double *__restrict__ cse_10, __global double *__restrict__ cse_116, __global double *__restrict__ cse_121, __
Loading
Sorry, something went wrong. Reload?
Sorry, we cannot display this file.
Sorry, this file is invalid so it cannot be displayed.
743a744,762
>
> # {{{ change meshmode requirements.txt
>
> if [[ "$proj_name" = "mirgecom" ]]; then
> sed -i "s/inducer\/meshmode.git/kaushikcfd\/meshmode.git@with_array_context/g" requirements.txt
> fi
>
> if [[ "$proj_name" = "grudge" ]]; then
> sed -i "s/inducer\/meshmode.git/kaushikcfd\/meshmode.git@with_array_context/g" requirements.txt
Loading
Sorry, something went wrong. Reload?
Sorry, we cannot display this file.
Sorry, this file is invalid so it cannot be displayed.
@kaushikcfd
kaushikcfd / isolator_3d_init_profile.svg
Last active April 23, 2022 18:25
t_final = dt, #Array-ops=6000. Total driver runtime: 13 minutes
Loading
Sorry, something went wrong. Reload?
Sorry, we cannot display this file.
Sorry, this file is invalid so it cannot be displayed.