From cfbf07c80031e4f604280d6141786d40a080f861 Mon Sep 17 00:00:00 2001 From: magpowell Date: Mon, 4 Nov 2024 13:42:25 -0800 Subject: [PATCH 01/11] add liq or ice cloud optics flags --- src_cuda_rt/Cloud_optics_rt.cu | 77 ++++++++++++++++++++-------------- src_test/test_rte_rrtmgp_rt.cu | 30 +++++++++---- 2 files changed, 67 insertions(+), 40 deletions(-) diff --git a/src_cuda_rt/Cloud_optics_rt.cu b/src_cuda_rt/Cloud_optics_rt.cu index c5a3f90a..06c88ec8 100644 --- a/src_cuda_rt/Cloud_optics_rt.cu +++ b/src_cuda_rt/Cloud_optics_rt.cu @@ -201,12 +201,16 @@ void Cloud_optics_rt::cloud_optics( dim3 block_m_gpu(block_col_m, block_lay_m); Array_gpu liqmsk({ncol, nlay}); - set_mask<<>>( - ncol, nlay, mask_min_value, liqmsk.ptr(), clwp.ptr()); - + if (clwp.ptr() != nullptr){ + set_mask<<>>( + ncol, nlay, mask_min_value, liqmsk.ptr(), clwp.ptr()); + } Array_gpu icemsk({ncol, nlay}); - set_mask<<>>( - ncol, nlay, mask_min_value, icemsk.ptr(), ciwp.ptr()); + if (ciwp.ptr() != nullptr){ + set_mask<<>>( + ncol, nlay, mask_min_value, icemsk.ptr(), ciwp.ptr()); + } + // Temporary arrays for storage. Array_gpu ltau ({ncol, nlay}); @@ -227,19 +231,22 @@ void Cloud_optics_rt::cloud_optics( dim3 block_gpu(block_col, block_lay); // Liquid water - compute_from_table_kernel<<>>( - ncol, nlay, ibnd-1, liqmsk.ptr(), clwp.ptr(), reliq.ptr(), - this->liq_nsteps, this->liq_step_size, this->radliq_lwr, - this->lut_extliq_gpu.ptr(), this->lut_ssaliq_gpu.ptr(), - this->lut_asyliq_gpu.ptr(), ltau.ptr(), ltaussa.ptr(), ltaussag.ptr()); + if (clwp.ptr() != nullptr){ + compute_from_table_kernel<<>>( + ncol, nlay, ibnd-1, liqmsk.ptr(), clwp.ptr(), reliq.ptr(), + this->liq_nsteps, this->liq_step_size, this->radliq_lwr, + this->lut_extliq_gpu.ptr(), this->lut_ssaliq_gpu.ptr(), + this->lut_asyliq_gpu.ptr(), ltau.ptr(), ltaussa.ptr(), ltaussag.ptr()); + } // Ice. - compute_from_table_kernel<<>>( - ncol, nlay, ibnd-1, icemsk.ptr(), ciwp.ptr(), reice.ptr(), - this->ice_nsteps, this->ice_step_size, this->radice_lwr, - this->lut_extice_gpu.ptr(), this->lut_ssaice_gpu.ptr(), - this->lut_asyice_gpu.ptr(), itau.ptr(), itaussa.ptr(), itaussag.ptr()); - + if (ciwp.ptr() != nullptr){ + compute_from_table_kernel<<>>( + ncol, nlay, ibnd-1, icemsk.ptr(), ciwp.ptr(), reice.ptr(), + this->ice_nsteps, this->ice_step_size, this->radice_lwr, + this->lut_extice_gpu.ptr(), this->lut_ssaice_gpu.ptr(), + this->lut_asyice_gpu.ptr(), itau.ptr(), itaussa.ptr(), itaussag.ptr()); + } constexpr Float eps = std::numeric_limits::epsilon(); combine_and_store_kernel<<>>( @@ -274,13 +281,15 @@ void Cloud_optics_rt::cloud_optics( dim3 block_m_gpu(block_col_m, block_lay_m); Array_gpu liqmsk({ncol, nlay}); - set_mask<<>>( - ncol, nlay, mask_min_value, liqmsk.ptr(), clwp.ptr()); - + if (clwp.ptr() != nullptr){ + set_mask<<>>( + ncol, nlay, mask_min_value, liqmsk.ptr(), clwp.ptr()); + } Array_gpu icemsk({ncol, nlay}); - set_mask<<>>( - ncol, nlay, mask_min_value, icemsk.ptr(), ciwp.ptr()); - + if (ciwp.ptr() != nullptr){ + set_mask<<>>( + ncol, nlay, mask_min_value, icemsk.ptr(), ciwp.ptr()); + } // Temporary arrays for storage. Array_gpu ltau ({ncol, nlay}); Array_gpu ltaussa ({ncol, nlay}); @@ -300,18 +309,22 @@ void Cloud_optics_rt::cloud_optics( dim3 block_gpu(block_col, block_lay); // Liquid water - compute_from_table_kernel<<>>( - ncol, nlay, ibnd-1, liqmsk.ptr(), clwp.ptr(), reliq.ptr(), - this->liq_nsteps, this->liq_step_size, this->radliq_lwr, - this->lut_extliq_gpu.ptr(), this->lut_ssaliq_gpu.ptr(), - this->lut_asyliq_gpu.ptr(), ltau.ptr(), ltaussa.ptr(), ltaussag.ptr()); + if (clwp.ptr() != nullptr){ + compute_from_table_kernel<<>>( + ncol, nlay, ibnd-1, liqmsk.ptr(), clwp.ptr(), reliq.ptr(), + this->liq_nsteps, this->liq_step_size, this->radliq_lwr, + this->lut_extliq_gpu.ptr(), this->lut_ssaliq_gpu.ptr(), + this->lut_asyliq_gpu.ptr(), ltau.ptr(), ltaussa.ptr(), ltaussag.ptr()); + } // Ice. - compute_from_table_kernel<<>>( - ncol, nlay, ibnd-1, icemsk.ptr(), ciwp.ptr(), reice.ptr(), - this->ice_nsteps, this->ice_step_size, this->radice_lwr, - this->lut_extice_gpu.ptr(), this->lut_ssaice_gpu.ptr(), - this->lut_asyice_gpu.ptr(), itau.ptr(), itaussa.ptr(), itaussag.ptr()); + if (ciwp.ptr() != nullptr){ + compute_from_table_kernel<<>>( + ncol, nlay, ibnd-1, icemsk.ptr(), ciwp.ptr(), reice.ptr(), + this->ice_nsteps, this->ice_step_size, this->radice_lwr, + this->lut_extice_gpu.ptr(), this->lut_ssaice_gpu.ptr(), + this->lut_asyice_gpu.ptr(), itau.ptr(), itaussa.ptr(), itaussag.ptr()); + } constexpr Float eps = std::numeric_limits::epsilon(); diff --git a/src_test/test_rte_rrtmgp_rt.cu b/src_test/test_rte_rrtmgp_rt.cu index fc7f204c..2e93cf9e 100644 --- a/src_test/test_rte_rrtmgp_rt.cu +++ b/src_test/test_rte_rrtmgp_rt.cu @@ -230,6 +230,8 @@ void solve_radiation(int argc, char** argv) {"raytracing" , { true, "Use raytracing for flux computation. '--raytracing 256': use 256 rays per pixel" }}, {"independent-column", { false, "run raytracer in independent column mode"}}, {"cloud-optics" , { false, "Enable cloud optics." }}, + {"liq-cloud-optics" , { false, "liquid only cloud optics." }}, + {"ice-cloud-optics" , { false, "ice only cloud optics." }}, {"cloud-mie" , { false, "Use Mie tables for cloud scattering in ray tracer" }}, {"aerosol-optics" , { false, "Enable aerosol optics." }}, {"single-gpt" , { false, "Output optical properties and fluxes for a single g-point. '--single-gpt 100': output 100th g-point" }}, @@ -250,6 +252,8 @@ void solve_radiation(int argc, char** argv) const bool switch_raytracing = command_line_switches.at("raytracing" ).first; const bool switch_independent_column= command_line_switches.at("independent-column").first; const bool switch_cloud_optics = command_line_switches.at("cloud-optics" ).first; + const bool switch_liq_cloud_optics = command_line_switches.at("liq-cloud-optics" ).first; + const bool switch_ice_cloud_optics = command_line_switches.at("ice-cloud-optics" ).first; const bool switch_cloud_mie = command_line_switches.at("cloud-mie" ).first; const bool switch_aerosol_optics = command_line_switches.at("aerosol-optics" ).first; const bool switch_single_gpt = command_line_switches.at("single-gpt" ).first; @@ -271,6 +275,11 @@ void solve_radiation(int argc, char** argv) throw std::runtime_error(error); } + if (switch_liq_cloud_optics && switch_ice_cloud_optics) { + std::string error = "Both liquid-only and ice-only cloud optics cannot be enabled simultaneously"; + throw std::runtime_error(error); + } + // Print the options to the screen. print_command_line_options(command_line_switches, command_line_ints); @@ -353,17 +362,22 @@ void solve_radiation(int argc, char** argv) if (switch_cloud_optics) { - lwp.set_dims({n_col, n_lay}); - lwp = std::move(input_nc.get_variable("lwp", {n_lay, n_col_y, n_col_x})); + + if(!switch_ice_cloud_optics){ + lwp.set_dims({n_col, n_lay}); + lwp = std::move(input_nc.get_variable("lwp", {n_lay, n_col_y, n_col_x})); - iwp.set_dims({n_col, n_lay}); - iwp = std::move(input_nc.get_variable("iwp", {n_lay, n_col_y, n_col_x})); + rel.set_dims({n_col, n_lay}); + rel = std::move(input_nc.get_variable("rel", {n_lay, n_col_y, n_col_x})); + } - rel.set_dims({n_col, n_lay}); - rel = std::move(input_nc.get_variable("rel", {n_lay, n_col_y, n_col_x})); + if(!switch_liq_cloud_optics){ + iwp.set_dims({n_col, n_lay}); + iwp = std::move(input_nc.get_variable("iwp", {n_lay, n_col_y, n_col_x})); - rei.set_dims({n_col, n_lay}); - rei = std::move(input_nc.get_variable("rei", {n_lay, n_col_y, n_col_x})); + rei.set_dims({n_col, n_lay}); + rei = std::move(input_nc.get_variable("rei", {n_lay, n_col_y, n_col_x})); + } } Array rh; From 8275529b6b304c0dafb35828cd5d83932c5f8fca Mon Sep 17 00:00:00 2001 From: magpowell Date: Wed, 6 Nov 2024 11:37:17 -0800 Subject: [PATCH 02/11] dont allocate temp arrays --- src_cuda_rt/Cloud_optics_rt.cu | 135 ++++++++++++++++++++++++++------- 1 file changed, 109 insertions(+), 26 deletions(-) diff --git a/src_cuda_rt/Cloud_optics_rt.cu b/src_cuda_rt/Cloud_optics_rt.cu index 06c88ec8..0546a630 100644 --- a/src_cuda_rt/Cloud_optics_rt.cu +++ b/src_cuda_rt/Cloud_optics_rt.cu @@ -87,6 +87,23 @@ namespace } } + __global__ + void combine_and_store_kernel_single_phase(const int ncol, const int nlay, const Float tmin, + Float* __restrict__ tau, + const Float* __restrict__ l_or_i_tau, const Float* __restrict__ l_or_i_taussa) + { + const int icol = blockIdx.x*blockDim.x + threadIdx.x; + const int ilay = blockIdx.y*blockDim.y + threadIdx.y; + + if ( (icol < ncol) && (ilay < nlay) ) + { + const int idx = icol + ilay*ncol; + const Float tau_t = (l_or_i_tau[idx] - l_or_i_taussa[idx]); + + tau[idx] = tau_t; + } + } + __global__ void combine_and_store_kernel(const int ncol, const int nlay, const Float tmin, Float* __restrict__ tau, Float* __restrict__ ssa, Float* __restrict__ g, @@ -109,6 +126,28 @@ namespace } } + __global__ + void combine_and_store_kernel_single_phase(const int ncol, const int nlay, const Float tmin, + Float* __restrict__ tau, Float* __restrict__ ssa, Float* __restrict__ g, + const Float* __restrict__ l_or_i_tau, const Float* __restrict__ l_or_i_taussa, const Float* __restrict__ l_or_i_taussag + ) + { + const int icol = blockIdx.x*blockDim.x + threadIdx.x; + const int ilay = blockIdx.y*blockDim.y + threadIdx.y; + + if ( (icol < ncol) && (ilay < nlay) ) + { + const int idx = icol + ilay*ncol; + const Float tau_t = l_or_i_tau[idx]; + const Float taussa = l_or_i_taussa[idx]; + const Float taussag = l_or_i_taussag[idx]; + + tau[idx] = tau_t; + ssa[idx] = taussa / max(tau_t, tmin); + g[idx] = taussag/ max(taussa, tmin); + } + } + __global__ void set_mask(const int ncol, const int nlay, const Float min_value, Bool* __restrict__ mask, const Float* __restrict__ values) @@ -200,26 +239,34 @@ void Cloud_optics_rt::cloud_optics( dim3 grid_m_gpu(grid_col_m, grid_lay_m); dim3 block_m_gpu(block_col_m, block_lay_m); - Array_gpu liqmsk({ncol, nlay}); + Array_gpu liqmsk, icemsk; + if (clwp.ptr() != nullptr){ + Array_gpu liqmsk({ncol, nlay}); set_mask<<>>( ncol, nlay, mask_min_value, liqmsk.ptr(), clwp.ptr()); } - Array_gpu icemsk({ncol, nlay}); + if (ciwp.ptr() != nullptr){ + Array_gpu icemsk({ncol, nlay}); set_mask<<>>( ncol, nlay, mask_min_value, icemsk.ptr(), ciwp.ptr()); } // Temporary arrays for storage. - Array_gpu ltau ({ncol, nlay}); - Array_gpu ltaussa ({ncol, nlay}); - Array_gpu ltaussag({ncol, nlay}); - - Array_gpu itau ({ncol, nlay}); - Array_gpu itaussa ({ncol, nlay}); - Array_gpu itaussag({ncol, nlay}); + Array_gpu ltau, ltaussa, ltaussag; + Array_gpu itau, itaussa, itaussag; + if (clwp.ptr() != nullptr){ + Array_gpu ltau ({ncol, nlay}); + Array_gpu ltaussa ({ncol, nlay}); + Array_gpu ltaussag({ncol, nlay}); + } + if (ciwp.ptr() != nullptr){ + Array_gpu itau ({ncol, nlay}); + Array_gpu itaussa ({ncol, nlay}); + Array_gpu itaussag({ncol, nlay}); + } const int block_col = 64; const int block_lay = 1; @@ -248,12 +295,27 @@ void Cloud_optics_rt::cloud_optics( this->lut_asyice_gpu.ptr(), itau.ptr(), itaussa.ptr(), itaussag.ptr()); } constexpr Float eps = std::numeric_limits::epsilon(); - - combine_and_store_kernel<<>>( + if ((ciwp.ptr() != nullptr) && (clwp.ptr() != nullptr)) + { + combine_and_store_kernel<<>>( ncol, nlay, eps, optical_props.get_tau().ptr(), optical_props.get_ssa().ptr(), optical_props.get_g().ptr(), ltau.ptr(), ltaussa.ptr(), ltaussag.ptr(), itau.ptr(), itaussa.ptr(), itaussag.ptr()); + } else if(ciwp.ptr() == nullptr) + { + combine_and_store_kernel_single_phase<<>>( + ncol, nlay, eps, + optical_props.get_tau().ptr(), optical_props.get_ssa().ptr(), optical_props.get_g().ptr(), + ltau.ptr(), ltaussa.ptr(), ltaussag.ptr()); + } else if (clwp.ptr() == nullptr) + { + combine_and_store_kernel_single_phase<<>>( + ncol, nlay, eps, + optical_props.get_tau().ptr(), optical_props.get_ssa().ptr(), optical_props.get_g().ptr(), + itau.ptr(), itaussa.ptr(), itaussag.ptr()); + } + } // 1scl variant of cloud optics. @@ -280,24 +342,29 @@ void Cloud_optics_rt::cloud_optics( dim3 grid_m_gpu(grid_col_m, grid_lay_m); dim3 block_m_gpu(block_col_m, block_lay_m); - Array_gpu liqmsk({ncol, nlay}); - if (clwp.ptr() != nullptr){ + // Temporary arrays for storage. + Array_gpu liqmsk, icemsk; + Array_gpu ltau, ltaussa, ltaussag; + Array_gpu itau, itaussa, itaussag; + if (clwp.ptr() != nullptr) + { + Array_gpu liqmsk({ncol, nlay}); set_mask<<>>( - ncol, nlay, mask_min_value, liqmsk.ptr(), clwp.ptr()); + ncol, nlay, mask_min_value, liqmsk.ptr(), clwp.ptr()); + + Array_gpu ltau ({ncol, nlay}); + Array_gpu ltaussa ({ncol, nlay}); + Array_gpu ltaussag({ncol, nlay}); } - Array_gpu icemsk({ncol, nlay}); - if (ciwp.ptr() != nullptr){ + if (ciwp.ptr() != nullptr) + { + Array_gpu icemsk({ncol, nlay}); set_mask<<>>( ncol, nlay, mask_min_value, icemsk.ptr(), ciwp.ptr()); + Array_gpu itau ({ncol, nlay}); + Array_gpu itaussa ({ncol, nlay}); + Array_gpu itaussag({ncol, nlay}); } - // Temporary arrays for storage. - Array_gpu ltau ({ncol, nlay}); - Array_gpu ltaussa ({ncol, nlay}); - Array_gpu ltaussag({ncol, nlay}); - - Array_gpu itau ({ncol, nlay}); - Array_gpu itaussa ({ncol, nlay}); - Array_gpu itaussag({ncol, nlay}); const int block_col = 64; const int block_lay = 1; @@ -327,11 +394,27 @@ void Cloud_optics_rt::cloud_optics( } constexpr Float eps = std::numeric_limits::epsilon(); - - combine_and_store_kernel<<>>( + if ((ciwp.ptr() != nullptr) && (clwp.ptr() != nullptr)) + { + combine_and_store_kernel<<>>( ncol, nlay, eps, optical_props.get_tau().ptr(), ltau.ptr(), ltaussa.ptr(), itau.ptr(), itaussa.ptr()); + } else if(ciwp.ptr() == nullptr) + { + combine_and_store_kernel_single_phase<<>>( + ncol, nlay, eps, + optical_props.get_tau().ptr(), + ltau.ptr(), ltaussa.ptr()); + } else if(clwp.ptr() == nullptr) + { + combine_and_store_kernel_single_phase<<>>( + ncol, nlay, eps, + optical_props.get_tau().ptr(), + itau.ptr(), itaussa.ptr()); + + } + } From f05983fefabc51b5a1dcac800b3ed25c787be9e3 Mon Sep 17 00:00:00 2001 From: magpowell Date: Wed, 6 Nov 2024 12:11:16 -0800 Subject: [PATCH 03/11] clean up cloud optics switches --- src_test/test_rte_rrtmgp_rt.cu | 11 ++++++++--- 1 file changed, 8 insertions(+), 3 deletions(-) diff --git a/src_test/test_rte_rrtmgp_rt.cu b/src_test/test_rte_rrtmgp_rt.cu index 2e93cf9e..b3931fe5 100644 --- a/src_test/test_rte_rrtmgp_rt.cu +++ b/src_test/test_rte_rrtmgp_rt.cu @@ -275,9 +275,14 @@ void solve_radiation(int argc, char** argv) throw std::runtime_error(error); } - if (switch_liq_cloud_optics && switch_ice_cloud_optics) { - std::string error = "Both liquid-only and ice-only cloud optics cannot be enabled simultaneously"; - throw std::runtime_error(error); + if (switch_cloud_optics) + { + switch_liq_cloud_optics = true; + switch_ice_cloud_optics = true; + } + if (switch_liq_cloud_optics || switch_ice_cloud_optics) + { + switch_cloud_optics = true; } // Print the options to the screen. From 67d548769bfeb4c3a2e1b48feba98f6ae043d982 Mon Sep 17 00:00:00 2001 From: magpowell Date: Wed, 6 Nov 2024 12:15:30 -0800 Subject: [PATCH 04/11] fix switch logic --- src_test/test_rte_rrtmgp_rt.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src_test/test_rte_rrtmgp_rt.cu b/src_test/test_rte_rrtmgp_rt.cu index b3931fe5..15407228 100644 --- a/src_test/test_rte_rrtmgp_rt.cu +++ b/src_test/test_rte_rrtmgp_rt.cu @@ -368,7 +368,7 @@ void solve_radiation(int argc, char** argv) if (switch_cloud_optics) { - if(!switch_ice_cloud_optics){ + if(switch_liq_cloud_optics){ lwp.set_dims({n_col, n_lay}); lwp = std::move(input_nc.get_variable("lwp", {n_lay, n_col_y, n_col_x})); @@ -376,7 +376,7 @@ void solve_radiation(int argc, char** argv) rel = std::move(input_nc.get_variable("rel", {n_lay, n_col_y, n_col_x})); } - if(!switch_liq_cloud_optics){ + if(switch_ice_cloud_optics){ iwp.set_dims({n_col, n_lay}); iwp = std::move(input_nc.get_variable("iwp", {n_lay, n_col_y, n_col_x})); From 9130adecfe27270a995f845a814c89494af79cf5 Mon Sep 17 00:00:00 2001 From: magpowell Date: Wed, 6 Nov 2024 12:17:45 -0800 Subject: [PATCH 05/11] clarify command line options --- src_test/test_rte_rrtmgp_rt.cu | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/src_test/test_rte_rrtmgp_rt.cu b/src_test/test_rte_rrtmgp_rt.cu index 15407228..47c28b3b 100644 --- a/src_test/test_rte_rrtmgp_rt.cu +++ b/src_test/test_rte_rrtmgp_rt.cu @@ -229,9 +229,9 @@ void solve_radiation(int argc, char** argv) {"fluxes" , { true, "Enable computation of fluxes." }}, {"raytracing" , { true, "Use raytracing for flux computation. '--raytracing 256': use 256 rays per pixel" }}, {"independent-column", { false, "run raytracer in independent column mode"}}, - {"cloud-optics" , { false, "Enable cloud optics." }}, - {"liq-cloud-optics" , { false, "liquid only cloud optics." }}, - {"ice-cloud-optics" , { false, "ice only cloud optics." }}, + {"cloud-optics" , { false, "Enable cloud optics (both liquid and ice)."}}, + {"liq-cloud-optics" , { false, "liquid only cloud optics." }}, + {"ice-cloud-optics" , { false, "ice only cloud optics." }}, {"cloud-mie" , { false, "Use Mie tables for cloud scattering in ray tracer" }}, {"aerosol-optics" , { false, "Enable aerosol optics." }}, {"single-gpt" , { false, "Output optical properties and fluxes for a single g-point. '--single-gpt 100': output 100th g-point" }}, From 0d5f8bdfc529fdadc0b4e65ce8261ed09f63d97f Mon Sep 17 00:00:00 2001 From: magpowell Date: Wed, 6 Nov 2024 12:35:21 -0800 Subject: [PATCH 06/11] non constant bools --- src_test/test_rte_rrtmgp_rt.cu | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/src_test/test_rte_rrtmgp_rt.cu b/src_test/test_rte_rrtmgp_rt.cu index 47c28b3b..2378afb6 100644 --- a/src_test/test_rte_rrtmgp_rt.cu +++ b/src_test/test_rte_rrtmgp_rt.cu @@ -251,9 +251,9 @@ void solve_radiation(int argc, char** argv) const bool switch_fluxes = command_line_switches.at("fluxes" ).first; const bool switch_raytracing = command_line_switches.at("raytracing" ).first; const bool switch_independent_column= command_line_switches.at("independent-column").first; - const bool switch_cloud_optics = command_line_switches.at("cloud-optics" ).first; - const bool switch_liq_cloud_optics = command_line_switches.at("liq-cloud-optics" ).first; - const bool switch_ice_cloud_optics = command_line_switches.at("ice-cloud-optics" ).first; + bool switch_cloud_optics = command_line_switches.at("cloud-optics" ).first; + bool switch_liq_cloud_optics = command_line_switches.at("liq-cloud-optics" ).first; + bool switch_ice_cloud_optics = command_line_switches.at("ice-cloud-optics" ).first; const bool switch_cloud_mie = command_line_switches.at("cloud-mie" ).first; const bool switch_aerosol_optics = command_line_switches.at("aerosol-optics" ).first; const bool switch_single_gpt = command_line_switches.at("single-gpt" ).first; From a15331dbc1171cb8035cdcd18f739c9fff641d90 Mon Sep 17 00:00:00 2001 From: magpowell Date: Wed, 6 Nov 2024 13:38:23 -0800 Subject: [PATCH 07/11] disable 2stream --- include_test/Radiation_solver_rt.h | 1 + src_test/Radiation_solver_rt.cu | 22 ++++--- src_test/test_rte_rrtmgp_rt.cu | 102 +++++++++++++++++------------ 3 files changed, 75 insertions(+), 50 deletions(-) diff --git a/include_test/Radiation_solver_rt.h b/include_test/Radiation_solver_rt.h index bb56d27f..4d59a976 100644 --- a/include_test/Radiation_solver_rt.h +++ b/include_test/Radiation_solver_rt.h @@ -99,6 +99,7 @@ class Radiation_solver_shortwave #ifdef __CUDACC__ void solve_gpu( const bool switch_fluxes, + const bool switch_disable_2s, const bool switch_raytracing, const bool switch_independent_column, const bool switch_cloud_optics, diff --git a/src_test/Radiation_solver_rt.cu b/src_test/Radiation_solver_rt.cu index 49348a22..883fe8eb 100644 --- a/src_test/Radiation_solver_rt.cu +++ b/src_test/Radiation_solver_rt.cu @@ -567,6 +567,7 @@ void Radiation_solver_shortwave::load_mie_tables( void Radiation_solver_shortwave::solve_gpu( const bool switch_fluxes, + const bool switch_disable_2s, const bool switch_raytracing, const bool switch_independent_column, const bool switch_cloud_optics, @@ -636,10 +637,13 @@ void Radiation_solver_shortwave::solve_gpu( if (switch_fluxes) { - Gas_optics_rrtmgp_kernels_cuda_rt::zero_array(n_lev, grid_cells.y, grid_cells.x, sw_flux_up.ptr()); - Gas_optics_rrtmgp_kernels_cuda_rt::zero_array(n_lev, grid_cells.y, grid_cells.x, sw_flux_dn.ptr()); - Gas_optics_rrtmgp_kernels_cuda_rt::zero_array(n_lev, grid_cells.y, grid_cells.x, sw_flux_dn_dir.ptr()); - Gas_optics_rrtmgp_kernels_cuda_rt::zero_array(n_lev, grid_cells.y, grid_cells.x, sw_flux_net.ptr()); + if (!switch_disable_2s) + { + Gas_optics_rrtmgp_kernels_cuda_rt::zero_array(n_lev, grid_cells.y, grid_cells.x, sw_flux_up.ptr()); + Gas_optics_rrtmgp_kernels_cuda_rt::zero_array(n_lev, grid_cells.y, grid_cells.x, sw_flux_dn.ptr()); + Gas_optics_rrtmgp_kernels_cuda_rt::zero_array(n_lev, grid_cells.y, grid_cells.x, sw_flux_dn_dir.ptr()); + Gas_optics_rrtmgp_kernels_cuda_rt::zero_array(n_lev, grid_cells.y, grid_cells.x, sw_flux_net.ptr()); + } if (switch_raytracing) { Gas_optics_rrtmgp_kernels_cuda_rt::zero_array(grid_cells.y, grid_cells.x, rt_flux_tod_up.ptr()); @@ -836,10 +840,12 @@ void Radiation_solver_shortwave::solve_gpu( } (*fluxes).net_flux(); - - Gpt_combine_kernels_cuda_rt::add_from_gpoint( - n_col, n_lev, sw_flux_up.ptr(), sw_flux_dn.ptr(), sw_flux_dn_dir.ptr(), sw_flux_net.ptr(), - (*fluxes).get_flux_up().ptr(), (*fluxes).get_flux_dn().ptr(), (*fluxes).get_flux_dn_dir().ptr(), (*fluxes).get_flux_net().ptr()); + if (!switch_disable_2s) + { + Gpt_combine_kernels_cuda_rt::add_from_gpoint( + n_col, n_lev, sw_flux_up.ptr(), sw_flux_dn.ptr(), sw_flux_dn_dir.ptr(), sw_flux_net.ptr(), + (*fluxes).get_flux_up().ptr(), (*fluxes).get_flux_dn().ptr(), (*fluxes).get_flux_dn_dir().ptr(), (*fluxes).get_flux_net().ptr()); + } if (switch_raytracing) { diff --git a/src_test/test_rte_rrtmgp_rt.cu b/src_test/test_rte_rrtmgp_rt.cu index 2378afb6..267c8dc3 100644 --- a/src_test/test_rte_rrtmgp_rt.cu +++ b/src_test/test_rte_rrtmgp_rt.cu @@ -227,6 +227,7 @@ void solve_radiation(int argc, char** argv) {"shortwave" , { true, "Enable computation of shortwave radiation."}}, {"longwave" , { false, "Enable computation of longwave radiation." }}, {"fluxes" , { true, "Enable computation of fluxes." }}, + {"disable-2s" , { false, "use raytracing onlu for flux computation. must be passed with raytracing" }}, {"raytracing" , { true, "Use raytracing for flux computation. '--raytracing 256': use 256 rays per pixel" }}, {"independent-column", { false, "run raytracer in independent column mode"}}, {"cloud-optics" , { false, "Enable cloud optics (both liquid and ice)."}}, @@ -249,6 +250,7 @@ void solve_radiation(int argc, char** argv) const bool switch_shortwave = command_line_switches.at("shortwave" ).first; const bool switch_longwave = command_line_switches.at("longwave" ).first; const bool switch_fluxes = command_line_switches.at("fluxes" ).first; + const bool switch_disable_2s = command_line_switches.at("disable-2s" ).first; const bool switch_raytracing = command_line_switches.at("raytracing" ).first; const bool switch_independent_column= command_line_switches.at("independent-column").first; bool switch_cloud_optics = command_line_switches.at("cloud-optics" ).first; @@ -275,6 +277,11 @@ void solve_radiation(int argc, char** argv) throw std::runtime_error(error); } + if (switch_disable_2s && !switch_raytracing) { + std::string error = "cannot disable two-stream for flux calculation without turning ray tracing on"; + throw std::runtime_error(error); + } + if (switch_cloud_optics) { switch_liq_cloud_optics = true; @@ -721,10 +728,13 @@ void solve_radiation(int argc, char** argv) if (switch_fluxes) { - sw_flux_up .set_dims({n_col, n_lev}); - sw_flux_dn .set_dims({n_col, n_lev}); - sw_flux_dn_dir.set_dims({n_col, n_lev}); - sw_flux_net .set_dims({n_col, n_lev}); + if(!switch_disable_2s) + { + sw_flux_up .set_dims({n_col, n_lev}); + sw_flux_dn .set_dims({n_col, n_lev}); + sw_flux_dn_dir.set_dims({n_col, n_lev}); + sw_flux_net .set_dims({n_col, n_lev}); + } if (switch_raytracing) { @@ -785,6 +795,7 @@ void solve_radiation(int argc, char** argv) rad_sw.solve_gpu( switch_fluxes, + switch_disable_2s, switch_raytracing, switch_independent_column, switch_cloud_optics, @@ -922,27 +933,31 @@ void solve_radiation(int argc, char** argv) if (switch_fluxes) { - auto nc_sw_flux_up = output_nc.add_variable("sw_flux_up" , {"lev", "y", "x"}); - auto nc_sw_flux_dn = output_nc.add_variable("sw_flux_dn" , {"lev", "y", "x"}); - auto nc_sw_flux_dn_dir = output_nc.add_variable("sw_flux_dn_dir", {"lev", "y", "x"}); - auto nc_sw_flux_net = output_nc.add_variable("sw_flux_net" , {"lev", "y", "x"}); + if (!switch_disable_2s) + { + auto nc_sw_flux_up = output_nc.add_variable("sw_flux_up" , {"lev", "y", "x"}); + auto nc_sw_flux_dn = output_nc.add_variable("sw_flux_dn" , {"lev", "y", "x"}); + auto nc_sw_flux_dn_dir = output_nc.add_variable("sw_flux_dn_dir", {"lev", "y", "x"}); + auto nc_sw_flux_net = output_nc.add_variable("sw_flux_net" , {"lev", "y", "x"}); + + nc_sw_flux_up .insert(sw_flux_up_cpu .v(), {0, 0, 0}); + nc_sw_flux_dn .insert(sw_flux_dn_cpu .v(), {0, 0, 0}); + nc_sw_flux_dn_dir.insert(sw_flux_dn_dir_cpu.v(), {0, 0, 0}); + nc_sw_flux_net .insert(sw_flux_net_cpu .v(), {0, 0, 0}); - nc_sw_flux_up .insert(sw_flux_up_cpu .v(), {0, 0, 0}); - nc_sw_flux_dn .insert(sw_flux_dn_cpu .v(), {0, 0, 0}); - nc_sw_flux_dn_dir.insert(sw_flux_dn_dir_cpu.v(), {0, 0, 0}); - nc_sw_flux_net .insert(sw_flux_net_cpu .v(), {0, 0, 0}); + nc_sw_flux_up.add_attribute("long_name","Upwelling shortwave fluxes (TwoStream solver)"); + nc_sw_flux_up.add_attribute("units","W m-2"); - nc_sw_flux_up.add_attribute("long_name","Upwelling shortwave fluxes (TwoStream solver)"); - nc_sw_flux_up.add_attribute("units","W m-2"); + nc_sw_flux_dn.add_attribute("long_name","Downwelling shortwave fluxes (TwoStream solver)"); + nc_sw_flux_dn.add_attribute("units","W m-2"); - nc_sw_flux_dn.add_attribute("long_name","Downwelling shortwave fluxes (TwoStream solver)"); - nc_sw_flux_dn.add_attribute("units","W m-2"); + nc_sw_flux_dn_dir.add_attribute("long_name","Downwelling direct shortwave fluxes (TwoStream solver)"); + nc_sw_flux_dn_dir.add_attribute("units","W m-2"); - nc_sw_flux_dn_dir.add_attribute("long_name","Downwelling direct shortwave fluxes (TwoStream solver)"); - nc_sw_flux_dn_dir.add_attribute("units","W m-2"); + nc_sw_flux_net.add_attribute("long_name","Net shortwave fluxes (TwoStream solver)"); + nc_sw_flux_net.add_attribute("units","W m-2"); - nc_sw_flux_net.add_attribute("long_name","Net shortwave fluxes (TwoStream solver)"); - nc_sw_flux_net.add_attribute("units","W m-2"); + } if (switch_raytracing) { @@ -983,28 +998,31 @@ void solve_radiation(int argc, char** argv) if (switch_single_gpt) - { - auto nc_sw_gpt_flux_up = output_nc.add_variable("sw_gpt_flux_up" , {"lev", "y", "x"}); - auto nc_sw_gpt_flux_dn = output_nc.add_variable("sw_gpt_flux_dn" , {"lev", "y", "x"}); - auto nc_sw_gpt_flux_dn_dir = output_nc.add_variable("sw_gpt_flux_dn_dir", {"lev", "y", "x"}); - auto nc_sw_gpt_flux_net = output_nc.add_variable("sw_gpt_flux_net" , {"lev", "y", "x"}); - - nc_sw_gpt_flux_up .insert(sw_gpt_flux_up_cpu .v(), {0, 0, 0}); - nc_sw_gpt_flux_dn .insert(sw_gpt_flux_dn_cpu .v(), {0, 0, 0}); - nc_sw_gpt_flux_dn_dir.insert(sw_gpt_flux_dn_dir_cpu.v(), {0, 0, 0}); - nc_sw_gpt_flux_net .insert(sw_gpt_flux_net_cpu .v(), {0, 0, 0}); - - nc_sw_gpt_flux_up.add_attribute("long_name","Upwelling shortwave fluxes for g-point "+std::to_string(single_gpt)+" (TwoStream solver)"); - nc_sw_gpt_flux_up.add_attribute("units","W m-2"); - - nc_sw_gpt_flux_dn.add_attribute("long_name","Downwelling shortwave fluxes for g-point "+std::to_string(single_gpt)+" (TwoStream solver)"); - nc_sw_gpt_flux_dn.add_attribute("units","W m-2"); - - nc_sw_gpt_flux_dn_dir.add_attribute("long_name","Downwelling direct shortwave fluxes for g-point "+std::to_string(single_gpt)+" (TwoStream solver)"); - nc_sw_gpt_flux_dn_dir.add_attribute("units","W m-2"); - - nc_sw_gpt_flux_net.add_attribute("long_name","Net shortwave fluxes for g-point "+std::to_string(single_gpt)+" (TwoStream solver)"); - nc_sw_gpt_flux_net.add_attribute("units","W m-2"); + { + if (!switch_disable_2s) + { + auto nc_sw_gpt_flux_up = output_nc.add_variable("sw_gpt_flux_up" , {"lev", "y", "x"}); + auto nc_sw_gpt_flux_dn = output_nc.add_variable("sw_gpt_flux_dn" , {"lev", "y", "x"}); + auto nc_sw_gpt_flux_dn_dir = output_nc.add_variable("sw_gpt_flux_dn_dir", {"lev", "y", "x"}); + auto nc_sw_gpt_flux_net = output_nc.add_variable("sw_gpt_flux_net" , {"lev", "y", "x"}); + + nc_sw_gpt_flux_up .insert(sw_gpt_flux_up_cpu .v(), {0, 0, 0}); + nc_sw_gpt_flux_dn .insert(sw_gpt_flux_dn_cpu .v(), {0, 0, 0}); + nc_sw_gpt_flux_dn_dir.insert(sw_gpt_flux_dn_dir_cpu.v(), {0, 0, 0}); + nc_sw_gpt_flux_net .insert(sw_gpt_flux_net_cpu .v(), {0, 0, 0}); + + nc_sw_gpt_flux_up.add_attribute("long_name","Upwelling shortwave fluxes for g-point "+std::to_string(single_gpt)+" (TwoStream solver)"); + nc_sw_gpt_flux_up.add_attribute("units","W m-2"); + + nc_sw_gpt_flux_dn.add_attribute("long_name","Downwelling shortwave fluxes for g-point "+std::to_string(single_gpt)+" (TwoStream solver)"); + nc_sw_gpt_flux_dn.add_attribute("units","W m-2"); + + nc_sw_gpt_flux_dn_dir.add_attribute("long_name","Downwelling direct shortwave fluxes for g-point "+std::to_string(single_gpt)+" (TwoStream solver)"); + nc_sw_gpt_flux_dn_dir.add_attribute("units","W m-2"); + + nc_sw_gpt_flux_net.add_attribute("long_name","Net shortwave fluxes for g-point "+std::to_string(single_gpt)+" (TwoStream solver)"); + nc_sw_gpt_flux_net.add_attribute("units","W m-2"); + } } } } From d85650ee4810fd17d5fc959fdc616a586b0b76b6 Mon Sep 17 00:00:00 2001 From: magpowell Date: Thu, 7 Nov 2024 13:42:47 -0800 Subject: [PATCH 08/11] ncol and nlay for ice or liq --- src_cuda_rt/Cloud_optics_rt.cu | 39 ++++++++++++++++++++++++---------- 1 file changed, 28 insertions(+), 11 deletions(-) diff --git a/src_cuda_rt/Cloud_optics_rt.cu b/src_cuda_rt/Cloud_optics_rt.cu index 0546a630..9e733e90 100644 --- a/src_cuda_rt/Cloud_optics_rt.cu +++ b/src_cuda_rt/Cloud_optics_rt.cu @@ -222,11 +222,19 @@ void Cloud_optics_rt::cloud_optics( const Array_gpu& reliq, const Array_gpu& reice, Optical_props_2str_rt& optical_props) { - const int ncol = clwp.dim(1); - const int nlay = clwp.dim(2); - - Optical_props_2str_rt clouds_liq(ncol, nlay, optical_props); - Optical_props_2str_rt clouds_ice(ncol, nlay, optical_props); + int ncol = -1; + int nlay = -1; + if (clwp.ptr() != nullptr) + { + const int ncol = clwp.dim(1); + const int nlay = clwp.dim(2); + Optical_props_1scl_rt clouds_liq(ncol, nlay, optical_props); + } else if (ciwp.ptr() != nullptr) + { + const int ncol = ciwp.dim(1); + const int nlay = ciwp.dim(2); + Optical_props_1scl_rt clouds_ice(ncol, nlay, optical_props); + } // Set the mask. constexpr Float mask_min_value = Float(0.); @@ -325,12 +333,21 @@ void Cloud_optics_rt::cloud_optics( const Array_gpu& reliq, const Array_gpu& reice, Optical_props_1scl_rt& optical_props) { - const int ncol = clwp.dim(1); - const int nlay = clwp.dim(2); - - Optical_props_1scl_rt clouds_liq(ncol, nlay, optical_props); - Optical_props_1scl_rt clouds_ice(ncol, nlay, optical_props); - + + int ncol = -1; + int nlay = -1; + if (clwp.ptr() != nullptr) + { + const int ncol = clwp.dim(1); + const int nlay = clwp.dim(2); + Optical_props_1scl_rt clouds_liq(ncol, nlay, optical_props); + } else if (ciwp.ptr() != nullptr) + { + const int ncol = ciwp.dim(1); + const int nlay = ciwp.dim(2); + Optical_props_1scl_rt clouds_ice(ncol, nlay, optical_props); + } + // Set the mask. constexpr Float mask_min_value = Float(0.); const int block_col_m = 16; From 5d43d519791156fe184d130a4692d4ac516fe310 Mon Sep 17 00:00:00 2001 From: magpowell Date: Fri, 8 Nov 2024 07:47:37 -0800 Subject: [PATCH 09/11] fix clr sky bug --- src_cuda_rt/Cloud_optics_rt.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src_cuda_rt/Cloud_optics_rt.cu b/src_cuda_rt/Cloud_optics_rt.cu index 9e733e90..06858ec0 100644 --- a/src_cuda_rt/Cloud_optics_rt.cu +++ b/src_cuda_rt/Cloud_optics_rt.cu @@ -228,12 +228,12 @@ void Cloud_optics_rt::cloud_optics( { const int ncol = clwp.dim(1); const int nlay = clwp.dim(2); - Optical_props_1scl_rt clouds_liq(ncol, nlay, optical_props); + Optical_props_2str_rt clouds_liq(ncol, nlay, optical_props); } else if (ciwp.ptr() != nullptr) { const int ncol = ciwp.dim(1); const int nlay = ciwp.dim(2); - Optical_props_1scl_rt clouds_ice(ncol, nlay, optical_props); + Optical_props_2str_rt clouds_ice(ncol, nlay, optical_props); } // Set the mask. From c1ae70558fd0ff2325de440d30297e81e6fc83fd Mon Sep 17 00:00:00 2001 From: magpowell Date: Fri, 8 Nov 2024 11:19:13 -0800 Subject: [PATCH 10/11] fix scope issues --- src_cuda_rt/Cloud_optics_rt.cu | 88 ++++++++++++++++++---------------- 1 file changed, 46 insertions(+), 42 deletions(-) diff --git a/src_cuda_rt/Cloud_optics_rt.cu b/src_cuda_rt/Cloud_optics_rt.cu index 06858ec0..5bcceddc 100644 --- a/src_cuda_rt/Cloud_optics_rt.cu +++ b/src_cuda_rt/Cloud_optics_rt.cu @@ -226,13 +226,13 @@ void Cloud_optics_rt::cloud_optics( int nlay = -1; if (clwp.ptr() != nullptr) { - const int ncol = clwp.dim(1); - const int nlay = clwp.dim(2); + ncol = clwp.dim(1); + nlay = clwp.dim(2); Optical_props_2str_rt clouds_liq(ncol, nlay, optical_props); } else if (ciwp.ptr() != nullptr) { - const int ncol = ciwp.dim(1); - const int nlay = ciwp.dim(2); + ncol = ciwp.dim(1); + nlay = ciwp.dim(2); Optical_props_2str_rt clouds_ice(ncol, nlay, optical_props); } @@ -247,38 +247,36 @@ void Cloud_optics_rt::cloud_optics( dim3 grid_m_gpu(grid_col_m, grid_lay_m); dim3 block_m_gpu(block_col_m, block_lay_m); - Array_gpu liqmsk, icemsk; + // Temporary arrays for storage. + Array_gpu liqmsk({0, 0}); + Array_gpu ltau ({0, 0}); + Array_gpu ltaussa ({0, 0}); + Array_gpu ltaussag({0, 0}); + Array_gpu icemsk({0, 0}); + Array_gpu itau ({0, 0}); + Array_gpu itaussa ({0, 0}); + Array_gpu itaussag({0, 0}); if (clwp.ptr() != nullptr){ - Array_gpu liqmsk({ncol, nlay}); + liqmsk.set_dims({ncol, nlay}); + ltau.set_dims({ncol, nlay}); + ltaussa.set_dims({ncol, nlay}); + ltaussag.set_dims({ncol, nlay}); + set_mask<<>>( ncol, nlay, mask_min_value, liqmsk.ptr(), clwp.ptr()); } - if (ciwp.ptr() != nullptr){ - Array_gpu icemsk({ncol, nlay}); + icemsk.set_dims({ncol, nlay}); + itau.set_dims({ncol, nlay}); + itaussa.set_dims({ncol, nlay}); + itaussag.set_dims({ncol, nlay}); + set_mask<<>>( ncol, nlay, mask_min_value, icemsk.ptr(), ciwp.ptr()); } - - - // Temporary arrays for storage. - Array_gpu ltau, ltaussa, ltaussag; - Array_gpu itau, itaussa, itaussag; - if (clwp.ptr() != nullptr){ - Array_gpu ltau ({ncol, nlay}); - Array_gpu ltaussa ({ncol, nlay}); - Array_gpu ltaussag({ncol, nlay}); - } - if (ciwp.ptr() != nullptr){ - Array_gpu itau ({ncol, nlay}); - Array_gpu itaussa ({ncol, nlay}); - Array_gpu itaussag({ncol, nlay}); - } - const int block_col = 64; const int block_lay = 1; - const int grid_col = ncol/block_col + (ncol%block_col > 0); const int grid_lay = nlay/block_lay + (nlay%block_lay > 0); @@ -333,18 +331,17 @@ void Cloud_optics_rt::cloud_optics( const Array_gpu& reliq, const Array_gpu& reice, Optical_props_1scl_rt& optical_props) { - int ncol = -1; int nlay = -1; if (clwp.ptr() != nullptr) { - const int ncol = clwp.dim(1); - const int nlay = clwp.dim(2); + ncol = clwp.dim(1); + nlay = clwp.dim(2); Optical_props_1scl_rt clouds_liq(ncol, nlay, optical_props); } else if (ciwp.ptr() != nullptr) { - const int ncol = ciwp.dim(1); - const int nlay = ciwp.dim(2); + ncol = ciwp.dim(1); + nlay = ciwp.dim(2); Optical_props_1scl_rt clouds_ice(ncol, nlay, optical_props); } @@ -360,27 +357,34 @@ void Cloud_optics_rt::cloud_optics( dim3 block_m_gpu(block_col_m, block_lay_m); // Temporary arrays for storage. - Array_gpu liqmsk, icemsk; - Array_gpu ltau, ltaussa, ltaussag; - Array_gpu itau, itaussa, itaussag; + Array_gpu liqmsk({0, 0}); + Array_gpu ltau ({0, 0}); + Array_gpu ltaussa ({0, 0}); + Array_gpu ltaussag({0, 0}); + Array_gpu icemsk({0, 0}); + Array_gpu itau ({0, 0}); + Array_gpu itaussa ({0, 0}); + Array_gpu itaussag({0, 0}); + if (clwp.ptr() != nullptr) { - Array_gpu liqmsk({ncol, nlay}); + liqmsk.set_dims({ncol, nlay}); + ltau.set_dims({ncol, nlay}); + ltaussa.set_dims({ncol, nlay}); + ltaussag.set_dims({ncol, nlay}); + set_mask<<>>( ncol, nlay, mask_min_value, liqmsk.ptr(), clwp.ptr()); - - Array_gpu ltau ({ncol, nlay}); - Array_gpu ltaussa ({ncol, nlay}); - Array_gpu ltaussag({ncol, nlay}); } if (ciwp.ptr() != nullptr) { - Array_gpu icemsk({ncol, nlay}); + icemsk.set_dims({ncol, nlay}); + itau.set_dims({ncol, nlay}); + itaussa.set_dims({ncol, nlay}); + itaussag.set_dims({ncol, nlay}); + set_mask<<>>( ncol, nlay, mask_min_value, icemsk.ptr(), ciwp.ptr()); - Array_gpu itau ({ncol, nlay}); - Array_gpu itaussa ({ncol, nlay}); - Array_gpu itaussag({ncol, nlay}); } const int block_col = 64; From dd1e611231b5f3dbd311d46aa540d0986696264f Mon Sep 17 00:00:00 2001 From: magpowell Date: Tue, 19 Nov 2024 08:59:53 -0800 Subject: [PATCH 11/11] add override sza and azi option --- src_test/test_rte_rrtmgp_rt.cu | 43 ++++++++++++++++++++++++++++++---- 1 file changed, 38 insertions(+), 5 deletions(-) diff --git a/src_test/test_rte_rrtmgp_rt.cu b/src_test/test_rte_rrtmgp_rt.cu index 267c8dc3..20d43f75 100644 --- a/src_test/test_rte_rrtmgp_rt.cu +++ b/src_test/test_rte_rrtmgp_rt.cu @@ -238,11 +238,15 @@ void solve_radiation(int argc, char** argv) {"single-gpt" , { false, "Output optical properties and fluxes for a single g-point. '--single-gpt 100': output 100th g-point" }}, {"profiling" , { false, "Perform additional profiling run." }}, {"delta-cloud" , { false, "delta-scaling of cloud optical properties" }}, - {"delta-aerosol" , { false, "delta-scaling of aerosol optical properties" }} }; + {"delta-aerosol" , { false, "delta-scaling of aerosol optical properties" }} , + {"override-sza" , { false, "override provided value of sza in input file. IN DEGREES. '--override-sza 50': use a sza of 50 degrees" }}, + {"override-azi" , { false, "override provided value of azi in input file. IN DEGREES. '--override-azi 240': use of azi of 240 degrees" }}}; std::map> command_line_ints { {"raytracing", {32, "Number of rays initialised at TOD per pixel per quadraute."}}, - {"single-gpt", {1 , "g-point to store optical properties and fluxes of" }} }; + {"single-gpt", {1 , "g-point to store optical properties and fluxes of" }}, + {"override-sza", {0, "solar zenith angle (theta) in degrees."}}, + {"override-azi", {0, "Solar azimuth angle in degrees."}} }; if (parse_command_line_options(command_line_switches, command_line_ints, argc, argv)) return; @@ -262,7 +266,8 @@ void solve_radiation(int argc, char** argv) const bool switch_profiling = command_line_switches.at("profiling" ).first; const bool switch_delta_cloud = command_line_switches.at("delta-cloud" ).first; const bool switch_delta_aerosol = command_line_switches.at("delta-aerosol" ).first; - + const bool override_sza = command_line_switches.at("override-sza" ).first; + const bool override_azi = command_line_switches.at("override-azi" ).first; Int photons_per_pixel = Int(command_line_ints.at("raytracing").first); if (Float(int(std::log2(Float(photons_per_pixel)))) != std::log2(Float(photons_per_pixel))) @@ -296,9 +301,21 @@ void solve_radiation(int argc, char** argv) print_command_line_options(command_line_switches, command_line_ints); int single_gpt = command_line_ints.at("single-gpt").first; + int sza_deg = Int(command_line_ints.at("override-sza").first); + int azi_deg = Int(command_line_ints.at("override-azi").first); Status::print_message("Using "+ std::to_string(photons_per_pixel) + " rays per pixel"); + if (override_sza) + { + Status::print_message("Using SZA of "+ std::to_string(sza_deg) + " degrees"); + } + + if (override_azi) + { + Status::print_message("Using azi of "+ std::to_string(azi_deg) + " degrees"); + } + ////// READ THE ATMOSPHERIC DATA ////// Status::print_message("Reading atmospheric input data from NetCDF."); @@ -665,9 +682,25 @@ void solve_radiation(int argc, char** argv) rad_sw.load_mie_tables("mie_lut_broadband.nc"); } + Array mu0({n_col}); + Array azi({n_col}); + + if (override_sza) { + Float mu0_in = cosf(sza_deg * 3.14159f / 180.0f); + for (int icol=1; icol<=n_col; ++icol) + mu0({icol}) = mu0_in; + } else { + mu0 = input_nc.get_variable("mu0", {n_col_y, n_col_x}); + } + + if (override_azi) { + Float azi_in = azi_deg * 3.14159f / 180.0f; + for (int icol=1; icol<=n_col; ++icol) + azi({icol}) = azi_in; + } else { + azi = input_nc.get_variable("azi", {n_col_y, n_col_x}); + } - Array mu0(input_nc.get_variable("mu0", {n_col_y, n_col_x}), {n_col}); - Array azi(input_nc.get_variable("azi", {n_col_y, n_col_x}), {n_col}); Array sfc_alb_dir(input_nc.get_variable("sfc_alb_dir", {n_col_y, n_col_x, n_bnd_sw}), {n_bnd_sw, n_col}); Array sfc_alb_dif(input_nc.get_variable("sfc_alb_dif", {n_col_y, n_col_x, n_bnd_sw}), {n_bnd_sw, n_col});