Skip to content

Commit 3589b64

Browse files
committed
switch scaling so two stream gets it too
1 parent 56ff4dd commit 3589b64

File tree

4 files changed

+36
-43
lines changed

4 files changed

+36
-43
lines changed

include_rt/Raytracer.h

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -24,8 +24,6 @@ class Raytracer
2424
void trace_rays(
2525
const int qrng_gpt_offset,
2626
const bool switch_independent_column,
27-
const bool switch_attenuate_path,
28-
const Float attenuate_scale_factor,
2927
const Int photons_per_pixel,
3028
const Raytracer_definitions::Vector<int> grid_cells,
3129
const Raytracer_definitions::Vector<Float> grid_d,

src_cuda_rt/Raytracer.cu

Lines changed: 0 additions & 35 deletions
Original file line numberDiff line numberDiff line change
@@ -164,35 +164,6 @@ namespace
164164
}
165165
}
166166

167-
168-
__global__
169-
void scale_kext_kernel(const Vector<int> grid_cells, Float* __restrict__ k_ext, Float scale_factor) {
170-
const int icol_x = blockIdx.x*blockDim.x + threadIdx.x;
171-
const int icol_y = blockIdx.y*blockDim.y + threadIdx.y;
172-
const int iz = blockIdx.z*blockDim.z + threadIdx.z;
173-
174-
if ( (icol_x < grid_cells.x) && (icol_y < grid_cells.y) && (iz < (grid_cells.z)) )
175-
{
176-
const int idx = icol_x + icol_y*grid_cells.x + iz*grid_cells.y*grid_cells.x;
177-
k_ext[idx] = k_ext[idx] * scale_factor;
178-
}
179-
}
180-
181-
__global__
182-
void scale_k_sca_kernel(const Vector<int> grid_cells, Optics_scat* __restrict__ scat_asy, Float scale_factor) {
183-
const int icol_x = blockIdx.x*blockDim.x + threadIdx.x;
184-
const int icol_y = blockIdx.y*blockDim.y + threadIdx.y;
185-
const int iz = blockIdx.z*blockDim.z + threadIdx.z;
186-
187-
if ( (icol_x < grid_cells.x) && (icol_y < grid_cells.y) && (iz < (grid_cells.z)) )
188-
{
189-
const int idx = icol_x + icol_y*grid_cells.x + iz*grid_cells.y*grid_cells.x;
190-
scat_asy[idx].k_sca_gas *= scale_factor;
191-
scat_asy[idx].k_sca_cld *= scale_factor;
192-
scat_asy[idx].k_sca_aer *= scale_factor;
193-
}
194-
}
195-
196167
__global__
197168
void count_to_flux_2d(
198169
const Vector<int> grid_cells, const Float photons_per_col, const Float toa_src,
@@ -259,8 +230,6 @@ Raytracer::Raytracer()
259230
void Raytracer::trace_rays(
260231
const int igpt,
261232
const bool switch_independent_column,
262-
const bool switch_attenuate_path,
263-
const Float attenuate_scale_factor,
264233
const Int photons_per_pixel,
265234
const Vector<int> grid_cells,
266235
const Vector<Float> grid_d,
@@ -325,10 +294,6 @@ void Raytracer::trace_rays(
325294
tau_aeros.ptr(), ssa_aeros.ptr(), asy_aeros.ptr(),
326295
k_ext.ptr(), scat_asy.ptr());
327296

328-
if (switch_attenuate_path){
329-
scale_kext_kernel<<<grid_3d, block_3d>>>(grid_cells, k_ext.ptr(), attenuate_scale_factor);
330-
scale_k_sca_kernel<<<grid_3d, block_3d>>>(grid_cells, scat_asy.ptr(), attenuate_scale_factor);
331-
}
332297

333298
// create k_null_grid
334299
const int block_kn_x = 8;

src_test/Radiation_solver_rt.cu

Lines changed: 36 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -42,6 +42,29 @@
4242

4343
namespace
4444
{
45+
__global__
46+
void scale_tau_kernel(Float* tau, const int ncol, const int nlay, Float scale_factor) {
47+
const int icol = blockIdx.x*blockDim.x + threadIdx.x;
48+
const int ilay = blockIdx.y*blockDim.y + threadIdx.y;
49+
50+
if ( (icol < ncol) && (ilay < nlay) )
51+
{
52+
const int idx = icol + ilay*ncol;
53+
tau[idx] = tau[idx] * scale_factor;
54+
}
55+
}
56+
57+
void scale_tau(Float* tau, const int ncol, const int nlay, Float scale_factor) {
58+
const int block_col = 64;
59+
const int block_lay = 1;
60+
const int grid_col = ncol/block_col + (ncol%block_col > 0);
61+
const int grid_lay = nlay/block_lay + (nlay%block_lay > 0);
62+
63+
dim3 grid_gpu(grid_col, grid_lay);
64+
dim3 block_gpu(block_col, block_lay);
65+
scale_tau_kernel<<<grid_gpu, block_gpu>>>(tau, ncol, nlay, scale_factor);
66+
}
67+
4568
std::vector<std::string> get_variable_string(
4669
const std::string& var_name,
4770
std::vector<int> i_count,
@@ -714,6 +737,11 @@ void Radiation_solver_shortwave::solve_gpu(
714737

715738
toa_src.fill(toa_src_temp({1}) * tsi_scaling({1}));
716739

740+
if (switch_attenuate_tica)
741+
{
742+
scale_tau(dynamic_cast<Optical_props_2str_rt&>(*optical_props).get_tau().ptr(), n_col, n_lay, attenuate_scale_factor);
743+
}
744+
717745
if (switch_aerosol_optics)
718746
{
719747
if (band > previous_band)
@@ -726,6 +754,10 @@ void Radiation_solver_shortwave::solve_gpu(
726754
*aerosol_optical_props);
727755
if (switch_delta_aerosol)
728756
aerosol_optical_props->delta_scale();
757+
if (switch_attenuate_tica)
758+
{
759+
scale_tau(dynamic_cast<Optical_props_2str_rt&>(*aerosol_optical_props).get_tau().ptr(), n_col, n_lay, attenuate_scale_factor);
760+
}
729761
}
730762

731763
// Add the cloud optical props to the gas optical properties.
@@ -754,6 +786,10 @@ void Radiation_solver_shortwave::solve_gpu(
754786

755787
if (switch_delta_cloud)
756788
cloud_optical_props->delta_scale();
789+
if (switch_attenuate_tica)
790+
{
791+
scale_tau(dynamic_cast<Optical_props_2str_rt&>(*cloud_optical_props).get_tau().ptr(), n_col, n_lay, attenuate_scale_factor);
792+
}
757793
}
758794
// Add the cloud optical props to the gas optical properties.
759795
add_to(
@@ -815,8 +851,6 @@ void Radiation_solver_shortwave::solve_gpu(
815851
raytracer.trace_rays(
816852
igpt,
817853
switch_independent_column,
818-
switch_attenuate_tica,
819-
attenuate_scale_factor,
820854
ray_count,
821855
grid_cells,
822856
grid_d,

src_test/test_rt_lite.cu

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -411,13 +411,9 @@ void solve_radiation(int argc, char** argv)
411411

412412
cudaEventRecord(start, 0);
413413
// do something.
414-
const bool switch_attenuate_path = false;
415-
const Float attenuate_scale_factor = 0;
416414
raytracer.trace_rays(
417415
0,
418416
switch_independent_column,
419-
switch_attenuate_path,
420-
attenuate_scale_factor,
421417
photons_per_pixel,
422418
grid_cells,
423419
grid_d,

0 commit comments

Comments
 (0)