From 1b9c42d219aa7eff0624b3c142807366568d76d4 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ole=20Sch=C3=BCtt?= Date: Mon, 30 Aug 2021 20:59:43 +0200 Subject: [PATCH] precommit: Add file extension .cc --- src/grid/hip/grid_hip_collocate.cc | 62 ++++----- src/grid/hip/grid_hip_context.cc | 75 +++++------ src/grid/hip/grid_hip_integrate.cc | 154 ++++++++++++----------- tools/precommit/check_file_properties.py | 2 +- tools/precommit/precommit.py | 2 +- 5 files changed, 154 insertions(+), 141 deletions(-) diff --git a/src/grid/hip/grid_hip_collocate.cc b/src/grid/hip/grid_hip_collocate.cc index 105c934c13..a4d25e2013 100644 --- a/src/grid/hip/grid_hip_collocate.cc +++ b/src/grid/hip/grid_hip_collocate.cc @@ -135,8 +135,9 @@ __global__ void calculate_coefficients(const kernel_params dev_) { calculate_coefficients. We only keep the non zero elements to same memory. */ - template -__global__ __launch_bounds__(64) void collocate_kernel(const kernel_params dev_) { +template +__global__ +__launch_bounds__(64) void collocate_kernel(const kernel_params dev_) { // Copy task from global to shared memory and precompute some stuff. __shared__ smem_task_reduced task; @@ -180,11 +181,11 @@ __global__ __launch_bounds__(64) void collocate_kernel(const kernel_params dev_) if (distributed__) { if (task.apply_border_mask) { - compute_window_size(dev_.grid_local_size_, dev_.grid_lower_corner_, - dev_.grid_full_size_, // also full size of the grid - dev_.tasks[dev_.first_task + blockIdx.x].border_mask, - dev_.grid_border_width_, &task.window_size, - &task.window_shift); + compute_window_size( + dev_.grid_local_size_, dev_.grid_lower_corner_, + dev_.grid_full_size_, // also full size of the grid + dev_.tasks[dev_.first_task + blockIdx.x].border_mask, + dev_.grid_border_width_, &task.window_size, &task.window_shift); } } } @@ -200,10 +201,10 @@ __global__ __launch_bounds__(64) void collocate_kernel(const kernel_params dev_) // check if the point is within the window if (task.apply_border_mask) { // this test is only relevant when the grid is split over several mpi - // ranks. in that case we take only the points contributing to local part - // of the grid. + // ranks. in that case we take only the points contributing to local + // part of the grid. if ((z2 < task.window_shift.z) || (z2 > task.window_size.z)) { - continue; + continue; } } } @@ -274,16 +275,15 @@ __global__ __launch_bounds__(64) void collocate_kernel(const kernel_params dev_) r3.y = (y + task.lb_cube.y + task.roffset.y) * dh_[4]; r3.z = (z + task.lb_cube.z + task.roffset.z) * dh_[8]; } else { - r3 = - compute_coordinates(dh_, (x + task.lb_cube.x + task.roffset.x), - (y + task.lb_cube.y + task.roffset.y), - (z + task.lb_cube.z + task.roffset.z)); + r3 = compute_coordinates(dh_, (x + task.lb_cube.x + task.roffset.x), + (y + task.lb_cube.y + task.roffset.y), + (z + task.lb_cube.z + task.roffset.z)); } if (distributed__) { - // check if the point is inside the sphere or not. Note that it does not - // apply for the orthorhombic case when the full sphere is inside the - // region of interest. + // check if the point is inside the sphere or not. Note that it does + // not apply for the orthorhombic case when the full sphere is inside + // the region of interest. if (((task.radius * task.radius) <= (r3.x * r3.x + r3.y * r3.y + r3.z * r3.z)) && @@ -408,32 +408,36 @@ void context_info::collocate_one_grid_level(const int level, if (func == GRID_FUNC_AB) { calculate_coefficients - <<>>( - params); + <<>>(params); } else { calculate_coefficients - <<>>( - params); + <<>>(params); } if (grid_[level].is_distributed()) { if (grid_[level].is_orthorhombic()) collocate_kernel - <<>>(params); + <<>>( + params); else collocate_kernel - <<>>(params); + <<>>( + params); } else { if (grid_[level].is_orthorhombic()) collocate_kernel - <<>>(params); + <<>>( + params); else collocate_kernel - <<>>(params); + <<>>( + params); } } } // namespace rocm_backend diff --git a/src/grid/hip/grid_hip_context.cc b/src/grid/hip/grid_hip_context.cc index 59eeb2e435..65cab6c0a0 100644 --- a/src/grid/hip/grid_hip_context.cc +++ b/src/grid/hip/grid_hip_context.cc @@ -188,7 +188,7 @@ extern "C" void grid_hip_create_task_list( tasks_host[i].block_transposed = (iatom > jatom); tasks_host[i].subblock_offset = - (tasks_host[i].block_transposed) + (tasks_host[i].block_transposed) ? (tasks_host[i].sgfa * tasks_host[i].nsgfb + tasks_host[i].sgfb) : (tasks_host[i].sgfb * tasks_host[i].nsgfa + tasks_host[i].sgfa); @@ -199,13 +199,15 @@ extern "C" void grid_hip_create_task_list( /* this block is only as temporary scratch for calculating the coefficients. * Doing this avoid a lot of atomic operations that are costly on hardware * that only have partial support of them. For better performance we should - * most probably align the offsets as well. it is 256 bytes on Mi100 and above */ + * most probably align the offsets as well. it is 256 bytes on Mi100 and + * above */ tasks_host[i].lp_max = tasks_host[i].lb_max + tasks_host[i].la_max + 6; if (i == 0) { tasks_host[i].coef_offset = 0; } else { - tasks_host[i].coef_offset = tasks_host[i - 1].coef_offset + - rocm_backend::ncoset(tasks_host[i - 1].lp_max); + tasks_host[i].coef_offset = + tasks_host[i - 1].coef_offset + + rocm_backend::ncoset(tasks_host[i - 1].lp_max); } coef_size += rocm_backend::ncoset(tasks_host[i].lp_max); @@ -215,44 +217,43 @@ extern "C" void grid_hip_create_task_list( tasks_host[i].apply_border_mask = (tasks_host[i].border_mask != 0); if (grid.is_orthorhombic() && (tasks_host[i].border_mask == 0)) { - tasks_host[i] - .discrete_radius = rocm_backend::compute_cube_properties( - tasks_host[i].radius, grid.dh(), grid.dh_inv(), - (double3 *)tasks_host[i].rp, // center of the gaussian - &tasks_host[i].roffset, // offset compared to the closest grid point - &tasks_host[i].cube_center, // center coordinates in grid space - &tasks_host[i].lb_cube, // lower boundary - &tasks_host[i].cube_size); + tasks_host[i].discrete_radius = + rocm_backend::compute_cube_properties( + tasks_host[i].radius, grid.dh(), grid.dh_inv(), + (double3 *)tasks_host[i].rp, // center of the gaussian + &tasks_host[i] + .roffset, // offset compared to the closest grid point + &tasks_host[i].cube_center, // center coordinates in grid space + &tasks_host[i].lb_cube, // lower boundary + &tasks_host[i].cube_size); } else { - tasks_host[i] - .discrete_radius = rocm_backend::compute_cube_properties( - tasks_host[i].radius, grid.dh(), grid.dh_inv(), - (double3 *)tasks_host[i].rp, // center of the gaussian - &tasks_host[i].roffset, // offset compared to the closest grid point - &tasks_host[i].cube_center, // center coordinates in grid space - &tasks_host[i].lb_cube, // lower boundary - &tasks_host[i].cube_size); + tasks_host[i].discrete_radius = + rocm_backend::compute_cube_properties( + tasks_host[i].radius, grid.dh(), grid.dh_inv(), + (double3 *)tasks_host[i].rp, // center of the gaussian + &tasks_host[i] + .roffset, // offset compared to the closest grid point + &tasks_host[i].cube_center, // center coordinates in grid space + &tasks_host[i].lb_cube, // lower boundary + &tasks_host[i].cube_size); } } // we need to sort the task list although I expect it to be sorted already -/* - * sorting with this lambda does not work -std::sort(tasks_host.begin(), tasks_host.end(), [](rocm_backend::task_info a, rocm_backend::task_info b) { - if (a.level == b.level) { - if (a.block_num <= b.block_num) - return true; - else - return false; - } else { - return (a.level < b.level); - } - }); -*/ + /* + * sorting with this lambda does not work + std::sort(tasks_host.begin(), tasks_host.end(), [](rocm_backend::task_info a, + rocm_backend::task_info b) { if (a.level == b.level) { if (a.block_num <= + b.block_num) return true; else return false; } else { return (a.level < + b.level); + } + }); + */ // it is a exclusive scan actually for (int level = 1; level < ctx->number_of_tasks_per_level_.size(); level++) { - ctx->first_task_per_level_[level] = ctx->first_task_per_level_[level - 1] + - ctx->number_of_tasks_per_level_[level - 1]; + ctx->first_task_per_level_[level] = + ctx->first_task_per_level_[level - 1] + + ctx->number_of_tasks_per_level_[level - 1]; } ctx->tasks_dev.clear(); @@ -323,7 +324,7 @@ std::sort(tasks_host.begin(), tasks_host.end(), [](rocm_backend::task_info a, ro ctx->num_tasks_per_block_dev_.resize(num_tasks_per_block.size()); ctx->num_tasks_per_block_dev_.copy_to_gpu(num_tasks_per_block); -// collect stats + // collect stats memset(ctx->stats, 0, 2 * 20 * sizeof(int)); for (int itask = 0; itask < ntasks; itask++) { const int iatom = iatom_list[itask] - 1; @@ -440,7 +441,7 @@ extern "C" void grid_hip_integrate_task_list( rocm_backend::context_info *ctx = (rocm_backend::context_info *)ptr; - if(ptr == nullptr) + if (ptr == nullptr) return; assert(ctx->nlevels == nlevels); diff --git a/src/grid/hip/grid_hip_integrate.cc b/src/grid/hip/grid_hip_integrate.cc index 2bdfb389de..1b3976fa8f 100644 --- a/src/grid/hip/grid_hip_integrate.cc +++ b/src/grid/hip/grid_hip_integrate.cc @@ -148,22 +148,23 @@ __global__ __launch_bounds__(64) void compute_hab_v2(const kernel_params dev_, if (CALCULATE_FORCES) { for (int k = 0; k < 3; k++) { fa[k] += block_val1 * sphia * sphib * - get_force_a(a, b, k, task.zeta, task.zetb, - task.n1, smem_cab); - fb[k] += block_val1 * sphia * sphib * - get_force_b(a, b, k, task.zeta, task.zetb, - task.rab, task.n1, smem_cab); + get_force_a( + a, b, k, task.zeta, task.zetb, task.n1, smem_cab); + fb[k] += + block_val1 * sphia * sphib * + get_force_b(a, b, k, task.zeta, task.zetb, + task.rab, task.n1, smem_cab); } if (dev_.ptr_dev[5] != nullptr) { for (int k = 0; k < 3; k++) { for (int l = 0; l < 3; l++) { virial[3 * k + l] += - (get_virial_a(a, b, k, l, task.zeta, - task.zetb, task.n1, - smem_cab) + - get_virial_b(a, b, k, l, task.zeta, - task.zetb, task.rab, task.n1, - smem_cab)) * + (get_virial_a(a, b, k, l, task.zeta, + task.zetb, task.n1, + smem_cab) + + get_virial_b(a, b, k, l, task.zeta, + task.zetb, task.rab, + task.n1, smem_cab)) * block_val1 * sphia * sphib; } } @@ -256,8 +257,10 @@ So most of the code is the same except the core of the routine that is specialized to the integration. ******************************************************************************/ - template -__global__ __launch_bounds__(64) void integrate_kernel(const kernel_params dev_) { +template +__global__ +__launch_bounds__(64) void integrate_kernel(const kernel_params dev_) { if (dev_.tasks[dev_.first_task + blockIdx.x].skip_task) return; @@ -283,11 +286,11 @@ __global__ __launch_bounds__(64) void integrate_kernel(const kernel_params dev_) if (distributed__) { if (task.apply_border_mask) { - compute_window_size(dev_.grid_local_size_, dev_.grid_lower_corner_, - dev_.grid_full_size_, /* also full size of the grid */ - dev_.tasks[dev_.first_task + blockIdx.x].border_mask, - dev_.grid_border_width_, &task.window_size, - &task.window_shift); + compute_window_size( + dev_.grid_local_size_, dev_.grid_lower_corner_, + dev_.grid_full_size_, /* also full size of the grid */ + dev_.tasks[dev_.first_task + blockIdx.x].border_mask, + dev_.grid_border_width_, &task.window_size, &task.window_shift); } } } @@ -347,7 +350,7 @@ __global__ __launch_bounds__(64) void integrate_kernel(const kernel_params dev_) y2 += dev_.grid_full_size_[1]; if (distributed__) { -/* check if the point is within the window */ + /* check if the point is within the window */ if (task.apply_border_mask) { if ((y2 < task.window_shift.y) || (y2 > task.window_size.y)) { continue; @@ -392,10 +395,9 @@ __global__ __launch_bounds__(64) void integrate_kernel(const kernel_params dev_) r3.y = (y + task.lb_cube.y + task.roffset.y) * dh_[4]; r3.z = (z + task.lb_cube.z + task.roffset.z) * dh_[8]; } else { - r3 = - compute_coordinates(dh_, (x + task.lb_cube.x + task.roffset.x), - (y + task.lb_cube.y + task.roffset.y), - (z + task.lb_cube.z + task.roffset.z)); + r3 = compute_coordinates(dh_, (x + task.lb_cube.x + task.roffset.x), + (y + task.lb_cube.y + task.roffset.y), + (z + task.lb_cube.z + task.roffset.z)); } // check if the point is inside the sphere or not. Note that it does // not apply for the orthorhombic case when the full sphere is inside @@ -605,49 +607,51 @@ __global__ __launch_bounds__(64) void integrate_kernel(const kernel_params dev_) } } -kernel_params context_info::set_kernel_parameters(const int level, const smem_parameters &smem_params) { - kernel_params params; - params.smem_cab_offset = smem_params.smem_cab_offset(); - params.smem_alpha_offset = smem_params.smem_alpha_offset(); - params.first_task = 0; +kernel_params +context_info::set_kernel_parameters(const int level, + const smem_parameters &smem_params) { + kernel_params params; + params.smem_cab_offset = smem_params.smem_cab_offset(); + params.smem_alpha_offset = smem_params.smem_alpha_offset(); + params.first_task = 0; - params.la_min_diff = smem_params.ldiffs().la_min_diff; - params.lb_min_diff = smem_params.ldiffs().lb_min_diff; + params.la_min_diff = smem_params.ldiffs().la_min_diff; + params.lb_min_diff = smem_params.ldiffs().lb_min_diff; - params.la_max_diff = smem_params.ldiffs().la_max_diff; - params.lb_max_diff = smem_params.ldiffs().lb_max_diff; - params.first_task = 0; - params.tasks = this->tasks_dev.data(); - params.task_sorted_by_blocks_dev = task_sorted_by_blocks_dev.data(); - params.sorted_blocks_offset_dev = sorted_blocks_offset_dev.data(); - params.num_tasks_per_block_dev = this->num_tasks_per_block_dev_.data(); - params.block_offsets = this->block_offsets_dev.data(); - params.la_min_diff = smem_params.ldiffs().la_min_diff; - params.lb_min_diff = smem_params.ldiffs().lb_min_diff; - params.la_max_diff = smem_params.ldiffs().la_max_diff; - params.lb_max_diff = smem_params.ldiffs().lb_max_diff; + params.la_max_diff = smem_params.ldiffs().la_max_diff; + params.lb_max_diff = smem_params.ldiffs().lb_max_diff; + params.first_task = 0; + params.tasks = this->tasks_dev.data(); + params.task_sorted_by_blocks_dev = task_sorted_by_blocks_dev.data(); + params.sorted_blocks_offset_dev = sorted_blocks_offset_dev.data(); + params.num_tasks_per_block_dev = this->num_tasks_per_block_dev_.data(); + params.block_offsets = this->block_offsets_dev.data(); + params.la_min_diff = smem_params.ldiffs().la_min_diff; + params.lb_min_diff = smem_params.ldiffs().lb_min_diff; + params.la_max_diff = smem_params.ldiffs().la_max_diff; + params.lb_max_diff = smem_params.ldiffs().lb_max_diff; - params.ptr_dev[0] = pab_block_.data(); + params.ptr_dev[0] = pab_block_.data(); - if (level >= 0) { - params.ptr_dev[1] = grid_[level].data(); - for (int i = 0; i < 3; i++) { - memcpy(params.dh_, grid_[level].dh(), 9 * sizeof(double)); - memcpy(params.dh_inv_, grid_[level].dh_inv(), 9 * sizeof(double)); - params.grid_full_size_[i] = grid_[level].full_size(i); - params.grid_local_size_[i] = grid_[level].local_size(i); - params.grid_lower_corner_[i] = grid_[level].lower_corner(i); - params.grid_border_width_[i] = grid_[level].border_width(i); - } - params.first_task = first_task_per_level_[level]; + if (level >= 0) { + params.ptr_dev[1] = grid_[level].data(); + for (int i = 0; i < 3; i++) { + memcpy(params.dh_, grid_[level].dh(), 9 * sizeof(double)); + memcpy(params.dh_inv_, grid_[level].dh_inv(), 9 * sizeof(double)); + params.grid_full_size_[i] = grid_[level].full_size(i); + params.grid_local_size_[i] = grid_[level].local_size(i); + params.grid_lower_corner_[i] = grid_[level].lower_corner(i); + params.grid_border_width_[i] = grid_[level].border_width(i); } - params.ptr_dev[2] = this->coef_dev_.data(); - params.ptr_dev[3] = hab_block_.data(); - params.ptr_dev[4] = forces_.data(); - params.ptr_dev[5] = virial_.data(); - params.sphi_dev = this->sphi_dev.data(); - return params; + params.first_task = first_task_per_level_[level]; } + params.ptr_dev[2] = this->coef_dev_.data(); + params.ptr_dev[3] = hab_block_.data(); + params.ptr_dev[4] = forces_.data(); + params.ptr_dev[5] = virial_.data(); + params.sphi_dev = this->sphi_dev.data(); + return params; +} /******************************************************************************* * \brief Launches the Cuda kernel that integrates all tasks of one grid level. @@ -678,18 +682,22 @@ void context_info::integrate_one_grid_level(const int level, int *lp_diff) { if (grid_[level].is_distributed()) { if (grid_[level].is_orthorhombic()) { integrate_kernel - <<>>(params); + <<>>(params); } else { integrate_kernel - <<>>(params); + <<>>(params); } } else { if (grid_[level].is_orthorhombic()) { integrate_kernel - <<>>(params); + <<>>(params); } else { integrate_kernel - <<>>(params); + <<>>(params); } } } @@ -700,7 +708,7 @@ void context_info::compute_hab_coefficients() { // Compute max angular momentum. const ldiffs_value ldiffs = - process_get_ldiffs(calculate_forces, calculate_virial, compute_tau); + process_get_ldiffs(calculate_forces, calculate_virial, compute_tau); smem_parameters smem_params(ldiffs, lmax()); init_constant_memory(); @@ -715,28 +723,28 @@ void context_info::compute_hab_coefficients() { if (!compute_tau && !calculate_forces) { compute_hab_v2 - <<nblocks, threads_per_block, smem_params.smem_per_block(), this->main_stream>>>( - params, this->ntasks); + <<nblocks, threads_per_block, smem_params.smem_per_block(), + this->main_stream>>>(params, this->ntasks); return; } if (!compute_tau && calculate_forces) { compute_hab_v2 - <<nblocks, threads_per_block, smem_params.smem_per_block(), this->main_stream>>>( - params, this->ntasks); + <<nblocks, threads_per_block, smem_params.smem_per_block(), + this->main_stream>>>(params, this->ntasks); return; } if (compute_tau && calculate_forces) { compute_hab_v2 - <<nblocks, threads_per_block, smem_params.smem_per_block(), this->main_stream>>>( - params, this->ntasks); + <<nblocks, threads_per_block, smem_params.smem_per_block(), + this->main_stream>>>(params, this->ntasks); } if (compute_tau && !calculate_forces) { compute_hab_v2 - <<nblocks, threads_per_block, smem_params.smem_per_block(), this->main_stream>>>( - params, this->ntasks); + <<nblocks, threads_per_block, smem_params.smem_per_block(), + this->main_stream>>>(params, this->ntasks); } } }; // namespace rocm_backend diff --git a/tools/precommit/check_file_properties.py b/tools/precommit/check_file_properties.py index 89275969a4..5edc1f8636 100755 --- a/tools/precommit/check_file_properties.py +++ b/tools/precommit/check_file_properties.py @@ -94,7 +94,7 @@ BANNER_C = """\ /*----------------------------------------------------------------------------*/ """ -C_FAMILY_EXTENSIONS = (".c", ".cu", ".cpp", ".h", ".hpp") +C_FAMILY_EXTENSIONS = (".c", ".cu", ".cpp", ".cc", ".h", ".hpp") @lru_cache(maxsize=None) diff --git a/tools/precommit/precommit.py b/tools/precommit/precommit.py index c29eb462b6..7f90dd2a97 100755 --- a/tools/precommit/precommit.py +++ b/tools/precommit/precommit.py @@ -183,7 +183,7 @@ def process_file(fn, allow_modifications): run_local_tool("./tools/doxify/doxify.sh", fn) run_prettify(fn) - if re.match(r".*\.(c|cu|cpp|h|hpp)$", fn): + if re.match(r".*\.(c|cu|cpp|cc|h|hpp)$", fn): run_remote_tool("clangformat", fn) if re.match(r"(.*/PACKAGE)|(.*\.py)$", fn):