Skip to content

Commit

Permalink
Auto-format code changes (#725)
Browse files Browse the repository at this point in the history 
Auto-format code using Clang-Format

Co-authored-by: GitHub Actions <actions@github.com>
  • Loading branch information
@github-actions @actions-user
github-actions[bot] and actions-user authored Feb 11, 2025
1 parent cbdc7bf commit 1b937ab
Show file tree
Hide file tree
Showing 5 changed files with 124 additions and 198 deletions.
2 changes: 1 addition & 1 deletion include/micm/cuda/process/cuda_process_set.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -115,7 +115,7 @@ namespace micm
ProcessSetParam hoststruct;
std::vector<ProcessInfoParam> jacobian_process_info(this->jacobian_process_info_.size());
std::size_t i_process = 0;
for (const auto & process_info : this->jacobian_process_info_)
for (const auto& process_info : this->jacobian_process_info_)
{
jacobian_process_info[i_process].process_id_ = process_info.process_id_;
jacobian_process_info[i_process].independent_id_ = process_info.independent_id_;
Expand Down
6 changes: 3 additions & 3 deletions include/micm/jit/process/jit_process_set.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -225,7 +225,7 @@ namespace micm
auto react_ids = jacobian_reactant_ids_.begin();
auto yields = jacobian_yields_.begin();
auto flat_id = jacobian_flat_ids_.begin();
for (const auto& process_info : jacobian_process_info_)
for (const auto &process_info : jacobian_process_info_)
{
llvm::Value *rc_start = llvm::ConstantInt::get(*(func.context_), llvm::APInt(64, process_info.process_id_ * L));
// save rate constant in d_rate_d_ind for each grid cell
Expand Down Expand Up @@ -256,7 +256,7 @@ namespace micm
}

// set jacobian terms for each reactant jac[i_react][i_ind][i_cell] += d_rate_d_ind[i_cell]
for (std::size_t i_dep = 0; i_dep < process_info.number_of_dependent_reactants_+1; ++i_dep)
for (std::size_t i_dep = 0; i_dep < process_info.number_of_dependent_reactants_ + 1; ++i_dep)
{
loop = func.StartLoop("reactant term", 0, L);
llvm::Value *dep_id = llvm::ConstantInt::get(*(func.context_), llvm::APInt(64, *(flat_id++)));
Expand All @@ -273,7 +273,7 @@ namespace micm

// set jacobian terms for each product jac[i_prod][i_ind][i_cell] -= yield * d_rate_d_ind[i_cell]
for (std::size_t i_dep = 0; i_dep < process_info.number_of_products_; ++i_dep)
{
{
loop = func.StartLoop("product term", 0, L);
llvm::Value *dep_id = llvm::ConstantInt::get(*(func.context_), llvm::APInt(64, *(flat_id++)));
ptr_index[0] = func.builder_->CreateNSWAdd(loop.index_, dep_id);
Expand Down
8 changes: 3 additions & 5 deletions include/micm/process/process_set.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -189,9 +189,7 @@ namespace micm

// Set up process information for Jacobian calculations
std::vector<std::pair<std::string, std::size_t>> sorted_names(variable_map.begin(), variable_map.end());
std::sort(sorted_names.begin(), sorted_names.end(), [](const auto& a, const auto& b) {
return a.second < b.second;
});
std::sort(sorted_names.begin(), sorted_names.end(), [](const auto& a, const auto& b) { return a.second < b.second; });
for (const auto& independent_variable : sorted_names)
{
for (std::size_t i_process = 0; i_process < processes.size(); ++i_process)
Expand Down Expand Up @@ -415,7 +413,7 @@ namespace micm
double d_rate_d_ind = cell_rate_constants[process_info.process_id_];
for (std::size_t i_react = 0; i_react < process_info.number_of_dependent_reactants_; ++i_react)
d_rate_d_ind *= cell_state[*(react_id++)];
for (std::size_t i_dep = 0; i_dep < process_info.number_of_dependent_reactants_+1; ++i_dep)
for (std::size_t i_dep = 0; i_dep < process_info.number_of_dependent_reactants_ + 1; ++i_dep)
cell_jacobian[*(flat_id++)] += d_rate_d_ind;
for (std::size_t i_dep = 0; i_dep < process_info.number_of_products_; ++i_dep)
cell_jacobian[*(flat_id++)] -= *(yield++) * d_rate_d_ind;
Expand Down Expand Up @@ -463,7 +461,7 @@ namespace micm
for (std::size_t i_cell = 0; i_cell < L; ++i_cell)
*(v_d_rate_d_ind_it++) *= *(v_state_variables_it++);
}
for (std::size_t i_dep = 0; i_dep < process_info.number_of_dependent_reactants_+1; ++i_dep)
for (std::size_t i_dep = 0; i_dep < process_info.number_of_dependent_reactants_ + 1; ++i_dep)
{
auto v_jacobian_it = v_jacobian.begin() + offset_jacobian + *flat_id;
auto v_d_rate_d_ind_it = d_rate_d_ind.begin();
Expand Down
289 changes: 108 additions & 181 deletions src/process/process_set.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -95,7 +95,7 @@ namespace micm
{
d_rate_d_ind *= d_state_variables[(d_reactant_ids[react_ids_offset + i_react] * number_of_grid_cells) + tid];
}
for (size_t i_dep = 0; i_dep < process_info.number_of_dependent_reactants_+1; ++i_dep)
for (size_t i_dep = 0; i_dep < process_info.number_of_dependent_reactants_ + 1; ++i_dep)
{
size_t jacobian_idx = d_jacobian_flat_ids[flat_id_offset] + tid;
d_jacobian[jacobian_idx] += d_rate_d_ind;
Expand Down Expand Up @@ -130,75 +130,43 @@ namespace micm

auto cuda_stream_id = micm::cuda::CudaStreamSingleton::GetInstance().GetCudaStream(0);

/// Allocate memory space on the device
CHECK_CUDA_ERROR(
cudaMallocAsync(
&(devstruct.number_of_reactants_),
number_of_reactants_bytes,
cuda_stream_id),
"cudaMalloc");
CHECK_CUDA_ERROR(
cudaMallocAsync(
&(devstruct.reactant_ids_),
reactant_ids_bytes,
cuda_stream_id),
"cudaMalloc");
CHECK_CUDA_ERROR(
cudaMallocAsync(
&(devstruct.number_of_products_),
number_of_products_bytes,
cuda_stream_id),
"cudaMalloc");
CHECK_CUDA_ERROR(
cudaMallocAsync(
&(devstruct.product_ids_), product_ids_bytes, cuda_stream_id),
"cudaMalloc");
CHECK_CUDA_ERROR(
cudaMallocAsync(
&(devstruct.yields_), yields_bytes, cuda_stream_id),
"cudaMalloc");
/// Allocate memory space on the device
CHECK_CUDA_ERROR(
cudaMallocAsync(&(devstruct.number_of_reactants_), number_of_reactants_bytes, cuda_stream_id), "cudaMalloc");
CHECK_CUDA_ERROR(cudaMallocAsync(&(devstruct.reactant_ids_), reactant_ids_bytes, cuda_stream_id), "cudaMalloc");
CHECK_CUDA_ERROR(
cudaMallocAsync(&(devstruct.number_of_products_), number_of_products_bytes, cuda_stream_id), "cudaMalloc");
CHECK_CUDA_ERROR(cudaMallocAsync(&(devstruct.product_ids_), product_ids_bytes, cuda_stream_id), "cudaMalloc");
CHECK_CUDA_ERROR(cudaMallocAsync(&(devstruct.yields_), yields_bytes, cuda_stream_id), "cudaMalloc");

/// Copy the data from host to device
CHECK_CUDA_ERROR(
cudaMemcpyAsync(
devstruct.number_of_reactants_,
hoststruct.number_of_reactants_,
number_of_reactants_bytes,
cudaMemcpyHostToDevice,
cuda_stream_id),
"cudaMemcpy");
CHECK_CUDA_ERROR(
cudaMemcpyAsync(
devstruct.reactant_ids_,
hoststruct.reactant_ids_,
reactant_ids_bytes,
cudaMemcpyHostToDevice,
cuda_stream_id),
"cudaMemcpy");
CHECK_CUDA_ERROR(
cudaMemcpyAsync(
devstruct.number_of_products_,
hoststruct.number_of_products_,
number_of_products_bytes,
cudaMemcpyHostToDevice,
cuda_stream_id),
"cudaMemcpy");
CHECK_CUDA_ERROR(
cudaMemcpyAsync(
devstruct.product_ids_,
hoststruct.product_ids_,
product_ids_bytes,
cudaMemcpyHostToDevice,
cuda_stream_id),
"cudaMemcpy");
CHECK_CUDA_ERROR(
cudaMemcpyAsync(
devstruct.yields_,
hoststruct.yields_,
yields_bytes,
cudaMemcpyHostToDevice,
cuda_stream_id),
"cudaMemcpy");
/// Copy the data from host to device
CHECK_CUDA_ERROR(
cudaMemcpyAsync(
devstruct.number_of_reactants_,
hoststruct.number_of_reactants_,
number_of_reactants_bytes,
cudaMemcpyHostToDevice,
cuda_stream_id),
"cudaMemcpy");
CHECK_CUDA_ERROR(
cudaMemcpyAsync(
devstruct.reactant_ids_, hoststruct.reactant_ids_, reactant_ids_bytes, cudaMemcpyHostToDevice, cuda_stream_id),
"cudaMemcpy");
CHECK_CUDA_ERROR(
cudaMemcpyAsync(
devstruct.number_of_products_,
hoststruct.number_of_products_,
number_of_products_bytes,
cudaMemcpyHostToDevice,
cuda_stream_id),
"cudaMemcpy");
CHECK_CUDA_ERROR(
cudaMemcpyAsync(
devstruct.product_ids_, hoststruct.product_ids_, product_ids_bytes, cudaMemcpyHostToDevice, cuda_stream_id),
"cudaMemcpy");
CHECK_CUDA_ERROR(
cudaMemcpyAsync(devstruct.yields_, hoststruct.yields_, yields_bytes, cudaMemcpyHostToDevice, cuda_stream_id),
"cudaMemcpy");

devstruct.number_of_reactants_size_ = hoststruct.number_of_reactants_size_;
devstruct.reactant_ids_size_ = hoststruct.reactant_ids_size_;
Expand All @@ -214,8 +182,7 @@ namespace micm
void CopyJacobianParams(ProcessSetParam& hoststruct, ProcessSetParam& devstruct)
{
/// Calculate the memory space
size_t jacobian_process_info_bytes = sizeof(ProcessInfoParam) *
hoststruct.jacobian_process_info_size_;
size_t jacobian_process_info_bytes = sizeof(ProcessInfoParam) * hoststruct.jacobian_process_info_size_;
size_t jacobian_reactant_ids_bytes = sizeof(size_t) * hoststruct.jacobian_reactant_ids_size_;
size_t jacobian_product_ids_bytes = sizeof(size_t) * hoststruct.jacobian_product_ids_size_;
size_t jacobian_yields_bytes = sizeof(double) * hoststruct.jacobian_yields_size_;
Expand All @@ -224,78 +191,57 @@ namespace micm
auto cuda_stream_id = micm::cuda::CudaStreamSingleton::GetInstance().GetCudaStream(0);

/// Allocate memory space on the device
CHECK_CUDA_ERROR(
cudaMallocAsync(
&(devstruct.jacobian_process_info_),
jacobian_process_info_bytes,
cuda_stream_id),
"cudaMalloc");
CHECK_CUDA_ERROR(
cudaMallocAsync(
&(devstruct.jacobian_reactant_ids_),
jacobian_reactant_ids_bytes,
cuda_stream_id),
"cudaMalloc");
CHECK_CUDA_ERROR(
cudaMallocAsync(
&(devstruct.jacobian_product_ids_),
jacobian_product_ids_bytes,
cuda_stream_id),
"cudaMalloc");
CHECK_CUDA_ERROR(
cudaMallocAsync(
&(devstruct.jacobian_yields_),
jacobian_yields_bytes,
cuda_stream_id),
"cudaMalloc");
CHECK_CUDA_ERROR(
cudaMallocAsync(
&(devstruct.jacobian_flat_ids_),
jacobian_flat_ids_bytes,
cuda_stream_id),
"cudaMalloc");
CHECK_CUDA_ERROR(
cudaMallocAsync(&(devstruct.jacobian_process_info_), jacobian_process_info_bytes, cuda_stream_id), "cudaMalloc");
CHECK_CUDA_ERROR(
cudaMallocAsync(&(devstruct.jacobian_reactant_ids_), jacobian_reactant_ids_bytes, cuda_stream_id), "cudaMalloc");
CHECK_CUDA_ERROR(
cudaMallocAsync(&(devstruct.jacobian_product_ids_), jacobian_product_ids_bytes, cuda_stream_id), "cudaMalloc");
CHECK_CUDA_ERROR(cudaMallocAsync(&(devstruct.jacobian_yields_), jacobian_yields_bytes, cuda_stream_id), "cudaMalloc");
CHECK_CUDA_ERROR(
cudaMallocAsync(&(devstruct.jacobian_flat_ids_), jacobian_flat_ids_bytes, cuda_stream_id), "cudaMalloc");

/// Copy the data from host to device
CHECK_CUDA_ERROR(
cudaMemcpyAsync(
devstruct.jacobian_process_info_,
hoststruct.jacobian_process_info_,
jacobian_process_info_bytes,
cudaMemcpyHostToDevice,
cuda_stream_id),
"cudaMemcpy");
CHECK_CUDA_ERROR(
cudaMemcpyAsync(
devstruct.jacobian_reactant_ids_,
hoststruct.jacobian_reactant_ids_,
jacobian_reactant_ids_bytes,
cudaMemcpyHostToDevice,
cuda_stream_id),
"cudaMemcpy");
CHECK_CUDA_ERROR(
cudaMemcpyAsync(
devstruct.jacobian_product_ids_,
hoststruct.jacobian_product_ids_,
jacobian_product_ids_bytes,
cudaMemcpyHostToDevice,
cuda_stream_id),
"cudaMemcpy");
CHECK_CUDA_ERROR(
cudaMemcpyAsync(
devstruct.jacobian_yields_,
hoststruct.jacobian_yields_,
jacobian_yields_bytes,
cudaMemcpyHostToDevice,
cuda_stream_id),
"cudaMemcpy");
CHECK_CUDA_ERROR(
cudaMemcpyAsync(
devstruct.jacobian_flat_ids_,
hoststruct.jacobian_flat_ids_,
jacobian_flat_ids_bytes,
cudaMemcpyHostToDevice,
cuda_stream_id),
"cudaMemcpy");
/// Copy the data from host to device
CHECK_CUDA_ERROR(
cudaMemcpyAsync(
devstruct.jacobian_process_info_,
hoststruct.jacobian_process_info_,
jacobian_process_info_bytes,
cudaMemcpyHostToDevice,
cuda_stream_id),
"cudaMemcpy");
CHECK_CUDA_ERROR(
cudaMemcpyAsync(
devstruct.jacobian_reactant_ids_,
hoststruct.jacobian_reactant_ids_,
jacobian_reactant_ids_bytes,
cudaMemcpyHostToDevice,
cuda_stream_id),
"cudaMemcpy");
CHECK_CUDA_ERROR(
cudaMemcpyAsync(
devstruct.jacobian_product_ids_,
hoststruct.jacobian_product_ids_,
jacobian_product_ids_bytes,
cudaMemcpyHostToDevice,
cuda_stream_id),
"cudaMemcpy");
CHECK_CUDA_ERROR(
cudaMemcpyAsync(
devstruct.jacobian_yields_,
hoststruct.jacobian_yields_,
jacobian_yields_bytes,
cudaMemcpyHostToDevice,
cuda_stream_id),
"cudaMemcpy");
CHECK_CUDA_ERROR(
cudaMemcpyAsync(
devstruct.jacobian_flat_ids_,
hoststruct.jacobian_flat_ids_,
jacobian_flat_ids_bytes,
cudaMemcpyHostToDevice,
cuda_stream_id),
"cudaMemcpy");

devstruct.jacobian_process_info_size_ = hoststruct.jacobian_process_info_size_;
devstruct.jacobian_reactant_ids_size_ = hoststruct.jacobian_reactant_ids_size_;
Expand All @@ -310,45 +256,26 @@ namespace micm
{
auto cuda_stream_id = micm::cuda::CudaStreamSingleton::GetInstance().GetCudaStream(0);

if (devstruct.number_of_reactants_ != nullptr)
CHECK_CUDA_ERROR(
cudaFreeAsync(devstruct.number_of_reactants_, cuda_stream_id),
"cudaFree");
if (devstruct.reactant_ids_ != nullptr)
CHECK_CUDA_ERROR(
cudaFreeAsync(devstruct.reactant_ids_, cuda_stream_id),
"cudaFree");
if (devstruct.number_of_products_ != nullptr)
CHECK_CUDA_ERROR(
cudaFreeAsync(devstruct.number_of_products_, cuda_stream_id),
"cudaFree");
if (devstruct.product_ids_ != nullptr)
CHECK_CUDA_ERROR(
cudaFreeAsync(devstruct.product_ids_, cuda_stream_id),
"cudaFree");
if (devstruct.yields_ != nullptr)
CHECK_CUDA_ERROR(
cudaFreeAsync(devstruct.yields_, cuda_stream_id), "cudaFree");
if (devstruct.jacobian_process_info_ != nullptr)
CHECK_CUDA_ERROR(
cudaFreeAsync(devstruct.jacobian_process_info_, cuda_stream_id),
"cudaFree");
if (devstruct.jacobian_reactant_ids_ != nullptr)
CHECK_CUDA_ERROR(
cudaFreeAsync(devstruct.jacobian_reactant_ids_, cuda_stream_id),
"cudaFree");
if (devstruct.jacobian_product_ids_ != nullptr)
CHECK_CUDA_ERROR(
cudaFreeAsync(devstruct.jacobian_product_ids_, cuda_stream_id),
"cudaFree");
if (devstruct.jacobian_yields_ != nullptr)
CHECK_CUDA_ERROR(
cudaFreeAsync(devstruct.jacobian_yields_, cuda_stream_id),
"cudaFree");
if (devstruct.jacobian_flat_ids_ != nullptr)
CHECK_CUDA_ERROR(
cudaFreeAsync(devstruct.jacobian_flat_ids_, cuda_stream_id),
"cudaFree");
if (devstruct.number_of_reactants_ != nullptr)
CHECK_CUDA_ERROR(cudaFreeAsync(devstruct.number_of_reactants_, cuda_stream_id), "cudaFree");
if (devstruct.reactant_ids_ != nullptr)
CHECK_CUDA_ERROR(cudaFreeAsync(devstruct.reactant_ids_, cuda_stream_id), "cudaFree");
if (devstruct.number_of_products_ != nullptr)
CHECK_CUDA_ERROR(cudaFreeAsync(devstruct.number_of_products_, cuda_stream_id), "cudaFree");
if (devstruct.product_ids_ != nullptr)
CHECK_CUDA_ERROR(cudaFreeAsync(devstruct.product_ids_, cuda_stream_id), "cudaFree");
if (devstruct.yields_ != nullptr)
CHECK_CUDA_ERROR(cudaFreeAsync(devstruct.yields_, cuda_stream_id), "cudaFree");
if (devstruct.jacobian_process_info_ != nullptr)
CHECK_CUDA_ERROR(cudaFreeAsync(devstruct.jacobian_process_info_, cuda_stream_id), "cudaFree");
if (devstruct.jacobian_reactant_ids_ != nullptr)
CHECK_CUDA_ERROR(cudaFreeAsync(devstruct.jacobian_reactant_ids_, cuda_stream_id), "cudaFree");
if (devstruct.jacobian_product_ids_ != nullptr)
CHECK_CUDA_ERROR(cudaFreeAsync(devstruct.jacobian_product_ids_, cuda_stream_id), "cudaFree");
if (devstruct.jacobian_yields_ != nullptr)
CHECK_CUDA_ERROR(cudaFreeAsync(devstruct.jacobian_yields_, cuda_stream_id), "cudaFree");
if (devstruct.jacobian_flat_ids_ != nullptr)
CHECK_CUDA_ERROR(cudaFreeAsync(devstruct.jacobian_flat_ids_, cuda_stream_id), "cudaFree");
}

void SubtractJacobianTermsKernelDriver(
Expand Down
Loading

0 comments on commit 1b937ab

Please sign in to comment.