<commit_before>/////////////////////////////////////////////////////////////
//
// Monte Carlo simulation software for light propagation in fluorescent turbid media,
// accelerated by GPU (graphic processing unit).
// The code is based on previous work by Alerstam et al and Wang et al,
// with the addition of a voxelized medium without symmetries and with an
// inhomogeneous distribution of absorbers and fluorescent marker
//
///////////////////////////////////////////////////////////////

/*	This file is part of CUDAMCFL.

    CUDAMCFL is free software: you can redistribute it and/or modify
    it under the terms of the GNU General Public License as published by
    the Free Software Foundation, either version 3 of the License, or
    (at your option) any later version.

    CUDAMCFL is distributed in the hope that it will be useful,
    but WITHOUT ANY WARRANTY; without even the implied warranty of
    MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
    GNU General Public License for more details.

    You should have received a copy of the GNU General Public License
    along with CUDAMCFL.  If not, see <http://www.gnu.org/licenses/>.
*/

#include "CUDAMCFL.h"
#include "cutil.h"
//#include <float.h> //for FLT_MAX
#include <limits.h>
#include <stdio.h>
#include "cuda_profiler_api.h"

__device__ __constant__ unsigned long long num_photons_dc[1];
__device__ __constant__ unsigned int n_layers_dc[1];
__device__ __constant__ unsigned int n_bulks_dc[1];
__device__ __constant__ unsigned int start_weight_dc[1];
__device__ __constant__ LayerStruct layers_dc[MAX_LAYERS];
__device__ __constant__ BulkStruct bulks_dc[MAX_LAYERS];
__device__ __constant__ DetStruct det_dc[1];
__device__ __constant__ IncStruct inclusion_dc[1];
__device__ __constant__ unsigned int ignoreAdetection_dc[1];
__device__ __constant__ unsigned int fhd_activated_dc[1];
__device__ __constant__ unsigned int bulk_method_dc[1];
__device__ __constant__ float xi_dc[1];
__device__ __constant__ float yi_dc[1];
__device__ __constant__ float zi_dc[1];
__device__ __constant__ float dir_dc[1];
__device__ __constant__ float esp_dc[1];
__device__ __constant__ unsigned int grid_size_dc[1];

#include "CUDAMCFLio.cu"
#include "CUDAMCFLmem.cu"
#include "CUDAMCFLrng.cu"
#include "CUDAMCFLtransport.cu"

// wrapper for device code - FHD Simulation
unsigned long long DoOneSimulation(SimulationStruct *simulation, unsigned long long *x,
                     unsigned int *a, double *tempfhd) {
  MemStruct DeviceMem;
  MemStruct HostMem;
  unsigned int threads_active_total = 1;
  unsigned int i, ii;

  // Output matrix size
  const int num_x = (int)(4 * (simulation->esp) * (float)simulation->grid_size);
  const int num_y = (int)(4 * (simulation->esp) * (float)simulation->grid_size);
  const int num_z = (int)((simulation->esp) * (float)simulation->grid_size);
  const int fhd_size = num_x * num_y * num_z;

  cudaError_t cudastat;
  clock_t time1, time2;

  // Start the clock
  time1 = clock();

  // x and a are already initialised in memory
  HostMem.x = x;
  HostMem.a = a;

  InitMemStructs(&HostMem, &DeviceMem, simulation);

  InitDCMem(simulation);

  dim3 dimBlock(NUM_THREADS_PER_BLOCK);
  dim3 dimGrid(NUM_BLOCKS);
  int blockSize;   // The launch configurator returned block size
  int minGridSize; // The minimum grid size needed to achieve the
                   // maximum occupancy for a full device launch
  //int gridSize;    // The actual grid size needed, based on input size

  cudaOccupancyMaxPotentialBlockSize( &minGridSize, &blockSize,
                                      MCd3D, 0, 0);
  printf ("Grid size: %i, Block size: %i \n\n", minGridSize, blockSize);


  LaunchPhoton_Global<<<dimGrid, dimBlock>>>(DeviceMem);
  //LaunchPhoton_Global<<<minGridSize, blockSize>>>(DeviceMem);
  CUDA_SAFE_CALL(cudaThreadSynchronize()); // Wait for all threads to finish
  cudastat = cudaGetLastError();           // Check if there was an error
  if (cudastat)
    printf("Error code=%i, %s.\n", cudastat, cudaGetErrorString(cudastat));

  i = 0;
  while (threads_active_total > 0) {
    i++;
    // run the kernel
    if (simulation->bulk_method == 1){
      MCd<<<dimGrid, dimBlock>>>(DeviceMem);
      //MCd<<<minGridSize, blockSize>>>(DeviceMem);

    }
    else if (simulation->bulk_method == 2) {

      MCd3D<<<dimGrid, dimBlock>>>(DeviceMem);
      //MCd3D<<<minGridSize, blockSize>>>(DeviceMem);
    }

    CUDA_SAFE_CALL(cudaThreadSynchronize()); // Wait for all threads to finish
    cudastat = cudaGetLastError();           // Check if there was an error
    if (cudastat)
      printf("Error code=%i, %s.\n", cudastat, cudaGetErrorString(cudastat));

    // Copy thread_active from device to host
    CUDA_SAFE_CALL(cudaMemcpy(HostMem.thread_active, DeviceMem.thread_active,
                              NUM_THREADS * sizeof(unsigned int),
                              cudaMemcpyDeviceToHost));
    threads_active_total = 0;
    for (ii = 0; ii < NUM_THREADS; ii++)
      threads_active_total += HostMem.thread_active[ii];

    CUDA_SAFE_CALL(cudaMemcpy(HostMem.num_terminated_photons,
                              DeviceMem.num_terminated_photons,
                              sizeof(unsigned long long), cudaMemcpyDeviceToHost));
    if (i == 100)
      printf("Estimated PHD simulation time: %.0f secs.\n\n",
             (double)(clock() - time1) / CLOCKS_PER_SEC *
                 (double)(simulation->number_of_photons /
                          *HostMem.num_terminated_photons));
//    if (fmod(i, 200u) == 0) printf("."); fflush(stdout);
    if (i % 200 == 0) printf("."); fflush(stdout);
    if (i % 10000 == 0)
      printf("\nRun %u, %llu photons simulated\n", i,
             *HostMem.num_terminated_photons);
  }

  CopyDeviceToHostMem(&HostMem, &DeviceMem, simulation);

  time2 = clock();

  printf("\nSimulation time: %.2f sec\n\n",
         (double)(time2 - time1) / CLOCKS_PER_SEC);

  printf("Writing excitation results...\n");
  Write_Simulation_Results(&HostMem, simulation, time2-time1);
  printf("PHD Simulation done!\n");
<commit_msg>Increased verbosity of PHD simulation progress.<commit_after>/////////////////////////////////////////////////////////////
//
// Monte Carlo simulation software for light propagation in fluorescent turbid media,
// accelerated by GPU (graphic processing unit).
// The code is based on previous work by Alerstam et al and Wang et al,
// with the addition of a voxelized medium without symmetries and with an
// inhomogeneous distribution of absorbers and fluorescent marker
//
///////////////////////////////////////////////////////////////

/*	This file is part of CUDAMCFL.

    CUDAMCFL is free software: you can redistribute it and/or modify
    it under the terms of the GNU General Public License as published by
    the Free Software Foundation, either version 3 of the License, or
    (at your option) any later version.

    CUDAMCFL is distributed in the hope that it will be useful,
    but WITHOUT ANY WARRANTY; without even the implied warranty of
    MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
    GNU General Public License for more details.

    You should have received a copy of the GNU General Public License
    along with CUDAMCFL.  If not, see <http://www.gnu.org/licenses/>.
*/

#include "CUDAMCFL.h"
#include "cutil.h"
//#include <float.h> //for FLT_MAX
#include <limits.h>
#include <stdio.h>
#include "cuda_profiler_api.h"

__device__ __constant__ unsigned long long num_photons_dc[1];
__device__ __constant__ unsigned int n_layers_dc[1];
__device__ __constant__ unsigned int n_bulks_dc[1];
__device__ __constant__ unsigned int start_weight_dc[1];
__device__ __constant__ LayerStruct layers_dc[MAX_LAYERS];
__device__ __constant__ BulkStruct bulks_dc[MAX_LAYERS];
__device__ __constant__ DetStruct det_dc[1];
__device__ __constant__ IncStruct inclusion_dc[1];
__device__ __constant__ unsigned int ignoreAdetection_dc[1];
__device__ __constant__ unsigned int fhd_activated_dc[1];
__device__ __constant__ unsigned int bulk_method_dc[1];
__device__ __constant__ float xi_dc[1];
__device__ __constant__ float yi_dc[1];
__device__ __constant__ float zi_dc[1];
__device__ __constant__ float dir_dc[1];
__device__ __constant__ float esp_dc[1];
__device__ __constant__ unsigned int grid_size_dc[1];

#include "CUDAMCFLio.cu"
#include "CUDAMCFLmem.cu"
#include "CUDAMCFLrng.cu"
#include "CUDAMCFLtransport.cu"

// wrapper for device code - FHD Simulation
unsigned long long DoOneSimulation(SimulationStruct *simulation, unsigned long long *x,
                     unsigned int *a, double *tempfhd) {
  MemStruct DeviceMem;
  MemStruct HostMem;
  unsigned int threads_active_total = 1;
  unsigned int i, ii;

  // Output matrix size
  const int num_x = (int)(4 * (simulation->esp) * (float)simulation->grid_size);
  const int num_y = (int)(4 * (simulation->esp) * (float)simulation->grid_size);
  const int num_z = (int)((simulation->esp) * (float)simulation->grid_size);
  const int fhd_size = num_x * num_y * num_z;

  cudaError_t cudastat;
  clock_t time1, time2;

  // Start the clock
  time1 = clock();

  // x and a are already initialised in memory
  HostMem.x = x;
  HostMem.a = a;

  InitMemStructs(&HostMem, &DeviceMem, simulation);

  InitDCMem(simulation);

  dim3 dimBlock(NUM_THREADS_PER_BLOCK);
  dim3 dimGrid(NUM_BLOCKS);
  int blockSize;   // The launch configurator returned block size
  int minGridSize; // The minimum grid size needed to achieve the
                   // maximum occupancy for a full device launch
  //int gridSize;    // The actual grid size needed, based on input size

  cudaOccupancyMaxPotentialBlockSize( &minGridSize, &blockSize,
                                      MCd3D, 0, 0);
  printf ("Grid size: %i, Block size: %i \n\n", minGridSize, blockSize);


  LaunchPhoton_Global<<<dimGrid, dimBlock>>>(DeviceMem);
  //LaunchPhoton_Global<<<minGridSize, blockSize>>>(DeviceMem);
  CUDA_SAFE_CALL(cudaThreadSynchronize()); // Wait for all threads to finish
  cudastat = cudaGetLastError();           // Check if there was an error
  if (cudastat)
    printf("Error code=%i, %s.\n", cudastat, cudaGetErrorString(cudastat));

  i = 0;
  while (threads_active_total > 0) {
    i++;
    // run the kernel
    if (simulation->bulk_method == 1){
      MCd<<<dimGrid, dimBlock>>>(DeviceMem);
      //MCd<<<minGridSize, blockSize>>>(DeviceMem);

    }
    else if (simulation->bulk_method == 2) {

      MCd3D<<<dimGrid, dimBlock>>>(DeviceMem);
      //MCd3D<<<minGridSize, blockSize>>>(DeviceMem);
    }

    CUDA_SAFE_CALL(cudaThreadSynchronize()); // Wait for all threads to finish
    cudastat = cudaGetLastError();           // Check if there was an error
    if (cudastat)
      printf("Error code=%i, %s.\n", cudastat, cudaGetErrorString(cudastat));

    // Copy thread_active from device to host
    CUDA_SAFE_CALL(cudaMemcpy(HostMem.thread_active, DeviceMem.thread_active,
                              NUM_THREADS * sizeof(unsigned int),
                              cudaMemcpyDeviceToHost));
    threads_active_total = 0;
    for (ii = 0; ii < NUM_THREADS; ii++)
      threads_active_total += HostMem.thread_active[ii];

    CUDA_SAFE_CALL(cudaMemcpy(HostMem.num_terminated_photons,
                              DeviceMem.num_terminated_photons,
                              sizeof(unsigned long long), cudaMemcpyDeviceToHost));
    if (i == 100)
      printf("Estimated PHD simulation time: %.0f secs.\n\n",
             (double)(clock() - time1) / CLOCKS_PER_SEC *
                 (double)(simulation->number_of_photons /
                          *HostMem.num_terminated_photons));
//    if (fmod(i, 200u) == 0) printf("."); fflush(stdout);
    if (i % 100 == 0) printf("."); fflush(stdout);
    if (i % 2000 == 0)
      printf("\nRun %u, %llu photons simulated\n", i,
             *HostMem.num_terminated_photons);
  }

  CopyDeviceToHostMem(&HostMem, &DeviceMem, simulation);

  time2 = clock();

  printf("\nSimulation time: %.2f sec\n\n",
         (double)(time2 - time1) / CLOCKS_PER_SEC);

  printf("Writing excitation results...\n");
  Write_Simulation_Results(&HostMem, simulation, time2-time1);
  printf("PHD Simulation done!\n");

{'fraction_non_alphanumeric': 0.12206303724928366, 'fraction_numerical': 0.009742120343839542, 'mean_word_length': 3.208060626937651, 'pattern_counts': {'":': 0, '<': 49, '<?xml version=': 0, '>': 67, 'https://': 0, 'lorem ipsum': 0, 'www.': 2, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 2, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '2524679', 'n_tokens_mistral': 3941, 'n_tokens_neox': 3776, 'n_words': 1107}