diff options
Diffstat (limited to 'tangle/LLBM/kernel/smagorinsky_collide.h')
| -rw-r--r-- | tangle/LLBM/kernel/smagorinsky_collide.h | 183 |
1 files changed, 183 insertions, 0 deletions
diff --git a/tangle/LLBM/kernel/smagorinsky_collide.h b/tangle/LLBM/kernel/smagorinsky_collide.h new file mode 100644 index 0000000..3489479 --- /dev/null +++ b/tangle/LLBM/kernel/smagorinsky_collide.h @@ -0,0 +1,183 @@ +#pragma once +#include <LLBM/call_tag.h> + +struct SmagorinskyBgkCollideO { + +using call_tag = tag::call_by_cell_id; + +template <typename T, typename S> +__device__ static void apply(descriptor::D2Q9, S f_curr[9], S f_next[9], std::size_t gid, T tau, T smagorinsky) { + T x0 = f_curr[1] + f_curr[2]; + T x1 = f_curr[3] + f_curr[6]; + T x2 = x0 + x1 + f_curr[0] + f_curr[4] + f_curr[5] + f_curr[7] + f_curr[8]; + T x3 = f_curr[0] - f_curr[8]; + T x4 = T{1} / (x2); + T x5 = T{72.0000000000000}*f_curr[2]; + T x6 = T{72.0000000000000}*f_curr[6]; + T rho = x2; + T x31 = T{4.00000000000000}*rho; + T x40 = T{2.00000000000000}*rho; + T u_0 = -x4*(x0 + x3 - f_curr[6] - f_curr[7]); + T x7 = T{6.00000000000000}*u_0; + T x8 = -x7; + T x15 = u_0*u_0; + T x16 = T{3.00000000000000}*x15; + T x17 = -x16; + T x27 = x16 + T{-2.00000000000000}; + T u_1 = -x4*(x1 + x3 - f_curr[2] - f_curr[5]); + T x9 = u_0 - u_1; + T x10 = T{9.00000000000000}*(x9*x9); + T x11 = u_1*u_1; + T x12 = T{3.00000000000000}*x11; + T x13 = T{2.00000000000000} - x12; + T x14 = T{6.00000000000000}*u_1; + T x18 = x14 + x17; + T x19 = x13 + x18; + T x20 = x10 + x19 + x8; + T x21 = rho*x20; + T x22 = x10 + x13 - x14 + x17 + x7; + T x23 = rho*x22; + T x24 = u_0 + u_1; + T x25 = T{9.00000000000000}*(x24*x24); + T x26 = x19 + x25 + x7; + T x28 = x12 + x27; + T x29 = x14 - x25 + x28 + x7; + T x30 = rho*x26 - rho*x29 - T{72.0000000000000}*f_curr[0] - T{72.0000000000000}*f_curr[8]; + T x32 = x13 + T{6.00000000000000}*x15; + T x33 = x32 + x8; + T x34 = x32 + x7; + T x35 = x21 + x23 + x30 - x5 - x6; + T x36 = T{6.00000000000000}*x11; + T x37 = x14 + x27 - x36; + T x38 = x18 + x36 + T{2.00000000000000}; + T x39 = T{1} / (tau + sqrt(T{0.707106781186548}*(smagorinsky*smagorinsky)*sqrt((-x21 - x23 + x30 + x5 + x6)*(-x21 - x23 + x30 + x5 + x6) + T{0.500000000000000}*((x31*x33 + x31*x34 + x35 - 72*f_curr[1] - 72*f_curr[7])*(x31*x33 + x31*x34 + x35 - 72*f_curr[1] - 72*f_curr[7])) + T{0.500000000000000}*((-x31*x37 + x31*x38 + x35 - 72*f_curr[3] - 72*f_curr[5])*(-x31*x37 + x31*x38 + x35 - 72*f_curr[3] - 72*f_curr[5]))) + tau*tau)); + f_next[0] = -T{0.0138888888888889}*x39*(x29*x40 + T{144.000000000000}*f_curr[0]) + f_curr[0]; + f_next[1] = T{0.0555555555555556}*x39*(x33*x40 - T{36.0000000000000}*f_curr[1]) + f_curr[1]; + f_next[2] = T{0.0138888888888889}*x39*(x20*x40 - T{144.000000000000}*f_curr[2]) + f_curr[2]; + f_next[3] = -T{0.0555555555555556}*x39*(x37*x40 + T{36.0000000000000}*f_curr[3]) + f_curr[3]; + f_next[4] = -T{0.111111111111111}*x39*(T{4.00000000000000}*rho*x28 + T{18.0000000000000}*f_curr[4]) + f_curr[4]; + f_next[5] = T{0.0555555555555556}*x39*(x38*x40 - T{36.0000000000000}*f_curr[5]) + f_curr[5]; + f_next[6] = T{0.0138888888888889}*x39*(x22*x40 - T{144.000000000000}*f_curr[6]) + f_curr[6]; + f_next[7] = T{0.0555555555555556}*x39*(x34*x40 - T{36.0000000000000}*f_curr[7]) + f_curr[7]; + f_next[8] = T{0.0138888888888889}*x39*(x26*x40 - T{144.000000000000}*f_curr[8]) + f_curr[8]; + +} + +template <typename T, typename S> +__device__ static void apply(descriptor::D3Q19, S f_curr[19], S f_next[19], std::size_t gid, T tau, T smagorinsky) { + T x0 = f_curr[10] + f_curr[13] + f_curr[17]; + T x1 = f_curr[14] + f_curr[5] + f_curr[6]; + T x2 = f_curr[1] + f_curr[2] + f_curr[4]; + T x3 = x0 + x1 + x2 + f_curr[0] + f_curr[11] + f_curr[12] + f_curr[15] + f_curr[16] + f_curr[18] + f_curr[3] + f_curr[7] + f_curr[8] + f_curr[9]; + T x4 = -f_curr[11] + f_curr[7]; + T x5 = -f_curr[15] + f_curr[3]; + T x6 = T{1} / (x3); + T x7 = f_curr[0] - f_curr[18]; + T x8 = T{72.0000000000000}*f_curr[5]; + T x9 = T{72.0000000000000}*f_curr[13]; + T x42 = T{72.0000000000000}*f_curr[1]; + T x43 = T{72.0000000000000}*f_curr[17]; + T x61 = T{72.0000000000000}*f_curr[4]; + T x62 = T{72.0000000000000}*f_curr[14]; + T rho = x3; + T x74 = T{2.00000000000000}*rho; + T x89 = T{2.00000000000000}*rho; + T u_0 = x6*(x0 + x4 + x5 - f_curr[1] - f_curr[5] - f_curr[8]); + T x12 = -u_0; + T x14 = T{6.00000000000000}*u_0; + T x17 = u_0*u_0; + T x18 = T{3.00000000000000}*x17; + T x26 = -x14; + T x35 = T{2.00000000000000} - x18; + T x75 = T{6.00000000000000}*x17; + T u_1 = x6*(x1 + x4 + x7 - f_curr[12] - f_curr[13] - f_curr[4]); + T x10 = T{6.00000000000000}*u_1; + T x11 = -x10; + T x13 = x12 + u_1; + T x19 = u_1*u_1; + T x20 = T{3.00000000000000}*x19; + T x21 = x18 + x20 + T{-2.00000000000000}; + T x27 = -u_1; + T x28 = x27 + u_0; + T x32 = u_0 + u_1; + T x33 = T{9.00000000000000}*(x32*x32); + T x34 = -x20; + T x36 = x34 + x35; + T x81 = T{6.00000000000000}*x19; + T u_2 = x6*(x2 + x5 + x7 - f_curr[14] - f_curr[16] - f_curr[17]); + T x15 = u_2*u_2; + T x16 = T{3.00000000000000}*x15; + T x22 = x16 + x21; + T x23 = x14 + x22; + T x24 = x11 + x23 - T{9.00000000000000}*x13*x13; + T x25 = rho*x24; + T x29 = x10 + x22; + T x30 = x26 + x29 - T{9.00000000000000}*x28*x28; + T x31 = rho*x30; + T x37 = -x16; + T x38 = x10 + x37; + T x39 = x14 + x33 + x36 + x38; + T x40 = x14 + x29 - x33; + T x41 = rho*x39 - rho*x40 - T{72.0000000000000}*f_curr[11] - T{72.0000000000000}*f_curr[7]; + T x44 = T{6.00000000000000}*u_2; + T x45 = -x44; + T x46 = x12 + u_2; + T x47 = x23 + x45 - T{9.00000000000000}*x46*x46; + T x48 = rho*x47; + T x49 = -u_2; + T x50 = x49 + u_0; + T x51 = x22 + x44; + T x52 = x26 + x51 - T{9.00000000000000}*x50*x50; + T x53 = rho*x52; + T x54 = u_0 + u_2; + T x55 = T{9.00000000000000}*(x54*x54); + T x56 = x36 + x44; + T x57 = x14 + x37; + T x58 = x55 + x56 + x57; + T x59 = x14 + x51 - x55; + T x60 = rho*x58 - rho*x59 - T{72.0000000000000}*f_curr[15] - T{72.0000000000000}*f_curr[3]; + T x63 = x27 + u_2; + T x64 = x29 + x45 - T{9.00000000000000}*x63*x63; + T x65 = rho*x64; + T x66 = x49 + u_1; + T x67 = x11 + x51 - T{9.00000000000000}*x66*x66; + T x68 = rho*x67; + T x69 = u_1 + u_2; + T x70 = T{9.00000000000000}*(x69*x69); + T x71 = x38 + x56 + x70; + T x72 = x29 + x44 - x70; + T x73 = rho*x71 - rho*x72 - T{72.0000000000000}*f_curr[0] - T{72.0000000000000}*f_curr[18]; + T x76 = x16 + T{-2.00000000000000}; + T x77 = x14 + x20 - x75 + x76; + T x78 = x34 + x57 + x75 + T{2.00000000000000}; + T x79 = -x42 - x43 - x48 - x53 + x60; + T x80 = -x25 - x31 + x41 - x8 - x9; + T x82 = x10 + x18 + x76 - x81; + T x83 = x35 + x38 + x81; + T x84 = -x61 - x62 - x65 - x68 + x73; + T x85 = T{6.00000000000000}*x15; + T x86 = x21 + x44 - x85; + T x87 = x56 + x85; + T x88 = T{1} / (tau + sqrt(T{0.707106781186548}*(smagorinsky*smagorinsky)*sqrt((x25 + x31 + x41 + x8 + x9)*(x25 + x31 + x41 + x8 + x9) + (x42 + x43 + x48 + x53 + x60)*(x42 + x43 + x48 + x53 + x60) + (x61 + x62 + x65 + x68 + x73)*(x61 + x62 + x65 + x68 + x73) + T{0.500000000000000}*((-x74*x77 + x74*x78 + x79 + x80 - 72*f_curr[10] - 72*f_curr[8])*(-x74*x77 + x74*x78 + x79 + x80 - 72*f_curr[10] - 72*f_curr[8])) + T{0.500000000000000}*((-x74*x82 + x74*x83 + x80 + x84 - 72*f_curr[12] - 72*f_curr[6])*(-x74*x82 + x74*x83 + x80 + x84 - 72*f_curr[12] - 72*f_curr[6])) + T{0.500000000000000}*((-x74*x86 + x74*x87 + x79 + x84 - 72*f_curr[16] - 72*f_curr[2])*(-x74*x86 + x74*x87 + x79 + x84 - 72*f_curr[16] - 72*f_curr[2]))) + tau*tau)); + f_next[0] = T{0.0138888888888889}*x88*(x71*x89 - T{144.000000000000}*f_curr[0]) + f_curr[0]; + f_next[1] = -T{0.0138888888888889}*x88*(x47*x89 + T{144.000000000000}*f_curr[1]) + f_curr[1]; + f_next[2] = T{0.0277777777777778}*x88*(x87*x89 - T{72.0000000000000}*f_curr[2]) + f_curr[2]; + f_next[3] = T{0.0138888888888889}*x88*(x58*x89 - T{144.000000000000}*f_curr[3]) + f_curr[3]; + f_next[4] = -T{0.0138888888888889}*x88*(x64*x89 + T{144.000000000000}*f_curr[4]) + f_curr[4]; + f_next[5] = -T{0.0138888888888889}*x88*(x24*x89 + T{144.000000000000}*f_curr[5]) + f_curr[5]; + f_next[6] = T{0.0277777777777778}*x88*(x83*x89 - T{72.0000000000000}*f_curr[6]) + f_curr[6]; + f_next[7] = T{0.0138888888888889}*x88*(x39*x89 - T{144.000000000000}*f_curr[7]) + f_curr[7]; + f_next[8] = -T{0.0277777777777778}*x88*(x77*x89 + T{72.0000000000000}*f_curr[8]) + f_curr[8]; + f_next[9] = -T{0.166666666666667}*x88*(x22*x89 + T{12.0000000000000}*f_curr[9]) + f_curr[9]; + f_next[10] = T{0.0277777777777778}*x88*(x78*x89 - T{72.0000000000000}*f_curr[10]) + f_curr[10]; + f_next[11] = -T{0.0138888888888889}*x88*(x40*x89 + T{144.000000000000}*f_curr[11]) + f_curr[11]; + f_next[12] = -T{0.0277777777777778}*x88*(x82*x89 + T{72.0000000000000}*f_curr[12]) + f_curr[12]; + f_next[13] = -T{0.0138888888888889}*x88*(x30*x89 + T{144.000000000000}*f_curr[13]) + f_curr[13]; + f_next[14] = -T{0.0138888888888889}*x88*(x67*x89 + T{144.000000000000}*f_curr[14]) + f_curr[14]; + f_next[15] = -T{0.0138888888888889}*x88*(x59*x89 + T{144.000000000000}*f_curr[15]) + f_curr[15]; + f_next[16] = -T{0.0277777777777778}*x88*(x86*x89 + T{72.0000000000000}*f_curr[16]) + f_curr[16]; + f_next[17] = -T{0.0138888888888889}*x88*(x52*x89 + T{144.000000000000}*f_curr[17]) + f_curr[17]; + f_next[18] = -T{0.0138888888888889}*x88*(x72*x89 + T{144.000000000000}*f_curr[18]) + f_curr[18]; +} + +}; |