CelerisLab/tests/validation/multigrid/spacetime_gpu.cu

// Test-only arrival-f space-time coupling. No production scheduler changes.
#include <stdint.h>
#include <math.h>
#include "config.h"
#include "core/precision.cuh"
#include "core/descriptors.cuh"
#include "core/layout.cuh"
#include "operators/macro.cuh"
#include "operators/collision_srt.cuh"
#include "streaming/pull_double_buffer.cuh"

#if DIM != 2 || NQ != 9 || COLLISION_MODEL != 0 || STORE_PRECISION != 0 || USE_DDF_SHIFTING || USE_LES || STREAMING_MODEL != 0
#error "Space-time experiment requires unshifted FP32 D2Q9 SRT double buffers without LES"
#endif

__device__ __forceinline__ void collide(float* f, float omega) {
    float rho, ux, uy, feq[NQ];
    compute_rho_u(f, rho, ux, uy);
    compute_feq(rho, ux, uy, feq);
    collide_srt_no_force(f, feq, omega);
}

__device__ __forceinline__ void health(float value, unsigned int* counts) {
    if (!isfinite(value)) atomicAdd(counts + 4, 1u);
    if (value <= 0.0f) atomicAdd(counts + 5, 1u);
}

extern "C" __global__ void CoarseStart(
    const float* coarse, const float* fine, float* out, float omega,
    unsigned int x0, unsigned int y0, unsigned int x1, unsigned int y1,
    unsigned int* counts)
{
    unsigned int x, y; unsigned long k;
    index_from_thread(x, y, k);
    if (x >= NX || y >= NY) return;
    const bool covered = x >= x0 && x < x1 && y >= y0 && y < y1;
    float f[NQ];
    for (unsigned int q = 0; q < NQ; ++q) {
        const unsigned long s = q * 4ul * TOTAL_CELLS + 2ul * y * (2u * NX) + 2u * x;
        f[q] = covered ? 0.25f * (fine[s] + fine[s+1] + fine[s+2u*NX] + fine[s+2u*NX+1])
                       : coarse[index_f(k, q)];
    }
    collide(f, omega);
    for (int q = 0; q < NQ; ++q) if (!covered) health(f[q], counts);
    stream_pull_store(k, f, out);
}

extern "C" __global__ void Collide(
    const float* in, float* out, const unsigned char* active, float omega,
    unsigned int* counts)
{
    unsigned int x, y; unsigned long k;
    index_from_thread(x, y, k);
    if (x >= NX || y >= NY) return;
    float f[NQ];
    if (active == nullptr || active[k]) {
        for (int q = 0; q < NQ; ++q) f[q] = in[index_f(k, q)];
        collide(f, omega);
        if (active != nullptr) for (int q = 0; q < NQ; ++q) health(f[q], counts);
    } else {
        for (int q = 0; q < NQ; ++q) f[q] = NAN;
    }
    stream_pull_store(k, f, out);
}

extern "C" __global__ void Pull(
    const float* in, float* out, const unsigned char* active, unsigned int* counts)
{
    unsigned int x, y; unsigned long k;
    index_from_thread(x, y, k);
    if (x >= NX || y >= NY) return;
    float f[NQ];
    if (active == nullptr || active[k]) {
        unsigned long j[NQ];
        compute_neighbors(x, y, j);
        stream_pull_load(k, f, in, j);
        if (active != nullptr) for (int q = 0; q < NQ; ++q) health(f[q], counts);
    } else {
        for (int q = 0; q < NQ; ++q) f[q] = NAN;
    }
    stream_pull_store(k, f, out);
}

extern "C" __global__ void Interpolate(
    const float* g0, const float* gp, const uint32_t* samples,
    const uint32_t* origins, const int* sizes, const double* weights,
    double* pre, float* h, unsigned int groups, unsigned int frozen,
    unsigned int* counts)
{
    unsigned int group = blockIdx.x * blockDim.x + threadIdx.x;
    if (group >= groups) return;
    const int size = sizes[group];
    for (int child = 0; child < 4; ++child) {
        double value = frozen ? (double)g0[origins[group]] : 0.0;
        if (!frozen) for (int j = 0; j < size; ++j) {
            const float* source = j < size/2 ? g0 : gp;
            value += weights[(4u*group+child)*10u+j] * (double)source[samples[10u*group+j]];
        }
        const float stored = __double2float_rn(value);
        pre[4u*group+child] = value;
        h[4u*group+child] = stored;
        if (!isfinite(value)) atomicAdd(counts, 1u);
        if (value <= 0.0) atomicAdd(counts+1, 1u);
        if (!isfinite(stored)) atomicAdd(counts+2, 1u);
        if (stored <= 0.0f) atomicAdd(counts+3, 1u);
    }
}

extern "C" __global__ void Inject(
    float* fine, const float* h, const uint32_t* rows, unsigned int count,
    unsigned int omit_first)
{
    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
    if (row >= count || (omit_first && row == 0)) return;
    fine[rows[2u*row]] = h[rows[2u*row+1]];
}

extern "C" __global__ void Register(
    const float* predicted, const float* fine, const float* h,
    const uint32_t* targets, const uint32_t* sources, const int* times,
    double* registers, unsigned int count, int substep, int mutation)
{
    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
    if (row >= count) return;
    double value = substep == 0 ? -(double)predicted[targets[row]] : registers[row];
    for (int j = 0; j < 4; ++j) {
        const unsigned int slot = 4u*row+j;
        const int when = times[slot];
        if (when != substep && !(substep == 0 && when == -1)) continue;
        if ((mutation == 1 && when == -1) || (mutation == 2 && when == 0)
                || (mutation == 3 && when == 1)) continue;
        value += 0.25 * (double)(when == -1 ? h[sources[slot]] : fine[sources[slot]]);
    }
    registers[row] = value;
}

extern "C" __global__ void ApplyRegisters(
    float* predicted, const double* registers, const uint32_t* targets,
    unsigned int count, unsigned int* counts)
{
    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
    if (row >= count) return;
    const unsigned int slot = targets[row];
    predicted[slot] = __double2float_rn((double)predicted[slot] + registers[row]);
    health(predicted[slot], counts);
}