sfx-sheba.cu

#include <cuda.h>
#include <cuda_runtime_api.h>

#include "sfx-sheba.h"
#include "sfx-model-compute-subfunc.cuh"
#include "sfx-surface.cuh"
#include "sfx-memory-processing.cuh"

namespace sfx_kernel
{
    template<typename T>
    __global__ void compute_flux(sfxDataVecTypeC sfx,
        meteoDataVecTypeC meteo,
        const sfx_sheba_param_C model, 
        const sfx_surface_param surface,
        const sfx_sheba_numericsType_C numerics,
        const sfx_phys_constants phys,
        const int grid_size);
}

template<typename T>
__global__ void sfx_kernel::compute_flux(sfxDataVecTypeC sfx,
    meteoDataVecTypeC meteo,
    const sfx_sheba_param_C model, 
    const sfx_surface_param surface,
    const sfx_sheba_numericsType_C numerics,
    const sfx_phys_constants phys,
    const int grid_size)
{
    const int index = blockIdx.x * blockDim.x + threadIdx.x;
    T h, U, dT, Tsemi, dQ, z0_m;
    T z0_t, B, h0_m, h0_t, u_dyn0, Re, 
    zeta, Rib, Udyn, Tdyn, Qdyn, phi_m, phi_h,
    Km, Pr_t_inv, Cm, Ct;
    int surface_type;

    if(index < grid_size)
    {
        U = meteo.U[index];
        Tsemi = meteo.Tsemi[index];
        dT = meteo.dT[index];
        dQ = meteo.dQ[index];
        h = meteo.h[index];
        z0_m = meteo.z0_m[index];

        surface_type = z0_m < 0.0 ? surface.surface_ocean : surface.surface_land;

        if (surface_type == surface.surface_ocean) 
        {
            get_charnock_roughness(z0_m, u_dyn0, U, h, surface, numerics.maxiters_charnock);
            h0_m = h / z0_m;
        }
        if (surface_type == surface.surface_land) 
        {
            h0_m = h / z0_m;
            u_dyn0 = U * model.kappa / logf(h0_m);
        }

        Re = u_dyn0 * z0_m / phys.nu_air;
        get_thermal_roughness(z0_t, B, z0_m, Re, surface, surface_type);

        // --- define relative height [thermal]
        h0_t = h / z0_t;

        // --- define Ri-bulk
        Rib = (phys.g / Tsemi) * h * (dT + 0.61e0 * Tsemi * dQ) / (U*U);

        // --- get the fluxes
        // ----------------------------------------------------------------------------
        get_dynamic_scales(Udyn, Tdyn, Qdyn, zeta, 
            U, Tsemi, dT, dQ, h, z0_m, z0_t, (phys.g / Tsemi), model, 10);
        // ----------------------------------------------------------------------------

        get_phi(phi_m, phi_h, zeta, model);
        // ----------------------------------------------------------------------------

        // --- define transfer coeff. (momentum) & (heat)
        Cm = 0.0;
        if (U > 0.0)
            Cm = Udyn / U;
        Ct = 0.0;
        if (fabsf(dT) > 0.0) 
            Ct = Tdyn / dT;

        // --- define eddy viscosity & inverse Prandtl number
        Km = model.kappa * Cm * U * h / phi_m;
        Pr_t_inv = phi_m / phi_h;

        sfx.zeta[index]         = zeta;
        sfx.Rib[index]          = Rib;
        sfx.Re[index]           = Re;
        sfx.B[index]            = B;
        sfx.z0_m[index]         = z0_m;
        sfx.z0_t[index]         = z0_t;
        sfx.Rib_conv_lim[index] = T(0.0);
        sfx.Cm[index]           = Cm;
        sfx.Ct[index]           = Ct;
        sfx.Km[index]           = Km;
        sfx.Pr_t_inv[index]     = Pr_t_inv;
    }
}

template<typename T, MemType memIn, MemType memOut >
void FluxSheba<T, memIn, memOut, MemType::GPU>::compute_flux()
{
    const int BlockCount = int(ceil(float(grid_size) / 1024.0));
    dim3 cuBlock = dim3(1024, 1, 1);
	dim3 cuGrid = dim3(BlockCount, 1, 1);

    sfx_kernel::compute_flux<T><<<cuGrid, cuBlock>>>(sfx, meteo, model, 
                                                    surface, numerics, phys, grid_size);

    if(MemType::GPU != memOut)
    {
        const size_t new_size = grid_size * sizeof(T);
        memproc::memcopy<memOut, MemType::GPU>((void*&)res_sfx->zeta, (void*&)sfx.zeta, new_size);
        memproc::memcopy<memOut, MemType::GPU>((void*&)res_sfx->Rib, (void*&)sfx.Rib, new_size);
        memproc::memcopy<memOut, MemType::GPU>((void*&)res_sfx->Re, (void*&)sfx.Re, new_size);
        memproc::memcopy<memOut, MemType::GPU>((void*&)res_sfx->B, (void*&)sfx.B, new_size);
        memproc::memcopy<memOut, MemType::GPU>((void*&)res_sfx->z0_m, (void*&)sfx.z0_m, new_size);
        memproc::memcopy<memOut, MemType::GPU>((void*&)res_sfx->z0_t, (void*&)sfx.z0_t, new_size);
        memproc::memcopy<memOut, MemType::GPU>((void*&)res_sfx->Rib_conv_lim, (void*&)sfx.Rib_conv_lim, new_size);
        memproc::memcopy<memOut, MemType::GPU>((void*&)res_sfx->Cm, (void*&)sfx.Cm, new_size);
        memproc::memcopy<memOut, MemType::GPU>((void*&)res_sfx->Ct, (void*&)sfx.Ct, new_size);
        memproc::memcopy<memOut, MemType::GPU>((void*&)res_sfx->Km, (void*&)sfx.Km, new_size);
        memproc::memcopy<memOut, MemType::GPU>((void*&)res_sfx->Pr_t_inv, (void*&)sfx.Pr_t_inv, new_size);
    }
}

template class FluxShebaBase<float, MemType::GPU, MemType::GPU, MemType::GPU>;
template class FluxShebaBase<float, MemType::GPU, MemType::GPU, MemType::CPU>;
template class FluxShebaBase<float, MemType::GPU, MemType::CPU, MemType::GPU>;
template class FluxShebaBase<float, MemType::CPU, MemType::GPU, MemType::GPU>;
template class FluxShebaBase<float, MemType::CPU, MemType::CPU, MemType::GPU>;
template class FluxShebaBase<float, MemType::CPU, MemType::GPU, MemType::CPU>;
template class FluxShebaBase<float, MemType::GPU, MemType::CPU, MemType::CPU>;

template class FluxSheba<float, MemType::GPU, MemType::GPU, MemType::GPU>;
template class FluxSheba<float, MemType::GPU, MemType::GPU, MemType::CPU>;
template class FluxSheba<float, MemType::GPU, MemType::CPU, MemType::GPU>;
template class FluxSheba<float, MemType::CPU, MemType::GPU, MemType::GPU>;
template class FluxSheba<float, MemType::CPU, MemType::CPU, MemType::GPU>;
template class FluxSheba<float, MemType::CPU, MemType::GPU, MemType::CPU>;
template class FluxSheba<float, MemType::GPU, MemType::CPU, MemType::CPU>;