embedded3 R6-2: the moving mask's classification on the device (RTX_E3_MASK_DEVICE=1, needs RTX_E3_GEOM_DEVICE=1)
- e3_mask.cu: the changed set (touched by either build, dilated), the faces of the changed cells, per changed cell the fluid flag / space-time activity / GCL entry and the merging master, per face the kind, the step aperture and open flag and the open-part centroid shift (compute_face_shifts' cv geometry), the GCL table's wall cells / areas / table with the correction, the merged cells' CSR (slave ranks from the master's distinct face neighbours), the imposition band's faces; block scans for ascending stream compaction and the CSR's exclusive scan. fp64, FMA contraction off. - step/device/mask.rs DeviceMask: snapshots the previous apertures / volumes / evaluated cells before the R6-1 geometry kernels, runs the classification into DeviceCut's persistent tables (a, open, open_pred, active, active_pred, owner, shift, wall_flux, fold_ptr/fold_idx) and returns the compact values (MaskUpdate); the GCL sums in ascending order on the host. DeviceCut::update_after_device_mask: only the band's surface velocities remain (band list compacted on the device). - maskupdate.rs: Mask::from_update — the host mirror from the old mask's step arrays and merging map (moved) and the mask retired a step earlier (instantaneous arrays, pooled), the compact values scattered; incremental fluid count / anchor / side check; the fresh-cell refill over the changed set. - RTX_E3_BAND_CHECK=1 with the knob: the host rebuild on copies (from_cut, face shifts, step apertures, merging, GCL, refill) against the mirror, bit for bit, and every device table against a full build. - Knob off: the host path unchanged (Mask::from_parts, Solver::configure_mask and retire_mask are the same statements); profile laps under RTX_E3_MOVING_PROFILE only. Co-Authored-By: Claude Opus 5.5 (1M context) <[email protected]>
This commit is contained in:
Omar Sobh
co-authored by
Claude Opus 5.5
Claude Opus 5.5
parent
f5f0ffdd2a
commit
0150f79a4a
@@ -0,0 +1,428 @@
|
||||
/**
|
||||
* embedded3 R6-2: the moving mask's classification on the device, behind
|
||||
* `RTX_E3_MASK_DEVICE=1` (with `RTX_E3_GEOM_DEVICE=1`). From the device
|
||||
* geometry of the step (R6-1: the end-of-step apertures in DeviceCut's
|
||||
* predictor set, the volumes and wall vectors in DeviceGeom) and the
|
||||
* previous step's (snapshots taken before the geometry kernels overwrite
|
||||
* them), the host `Mask` expressions in fp64 with FMA contraction off:
|
||||
*
|
||||
* e3_mask_changed the changed set (touched by either build, dilated by
|
||||
* one; `set_step_apertures_with`);
|
||||
* e3_mask_face_flags the faces of the changed cells per component;
|
||||
* e3_mask_cells per changed cell: fluid (vol > 0), the space-time
|
||||
* activity, the GCL entry (`gcl_flux_table`);
|
||||
* e3_mask_merge per changed cell: the virtual-merging master
|
||||
* (`compute_merging` with the old mask);
|
||||
* e3_mask_faces per face of a changed cell: the kind, the step
|
||||
* aperture and open flag (`set_step_apertures`), the
|
||||
* open-part centroid shift (`compute_face_shifts`);
|
||||
* e3_mask_wall_flags / e3_mask_wall_area / e3_mask_wall_flux
|
||||
* the GCL table's wall cells, areas and the table
|
||||
* with the compatibility correction;
|
||||
* e3_mask_slave_flags / e3_mask_fold_count / e3_mask_fold_fill
|
||||
* the merged cells' CSR (slaves ascending per master);
|
||||
* e3_mask_band_flags the imposition band's faces (|d| <= band);
|
||||
* e3_scan_* block scans: stream compaction (ascending) and the
|
||||
* exclusive scan of counts.
|
||||
* Every per-entry value depends on its own inputs alone; the two serial
|
||||
* sums of the GCL table (in ascending order) are the host's.
|
||||
*/
|
||||
|
||||
typedef unsigned int u32;
|
||||
typedef unsigned char u8;
|
||||
|
||||
struct MaskGrid {
|
||||
int nx, ny, nz, periodic;
|
||||
double dx, dy, dz;
|
||||
};
|
||||
|
||||
#define SCAN_BLOCK 1024
|
||||
|
||||
/* Exclusive block scan of v over a 1024-thread block; *total = the block sum. */
|
||||
__device__ u32 block_excl_scan(u32 v, u32* total)
|
||||
{
|
||||
__shared__ u32 warp_sums[32];
|
||||
int lane = threadIdx.x & 31, wid = threadIdx.x >> 5;
|
||||
u32 x = v;
|
||||
for (int o = 1; o < 32; o <<= 1) {
|
||||
u32 y = __shfl_up_sync(0xffffffffu, x, o);
|
||||
if (lane >= o) x += y;
|
||||
}
|
||||
if (lane == 31) warp_sums[wid] = x;
|
||||
__syncthreads();
|
||||
if (wid == 0) {
|
||||
u32 s = warp_sums[lane];
|
||||
for (int o = 1; o < 32; o <<= 1) {
|
||||
u32 y = __shfl_up_sync(0xffffffffu, s, o);
|
||||
if (lane >= o) s += y;
|
||||
}
|
||||
warp_sums[lane] = s;
|
||||
}
|
||||
__syncthreads();
|
||||
u32 incl = x + (wid > 0 ? warp_sums[wid - 1] : 0u);
|
||||
*total = warp_sums[31];
|
||||
__syncthreads();
|
||||
return incl - v;
|
||||
}
|
||||
|
||||
/* Per block of 1024: the number of flags set. */
|
||||
extern "C" __global__ void e3_scan_count_flags(long long n, const u8* __restrict__ flags, u32* __restrict__ sums)
|
||||
{
|
||||
long long i = (long long) blockIdx.x * SCAN_BLOCK + threadIdx.x;
|
||||
u32 total;
|
||||
block_excl_scan(i < n ? (u32) (flags[i] != 0) : 0u, &total);
|
||||
if (threadIdx.x == 0) sums[blockIdx.x] = total;
|
||||
}
|
||||
|
||||
/* Per block of 1024: the sum of the values. */
|
||||
extern "C" __global__ void e3_scan_count_vals(long long n, const u32* __restrict__ vals, u32* __restrict__ sums)
|
||||
{
|
||||
long long i = (long long) blockIdx.x * SCAN_BLOCK + threadIdx.x;
|
||||
u32 total;
|
||||
block_excl_scan(i < n ? vals[i] : 0u, &total);
|
||||
if (threadIdx.x == 0) sums[blockIdx.x] = total;
|
||||
}
|
||||
|
||||
/* One block: the block sums scanned exclusively in place; sums[nb] = the total. */
|
||||
extern "C" __global__ void e3_scan_top(int nb, u32* __restrict__ sums)
|
||||
{
|
||||
__shared__ u32 carry;
|
||||
if (threadIdx.x == 0) carry = 0u;
|
||||
__syncthreads();
|
||||
for (int base = 0; base < nb; base += SCAN_BLOCK) {
|
||||
int i = base + threadIdx.x;
|
||||
u32 v = i < nb ? sums[i] : 0u;
|
||||
u32 total;
|
||||
u32 ex = block_excl_scan(v, &total);
|
||||
u32 c = carry;
|
||||
if (i < nb) sums[i] = c + ex;
|
||||
__syncthreads();
|
||||
if (threadIdx.x == 0) carry = c + total;
|
||||
__syncthreads();
|
||||
}
|
||||
if (threadIdx.x == 0) sums[nb] = carry;
|
||||
}
|
||||
|
||||
/* Stream compaction: the flagged indices, ascending, at their scanned offsets. */
|
||||
extern "C" __global__ void e3_scan_compact(long long n, const u8* __restrict__ flags, const u32* __restrict__ offsets, u32* __restrict__ out)
|
||||
{
|
||||
long long i = (long long) blockIdx.x * SCAN_BLOCK + threadIdx.x;
|
||||
u32 v = i < n ? (u32) (flags[i] != 0) : 0u;
|
||||
u32 total;
|
||||
u32 ex = block_excl_scan(v, &total);
|
||||
if (v) out[offsets[blockIdx.x] + ex] = (u32) i;
|
||||
}
|
||||
|
||||
/* The exclusive scan of the values. */
|
||||
extern "C" __global__ void e3_scan_values(long long n, const u32* __restrict__ vals, const u32* __restrict__ offsets, u32* __restrict__ out)
|
||||
{
|
||||
long long i = (long long) blockIdx.x * SCAN_BLOCK + threadIdx.x;
|
||||
u32 total;
|
||||
u32 ex = block_excl_scan(i < n ? vals[i] : 0u, &total);
|
||||
if (i < n) out[i] = offsets[blockIdx.x] + ex;
|
||||
}
|
||||
|
||||
/* The changed set: touched by either build, dilated by one (periodic z wraps). */
|
||||
extern "C" __global__ void e3_mask_changed(MaskGrid g, const u8* __restrict__ t_new, const u8* __restrict__ t_old, u8* __restrict__ flag)
|
||||
{
|
||||
long long idx = (long long) blockIdx.x * blockDim.x + threadIdx.x;
|
||||
long long nx = g.nx, ny = g.ny, nz = g.nz, nxy = nx * ny;
|
||||
if (idx >= nxy * nz) return;
|
||||
long long k = idx / nxy, j = (idx % nxy) / nx, i = idx % nx;
|
||||
#define T(q) (t_new[q] | t_old[q])
|
||||
u8 f = T(idx)
|
||||
|| (i + 1 < nx && T(idx + 1)) || (i > 0 && T(idx - 1))
|
||||
|| (j + 1 < ny && T(idx + nx)) || (j > 0 && T(idx - nx))
|
||||
|| (k + 1 < nz && T(idx + nxy)) || (k > 0 && T(idx - nxy))
|
||||
|| (g.periodic && nz > 1 && k + 1 == nz && T(idx - (nz - 1) * nxy))
|
||||
|| (g.periodic && nz > 1 && k == 0 && T(idx + (nz - 1) * nxy));
|
||||
#undef T
|
||||
flag[idx] = f;
|
||||
}
|
||||
|
||||
/* The faces of the changed cells, component c (a face is its two cells' face). */
|
||||
extern "C" __global__ void e3_mask_face_flags(MaskGrid g, int c, const u8* __restrict__ changed, u8* __restrict__ flag)
|
||||
{
|
||||
long long f = (long long) blockIdx.x * blockDim.x + threadIdx.x;
|
||||
long long ni = g.nx + (c == 0), nj = g.ny + (c == 1), nk = g.nz + (c == 2);
|
||||
if (f >= ni * nj * nk) return;
|
||||
long long k = f / (nj * ni), j = (f / ni) % nj, i = f % ni;
|
||||
long long nx = g.nx, nxy = (long long) g.nx * g.ny;
|
||||
long long cell = (k * g.ny + j) * nx + i;
|
||||
long long own = c == 0 ? i : (c == 1 ? j : k);
|
||||
long long lim = c == 0 ? g.nx : (c == 1 ? g.ny : g.nz);
|
||||
long long step = c == 0 ? 1 : (c == 1 ? nx : nxy);
|
||||
u8 m = 0;
|
||||
if (own < lim) m |= changed[cell];
|
||||
if (own > 0) m |= changed[cell - step];
|
||||
flag[f] = m;
|
||||
}
|
||||
|
||||
/* Per changed cell: fluid at the new geometry, the space-time activity (fluid
|
||||
* at either end), the GCL entry `(V^{n+1} − V^n) dv / dt` of an active cell;
|
||||
* the device tables written (activity: step and predictor sets). Out flags:
|
||||
* bit 0 fluid, bit 1 active. */
|
||||
extern "C" __global__ void e3_mask_cells(
|
||||
int n, const u32* __restrict__ list, const double* __restrict__ vol, const double* __restrict__ vol_prev,
|
||||
double dv, double dt, int* __restrict__ active, int* __restrict__ active_pred,
|
||||
u8* __restrict__ out_flags, double* __restrict__ out_entry)
|
||||
{
|
||||
int t = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (t >= n) return;
|
||||
u32 idx = list[t];
|
||||
int fn = vol[idx] > 0.0;
|
||||
int fo = vol_prev[idx] > 0.0;
|
||||
int act = fn || fo;
|
||||
active[idx] = act;
|
||||
active_pred[idx] = fn;
|
||||
out_flags[t] = (u8) (fn | (act << 1));
|
||||
out_entry[t] = act ? (vol[idx] - vol_prev[idx]) * dv / dt : 0.0;
|
||||
}
|
||||
|
||||
/* lat.cell: the cell at (i, j, k) with the periodic wrap in z; -1 outside. */
|
||||
__device__ __forceinline__ long long lat_cell(const MaskGrid& g, long long i, long long j, long long k)
|
||||
{
|
||||
if (i < 0 || i >= g.nx || j < 0 || j >= g.ny) return -1;
|
||||
if (g.periodic) {
|
||||
k %= g.nz;
|
||||
if (k < 0) k += g.nz;
|
||||
} else if (k < 0 || k >= g.nz) {
|
||||
return -1;
|
||||
}
|
||||
return (k * g.ny + j) * g.nx + i;
|
||||
}
|
||||
|
||||
__device__ __forceinline__ int is_small(long long q, const double* vol, const double* vol_prev, const int* active, double threshold)
|
||||
{
|
||||
return active[q] && fmax(vol[q], vol_prev[q]) < threshold;
|
||||
}
|
||||
|
||||
/* Per changed cell: the master (the active, not small face neighbour of
|
||||
* largest new volume, the first on ties) of a small cell. */
|
||||
extern "C" __global__ void e3_mask_merge(
|
||||
MaskGrid g, int n, const u32* __restrict__ list, const double* __restrict__ vol,
|
||||
const double* __restrict__ vol_prev, const int* __restrict__ active, double threshold,
|
||||
u32* __restrict__ owner, u32* __restrict__ out_master)
|
||||
{
|
||||
int t = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (t >= n) return;
|
||||
long long idx = list[t];
|
||||
long long nxy = (long long) g.nx * g.ny;
|
||||
long long k = idx / nxy, j = (idx % nxy) / g.nx, i = idx % g.nx;
|
||||
u32 master = 0xffffffffu;
|
||||
if (is_small(idx, vol, vol_prev, active, threshold)) {
|
||||
int have = 0;
|
||||
double best = 0.0;
|
||||
long long bi = 0;
|
||||
for (int d = 0; d < 3; ++d) {
|
||||
for (int side = -1; side <= 1; side += 2) {
|
||||
long long q = lat_cell(g, i + (d == 0) * side, j + (d == 1) * side, k + (d == 2) * side);
|
||||
if (q < 0) continue;
|
||||
double v = vol[q];
|
||||
if (active[q] && !is_small(q, vol, vol_prev, active, threshold) && (!have || v > best)) {
|
||||
have = 1;
|
||||
best = v;
|
||||
bi = q;
|
||||
}
|
||||
}
|
||||
}
|
||||
if (have) master = (u32) bi;
|
||||
}
|
||||
owner[idx] = master == 0xffffffffu ? (u32) idx : master;
|
||||
out_master[t] = master;
|
||||
}
|
||||
|
||||
/* lat.face: the index of the face of component c at (i, j, k), -1 outside
|
||||
* (periodic z wraps by nz whatever the component). */
|
||||
__device__ __forceinline__ long long lat_face(const MaskGrid& g, int c, long long i, long long j, long long k)
|
||||
{
|
||||
long long ni = g.nx + (c == 0), nj = g.ny + (c == 1), nk = g.nz + (c == 2);
|
||||
if (i < 0 || i >= ni || j < 0 || j >= nj) return -1;
|
||||
if (g.periodic) {
|
||||
k %= g.nz;
|
||||
if (k < 0) k += g.nz;
|
||||
} else if (k < 0 || k >= nk) {
|
||||
return -1;
|
||||
}
|
||||
if (c == 0) return (k * g.ny + j) * (g.nx + 1) + i;
|
||||
if (c == 1) return (k * (g.ny + 1) + j) * g.nx + i;
|
||||
return (k * g.ny + j) * g.nx + i;
|
||||
}
|
||||
|
||||
struct Apertures {
|
||||
const double* a[3];
|
||||
};
|
||||
|
||||
/* aperture(c, p).unwrap_or(dflt) */
|
||||
__device__ __forceinline__ double ap_or(const MaskGrid& g, const Apertures& A, int c, const long long p[3], double dflt)
|
||||
{
|
||||
long long f = lat_face(g, c, p[0], p[1], p[2]);
|
||||
return f < 0 ? dflt : A.a[c][f];
|
||||
}
|
||||
|
||||
/* Per face of a changed cell, component c: the kind (bit 0: fluid), the step
|
||||
* aperture ½(αⁿ⁺¹ + αⁿ) and its open flag (bit 1), the open-part centroid
|
||||
* shift (`compute_face_shifts`, with `has_shift`). */
|
||||
extern "C" __global__ void e3_mask_faces(
|
||||
MaskGrid g, int c, int n, const u32* __restrict__ list, Apertures A, const double* __restrict__ a_prev,
|
||||
double* __restrict__ a_step, int* __restrict__ open, int* __restrict__ open_pred, int has_shift,
|
||||
double* __restrict__ shift, u8* __restrict__ out_flags, double* __restrict__ out_a, double* __restrict__ out_shift)
|
||||
{
|
||||
int t = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (t >= n) return;
|
||||
long long f = list[t];
|
||||
long long ni = g.nx + (c == 0), nj = g.ny + (c == 1);
|
||||
long long k = f / (nj * ni), j = (f / ni) % nj, i = f % ni;
|
||||
double a_new = A.a[c][f];
|
||||
int side;
|
||||
if (c == 0) side = i == 0 || i == g.nx;
|
||||
else if (c == 1) side = j == 0 || j == g.ny;
|
||||
else side = !g.periodic && (k == 0 || k == g.nz);
|
||||
int fluid = side || a_new > 0.0;
|
||||
double as = 0.5 * (a_new + a_prev[f]);
|
||||
int op = as > 0.0;
|
||||
a_step[f] = as;
|
||||
open[f] = op;
|
||||
open_pred[f] = fluid;
|
||||
out_flags[t] = (u8) (fluid | (op << 1));
|
||||
out_a[t] = as;
|
||||
if (!has_shift) return;
|
||||
double out[3] = {0.0, 0.0, 0.0};
|
||||
long long p[3] = {i, j, k};
|
||||
double h[3] = {g.dx, g.dy, g.dz};
|
||||
long long nlim[3] = {g.nx, g.ny, g.nz};
|
||||
long long fa = lat_face(g, c, p[0], p[1], p[2]);
|
||||
/* `self.aperture(c, p)` is Some for a face inside the grid */
|
||||
double alpha = A.a[c][fa];
|
||||
int on_side = p[c] == 0 || p[c] == nlim[c];
|
||||
if (alpha > 0.0 && alpha < 1.0 && !(on_side && !(c == 2 && g.periodic))) {
|
||||
/* cv_geometry(c, p): the control volume's side apertures and wall */
|
||||
double area[3] = {g.dy * g.dz, g.dx * g.dz, g.dx * g.dy};
|
||||
long long cm[3] = {p[0], p[1], p[2]};
|
||||
cm[c] -= 1;
|
||||
double ap[3][2];
|
||||
for (int d = 0; d < 3; ++d) {
|
||||
if (d == c) {
|
||||
long long q[3] = {p[0], p[1], p[2]};
|
||||
q[c] -= 1;
|
||||
double am = ap_or(g, A, c, q, alpha);
|
||||
q[c] += 2;
|
||||
double apl = ap_or(g, A, c, q, alpha);
|
||||
ap[d][0] = 0.5 * (am + alpha);
|
||||
ap[d][1] = 0.5 * (alpha + apl);
|
||||
} else {
|
||||
long long cmd[3] = {cm[0], cm[1], cm[2]};
|
||||
long long cpd[3] = {p[0], p[1], p[2]};
|
||||
cmd[d] += 1;
|
||||
cpd[d] += 1;
|
||||
double minus = 0.5 * (ap_or(g, A, d, cm, 1.0) + ap_or(g, A, d, p, 1.0));
|
||||
double plus = 0.5 * (ap_or(g, A, d, cmd, 1.0) + ap_or(g, A, d, cpd, 1.0));
|
||||
ap[d][0] = minus;
|
||||
ap[d][1] = plus;
|
||||
}
|
||||
}
|
||||
double nw[3];
|
||||
for (int d = 0; d < 3; ++d) nw[d] = -(ap[d][1] - ap[d][0]) * area[d];
|
||||
nw[c] = 0.0;
|
||||
double a = sqrt(nw[0] * nw[0] + nw[1] * nw[1] + nw[2] * nw[2]);
|
||||
if (a != 0.0) {
|
||||
for (int d = 0; d < 3; ++d) out[d] = -0.5 * h[d] * (1.0 - alpha) * nw[d] / a;
|
||||
}
|
||||
}
|
||||
for (int d = 0; d < 3; ++d) {
|
||||
shift[3 * f + d] = out[d];
|
||||
out_shift[3 * t + d] = out[d];
|
||||
}
|
||||
}
|
||||
|
||||
/* The GCL table's wall cells: active with a nonzero wall vector. */
|
||||
extern "C" __global__ void e3_mask_wall_flags(long long nc, const int* __restrict__ active, const double* __restrict__ wall, u8* __restrict__ flag)
|
||||
{
|
||||
long long idx = (long long) blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (idx >= nc) return;
|
||||
const double* w = wall + 3 * idx;
|
||||
flag[idx] = active[idx] && (w[0] != 0.0 || w[1] != 0.0 || w[2] != 0.0);
|
||||
}
|
||||
|
||||
/* The wall areas of the listed cells. */
|
||||
extern "C" __global__ void e3_mask_wall_area(int n, const u32* __restrict__ list, const double* __restrict__ wall, double* __restrict__ out)
|
||||
{
|
||||
int t = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (t >= n) return;
|
||||
const double* w = wall + 3 * (long long) list[t];
|
||||
out[t] = sqrt(w[0] * w[0] + w[1] * w[1] + w[2] * w[2]);
|
||||
}
|
||||
|
||||
/* The GCL wall-flux table: the entries of the changed active cells, less the
|
||||
* correction per unit wall area on every cell with a wall. */
|
||||
extern "C" __global__ void e3_mask_wall_flux(
|
||||
long long nc, const u8* __restrict__ changed, const int* __restrict__ active, const double* __restrict__ vol,
|
||||
const double* __restrict__ vol_prev, const double* __restrict__ wall, double dv, double dt, double correction,
|
||||
double* __restrict__ table)
|
||||
{
|
||||
long long idx = (long long) blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (idx >= nc) return;
|
||||
double v = 0.0;
|
||||
if (changed[idx] && active[idx]) v = (vol[idx] - vol_prev[idx]) * dv / dt;
|
||||
if (correction != 0.0) {
|
||||
const double* w = wall + 3 * idx;
|
||||
if (w[0] != 0.0 || w[1] != 0.0 || w[2] != 0.0) {
|
||||
double a = sqrt(w[0] * w[0] + w[1] * w[1] + w[2] * w[2]);
|
||||
v -= correction * a;
|
||||
}
|
||||
}
|
||||
table[idx] = v;
|
||||
}
|
||||
|
||||
/* The merged cells: owner != self. */
|
||||
extern "C" __global__ void e3_mask_slave_flags(long long nc, const u32* __restrict__ owner, u8* __restrict__ flag)
|
||||
{
|
||||
long long idx = (long long) blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (idx >= nc) return;
|
||||
flag[idx] = owner[idx] != (u32) idx;
|
||||
}
|
||||
|
||||
/* The slave count per master (counts zeroed by the caller). */
|
||||
extern "C" __global__ void e3_mask_fold_count(int n, const u32* __restrict__ slaves, const u32* __restrict__ owner, u32* __restrict__ counts)
|
||||
{
|
||||
int t = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (t >= n) return;
|
||||
atomicAdd(&counts[owner[slaves[t]]], 1u);
|
||||
}
|
||||
|
||||
/* The CSR's slave lists, ascending per master: a slave's rank = its master's
|
||||
* slaves (face neighbours of the master, distinct) of smaller index. */
|
||||
extern "C" __global__ void e3_mask_fold_fill(
|
||||
MaskGrid g, int n, const u32* __restrict__ slaves, const u32* __restrict__ owner, const u32* __restrict__ fold_ptr,
|
||||
u32* __restrict__ fold_idx)
|
||||
{
|
||||
int t = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (t >= n) return;
|
||||
long long s = slaves[t];
|
||||
long long m = owner[s];
|
||||
long long nxy = (long long) g.nx * g.ny;
|
||||
long long k = m / nxy, j = (m % nxy) / g.nx, i = m % g.nx;
|
||||
long long seen[6];
|
||||
int ns = 0;
|
||||
u32 rank = 0;
|
||||
for (int d = 0; d < 3; ++d) {
|
||||
for (int side = -1; side <= 1; side += 2) {
|
||||
long long q = lat_cell(g, i + (d == 0) * side, j + (d == 1) * side, k + (d == 2) * side);
|
||||
if (q < 0 || q == m) continue;
|
||||
int dup = 0;
|
||||
for (int r = 0; r < ns; ++r) dup |= seen[r] == q;
|
||||
if (dup) continue;
|
||||
seen[ns++] = q;
|
||||
if (q < s && owner[q] == (u32) m) ++rank;
|
||||
}
|
||||
}
|
||||
fold_idx[fold_ptr[m] + rank] = (u32) s;
|
||||
}
|
||||
|
||||
/* The imposition band's faces: |d| <= band. */
|
||||
extern "C" __global__ void e3_mask_band_flags(long long n, const double* __restrict__ d, double band, u8* __restrict__ flag)
|
||||
{
|
||||
long long f = (long long) blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (f >= n) return;
|
||||
flag[f] = fabs(d[f]) <= band;
|
||||
}
|
||||
Reference in New Issue
Block a user