added tid_array for debugging

NCAR · Jun 24, 2023 · 4d5674c · 4d5674c
1 parent 296989f
commit 4d5674c
Show file tree

Hide file tree

Showing 2 changed files with 34 additions and 6 deletions.
diff --git a/src/process/process_set.cu b/src/process/process_set.cu
@@ -8,6 +8,7 @@ namespace micm {
     //one thread per reaction
     //passing all device pointers 
     __global__ void AddForcingTerms_kernel(
+        int* tid_array,
         double* rate_constants, 
         double* state_variables, 
         double* forcing, 
@@ -28,14 +29,16 @@ namespace micm {
     int rate_constants_size = matrix_rows * rate_constants_columns; 
 
     if (tid < rate_constants_size){
+        tid_array[tid] = tid; 
         int rate = rate_constants[tid];
         int rate_constants_col_index = tid % rate_constants_columns; 
         int row_index = (tid - rate_constants_col_index)/rate_constants_columns;
-        int reactant_num = number_of_reactants_[tid % rate_constants_columns]; //number of reactants of the reaction
-        int product_num = number_of_products_[tid % rate_constants_columns]; //number of products of the reaction 
-        int initial_reactant_ids_index = accumulated_n_reactants[tid % rate_constants_columns];
-        int initial_product_ids_index = accumulated_n_products[tid % rate_constants_columns];
-        int initial_yields_index = accumulated_n_products[tid % rate_constants_columns]; 
+        int reactant_num = number_of_reactants_[rate_constants_col_index]; //number of reactants of the reaction
+        int product_num = number_of_products_[rate_constants_col_index]; //number of products of the reaction 
+
+        int initial_reactant_ids_index = accumulated_n_reactants[rate_constants_col_index];
+        int initial_product_ids_index = accumulated_n_products[rate_constants_col_index];
+        int initial_yields_index = accumulated_n_products[rate_constants_col_index]; 
 
 
         for (int i_reactant = 0; i_reactant < reactant_num; i_reactant++){
@@ -108,6 +111,11 @@ namespace micm {
         size_t* d_accumulated_n_products; 
         size_t* d_product_ids_; 
         size_t* d_yields_; 
+        //debugging
+        int* d_tid_array;
+        int* tid_array; 
+        cudaMalloc(&d_tid_array, sizeof(int) * matrix_rows * rate_constants_columns); 
+        tid_array = (int*)malloc(sizeof(int) * matrix_rows * rate_constants_columns);
 
         //allocate device memory
         size_t rate_constants_bytes = sizeof(double) * (matrix_rows * rate_constants_columns); 
@@ -148,6 +156,7 @@ namespace micm {
         //kernel function call
 
         AddForcingTerms_kernel<<<num_block, block_size>>>(
+            d_tid_array,
             d_rate_constants, 
             d_state_variables, 
             d_forcing, 
@@ -163,7 +172,13 @@ namespace micm {
             d_yields_);
         cudaDeviceSynchronize(); 
         cudaMemcpy(forcing_data, d_forcing, state_forcing_bytes, cudaMemcpyDeviceToHost);
-
+
+        //debugging 
+        int tid_array_size = matrix_rows * rate_constants_columns; 
+        cudaMemcpy(tid_array, d_tid_array, sizeof(int)*tid_array_size, cudaMemcpyDeviceToHost );
+        for (int k = 0; k < tid_array_size; k++){
+            std::cout << tid_array[k]<<std::endl; 
+        }
         cudaFree(d_rate_constants); 
         cudaFree(d_state_variables); 
         cudaFree(d_forcing);

diff --git a/test/unit/process/test_cuda_process_set.cpp b/test/unit/process/test_cuda_process_set.cpp
@@ -47,6 +47,13 @@ TEST(ProcessSet, Constructor)
   rate_constants[1] = { 110.0, 120.0, 130.0 };
 
   micm::Matrix<double> forcing{ 2, 5, 1000.0 };
+
+  std::cout << "rate_constants data: "<<std::endl; 
+  double* rate_constants_data = rate_constants.AsVector().data(); 
+  for (int i = 0; i < rate_constants.AsVector().size(); i++){
+      std::cout <<rate_constants_data[i]<<std::endl; 
+  } 
+
 
   //debugging 
   std::cout<< "Before operation"<<std::endl; 
@@ -67,6 +74,12 @@ TEST(ProcessSet, Constructor)
   const double* yields = set.yields_vector().data();
   int yields_size = set.yields_vector().size(); 
 
+
+  // //print reactant_ids_ 
+  // std::cout <<
+  // for (int m = 0; m < reactant_ids_size; m++){
+  //   std::cout << reactant_ids[m]<<std::endl; 
+  // }
   micm::cuda::AddForcingTerms_kernelSetup(
     number_of_reactants,
     number_of_reactants_size,