CelerisLab/src/CelerisLab/lbm/kernels/step/one_step_double.cu

// CelerisLab – step/one_step_double.cu
// Double-buffer PULL streaming step kernel.
// Included by kernel_v2.cu.  No standalone compilation.
// ============================================================================

#ifndef CELERIS_STEP_ONE_STEP_DOUBLE_CU
#define CELERIS_STEP_ONE_STEP_DOUBLE_CU

#ifndef Y_WALL_BC
#define Y_WALL_BC 0
#endif

// ---- Helper: apply inlet/outlet/wall BC in pull-double mode ----
#if DIM == 2
__device__ __forceinline__ void apply_boundary_pull(
    float* f, uint16_t fl, unsigned int x, unsigned int y,
    const fpxx* fi_in, unsigned long k)
{
    // Curved-BC nodes are handled by CurvedBoundaryKernel before OneStep (see stepper).
    if (is_curved(fl)) return;

    bool interior_y = (y > 0u) && (y < (unsigned int)(NY - 1));

    if (is_inlet(fl) && interior_y) {
        apply_inlet_pull_d2q9(f, x, y, fi_in);
    }
    else if (is_outlet(fl) && interior_y) {
        apply_outlet_pull_d2q9(f, x, y, fi_in);
    }
    else {
        bounce_back_swap(f);
    }
}
#endif // DIM == 2

#if DIM == 3
__device__ __forceinline__ void apply_boundary_pull_3d(
    float* f, uint16_t fl, unsigned int x, unsigned int y, unsigned int z,
    const fpxx* fi_in, unsigned long k)
{
    // Curved-BC nodes are handled by CurvedBoundaryKernel before OneStep (see stepper).
    if (is_curved(fl)) return;

    bool interior_y = (y > 0u) && (y < (unsigned int)(NY - 1));

    if (is_inlet(fl) && interior_y) {
        apply_inlet_pull_d3q19(f, x, y, z, fi_in);
    }
    else if (is_outlet(fl) && interior_y) {
        apply_outlet_pull_d3q19(f, x, y, z, fi_in);
    }
    else {
        bounce_back_swap(f);
    }
}
#endif

// ============================================================================
// OneStep — main pull double-buffer step kernel
// ============================================================================
__global__ __launch_bounds__(NT, 2)
void OneStep(
    uint16_t* flag, fpxx* fi_in, fpxx* fi_out)
{
#if DIM == 2
    unsigned int x, y; unsigned long k;
    index_from_thread(x, y, k);
    if (x >= (unsigned int)NX || y >= (unsigned int)NY) return;

    uint16_t fl = flag[k];
    unsigned long j[NQ];
    const bool interior_y = (y > 0u) && (y < (unsigned int)(NY - 1));
    compute_neighbors(x, y, j);

    float f[NQ];
    stream_pull_load(k, f, fi_in, j);

    // Boundary handling
    // has_bc(fl): BC subtype != NONE -> wall/inlet/outlet nodes only
    // Obstacle nodes (FLAG_OBSTACLE, BC subtype=NONE) are excluded here
    // and handled solely by the is_obstacle branch below.
    // TODO: BC_MOVING and BC_PERIODIC are defined in flags.cuh but NOT
    //       dispatched here — they silently fall through to bounce_back_swap.
    if (is_solid(fl) && has_bc(fl))
        apply_boundary_pull(f, fl, x, y, fi_in, k);

    if (is_obstacle(fl) && !is_curved(fl))
        bounce_back_swap(f);

    if (is_fluid(fl) && (y == 1u || y == (unsigned int)(NY - 2))) {
#if Y_WALL_BC == 2
        // Reconstruct f(t+dt) on the physical plane, then collide below.
        // This replaces, rather than compounds, the halfway wall correction.
        if (y == 1u) reconstruct_bottom_velocity(f, 0.0f, bottom_wall_normal_velocity(x));
#if TOP_WALL_BC == 1
        else reconstruct_top_reservoir(f, RHO);
#else
        else reconstruct_top_velocity(f, U0, TOP_WALL_NORMAL_VELOCITY);
#endif
#elif Y_WALL_BC == 1
        apply_wall_freeslip_d2q9_y_pull(y, f, fi_in, k);
#else
        apply_wall_bb_d2q9_y_pull(y, f, fi_in, k);
#endif
    }

    // Collision
    // Normal path: fluid nodes collide.
    // Path A repair: when using the local Zou-He inlet, the west inlet nodes are
    // still tagged SOLID|BC_INLET ghost nodes, but their post-BC state is later
    // used as a pull source for x=1. Leaving that ghost state uncollided was the
    // main donor/ghost semantic bug behind the observed inlet blow-ups.
    const bool collide_inlet_ghost = is_inlet(fl) && interior_y
                                   && inlet_scheme_uses_post_collision_ghost();
    if (is_fluid(fl) || collide_inlet_ghost) {
        float rho_n, ux, uy;
        compute_rho_u(f, rho_n, ux, uy);
        collide_dispatch(f, rho_n, ux, uy);
    }

    stream_pull_store(k, f, fi_out);

#elif DIM == 3
    unsigned int x, y, z; unsigned long k;
    index_from_thread(x, y, z, k);
    if (x >= (unsigned int)NX || y >= (unsigned int)NY || z >= (unsigned int)NZ) return;

    uint16_t fl = flag[k];
    unsigned long j[NQ];
    const bool interior_y = (y > 0u) && (y < (unsigned int)(NY - 1));
    compute_neighbors(k, j);

    float f[NQ];
    stream_pull_load(k, f, fi_in, j);

    if (is_solid(fl) && has_bc(fl))
        apply_boundary_pull_3d(f, fl, x, y, z, fi_in, k);

    if (is_obstacle(fl) && !is_curved(fl))
        bounce_back_swap(f);

    if (is_fluid(fl) && (y == 1u || y == (unsigned int)(NY - 2)))
        apply_wall_bb_d3q19_y_pull(y, f, fi_in, k);

    const bool collide_inlet_ghost = is_inlet(fl) && interior_y
                                   && inlet_scheme_uses_post_collision_ghost();
    if (is_fluid(fl) || collide_inlet_ghost) {
        float rho_n, ux, uy, uz;
        compute_rho_u(f, rho_n, ux, uy, uz);
        collide_dispatch(f, rho_n, ux, uy, uz);
    }

    stream_pull_store(k, f, fi_out);
#endif
}

#endif // CELERIS_STEP_ONE_STEP_DOUBLE_CU