summaryrefslogtreecommitdiff
path: root/tangle/LLBM/kernel/collide.h
blob: 7ca02bf1d7f603680063a1e294636b476e3733f9 (plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
#pragma once
#include <LLBM/call_tag.h>

struct BgkCollideO {

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 m0 = f_curr[1] + f_curr[2];
  T m1 = f_curr[3] + f_curr[6];
  T m2 = m0 + m1 + f_curr[0] + f_curr[4] + f_curr[5] + f_curr[7] + f_curr[8];
  T m3 = f_curr[0] - f_curr[8];
  T m4 = T{1} / (m2);
  T rho = m2;
  T u_0 = -m4*(m0 + m3 - f_curr[6] - f_curr[7]);
  T u_1 = -m4*(m1 + m3 - f_curr[2] - f_curr[5]);
  T x0 = T{1} / (tau);
  T x1 = T{0.0138888888888889}*x0;
  T x2 = T{6.00000000000000}*u_1;
  T x3 = T{6.00000000000000}*u_0;
  T x4 = u_0 + u_1;
  T x5 = T{9.00000000000000}*(x4*x4);
  T x6 = u_1*u_1;
  T x7 = T{3.00000000000000}*x6;
  T x8 = u_0*u_0;
  T x9 = T{3.00000000000000}*x8;
  T x10 = x9 + T{-2.00000000000000};
  T x11 = x10 + x7;
  T x12 = T{0.0555555555555556}*x0;
  T x13 = -x3;
  T x14 = T{2.00000000000000} - x7;
  T x15 = x14 + T{6.00000000000000}*x8;
  T x16 = -x9;
  T x17 = x16 + x2;
  T x18 = u_0 - u_1;
  T x19 = x14 + T{9.00000000000000}*(x18*x18);
  T x20 = T{6.00000000000000}*x6;
  f_next[0] = -x1*(rho*(x11 + x2 + x3 - x5) + T{72.0000000000000}*f_curr[0]) + f_curr[0];
  f_next[1] = x12*(rho*(x13 + x15) - T{18.0000000000000}*f_curr[1]) + f_curr[1];
  f_next[2] = x1*(rho*(x13 + x17 + x19) - T{72.0000000000000}*f_curr[2]) + f_curr[2];
  f_next[3] = -x12*(rho*(x10 + x2 - x20) + T{18.0000000000000}*f_curr[3]) + f_curr[3];
  f_next[4] = -T{0.111111111111111}*x0*(T{2.00000000000000}*rho*x11 + T{9.00000000000000}*f_curr[4]) + f_curr[4];
  f_next[5] = x12*(rho*(x17 + x20 + T{2.00000000000000}) - T{18.0000000000000}*f_curr[5]) + f_curr[5];
  f_next[6] = x1*(rho*(x16 + x19 - x2 + x3) - T{72.0000000000000}*f_curr[6]) + f_curr[6];
  f_next[7] = x12*(rho*(x15 + x3) - T{18.0000000000000}*f_curr[7]) + f_curr[7];
  f_next[8] = x1*(rho*(x14 + x17 + x3 + x5) - T{72.0000000000000}*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 m0 = f_curr[10] + f_curr[13] + f_curr[17];
  T m1 = f_curr[14] + f_curr[5] + f_curr[6];
  T m2 = f_curr[1] + f_curr[2] + f_curr[4];
  T m3 = m0 + m1 + m2 + 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 m4 = -f_curr[11] + f_curr[7];
  T m5 = -f_curr[15] + f_curr[3];
  T m6 = T{1} / (m3);
  T m7 = f_curr[0] - f_curr[18];
  T rho = m3;
  T u_0 = m6*(m0 + m4 + m5 - f_curr[1] - f_curr[5] - f_curr[8]);
  T u_1 = m6*(m1 + m4 + m7 - f_curr[12] - f_curr[13] - f_curr[4]);
  T u_2 = m6*(m2 + m5 + m7 - f_curr[14] - f_curr[16] - f_curr[17]);
  T x0 = T{1} / (tau);
  T x1 = T{0.0138888888888889}*x0;
  T x2 = u_1 + u_2;
  T x3 = T{9.00000000000000}*(x2*x2);
  T x4 = T{6.00000000000000}*u_2;
  T x5 = u_1*u_1;
  T x6 = T{3.00000000000000}*x5;
  T x7 = -x6;
  T x8 = u_0*u_0;
  T x9 = T{3.00000000000000}*x8;
  T x10 = T{2.00000000000000} - x9;
  T x11 = x10 + x7;
  T x12 = x11 + x4;
  T x13 = u_2*u_2;
  T x14 = T{3.00000000000000}*x13;
  T x15 = -x14;
  T x16 = T{6.00000000000000}*u_1;
  T x17 = x15 + x16;
  T x18 = T{6.00000000000000}*u_0;
  T x19 = -u_0;
  T x20 = x19 + u_2;
  T x21 = x14 + x6 + T{-2.00000000000000};
  T x22 = x21 + x9;
  T x23 = x22 - x4;
  T x24 = T{0.0277777777777778}*x0;
  T x25 = T{6.00000000000000}*x13;
  T x26 = u_0 + u_2;
  T x27 = T{9.00000000000000}*(x26*x26);
  T x28 = x15 + x18;
  T x29 = -u_1;
  T x30 = x29 + u_2;
  T x31 = x19 + u_1;
  T x32 = -x16;
  T x33 = x18 + x22;
  T x34 = T{6.00000000000000}*x5;
  T x35 = u_0 + u_1;
  T x36 = T{9.00000000000000}*(x35*x35);
  T x37 = T{6.00000000000000}*x8;
  T x38 = x9 + T{-2.00000000000000};
  T x39 = x29 + u_0;
  T x40 = -x18 + x22;
  T x41 = -u_2;
  T x42 = x41 + u_1;
  T x43 = x22 + x4;
  T x44 = x41 + u_0;
  f_next[0] = x1*(rho*(x12 + x17 + x3) - T{72.0000000000000}*f_curr[0]) + f_curr[0];
  f_next[1] = -x1*(rho*(x18 + x23 - T{9.00000000000000}*x20*x20) + T{72.0000000000000}*f_curr[1]) + f_curr[1];
  f_next[2] = x24*(rho*(x12 + x25) - T{36.0000000000000}*f_curr[2]) + f_curr[2];
  f_next[3] =