add cuda example

.gitignore

@@ -16,7 +16,7 @@ G4STORK/Build-xcode-serial
 geant4.10.00.p02-build-xcode
 
 CUDA_simpleTest/bin
-CUDA_Hadr04
+CUDA_Hadr04/bin
 
 # Geant / G4STORK build files
 G4STORK/Build/*

CUDA_Hadr04/src/CMakeLists.txt

@@ -0,0 +1,14 @@
+cmake_minimum_required (VERSION 2.8)
+
+option (USE_GPU "Use CUDA to run simulations on the GPU" ON)
+add_subdirectory(Hadr04)
+
+if (USE_GPU)
+	find_package (CUDA REQUIRED)
+
+	include_directories(/usr/local/cuda/include)
+	cuda_add_library (CUDA_G4NeutronHPVector SHARED CUDA_G4NeutronHPVector.h CUDA_G4NeutronHPVector.cu)
+
+	LIST(APPEND CUDA_NVCC_FLAGS --compiler-options -fno-strict-aliasing -lineinfo -use_fast_math -Xptxas -dlcm=cg)
+	LIST(APPEND CUDA_NVCC_FLAGS -gencode arch=compute_30,code=sm_30)
+endif()

CUDA_Hadr04/src/CUDA_G4NeutronHPVector.cu

@@ -0,0 +1,32 @@
+#include <cuda_runtime.h>
+#include <cuda.h>
+#include <stdio.h>
+#include "CUDA_G4NeutronHPVector.h"
+
+// CUDA kernel
+__global__
+void sumArrays(int* arr1, int* arr2, int* res, int n)
+{
+  int tid = blockIdx.x;
+  if (tid < n) 
+      res[tid] = arr1[tid] + arr2[tid];
+}
+
+void CUDA_sumArrays(int* arr1, int* arr2, int* res, int n) {
+    int *gpu_arr1, *gpu_arr2, *gpu_res;
+
+    cudaMalloc((void**)&gpu_arr1, n*sizeof(int));
+    cudaMalloc((void**)&gpu_arr2, n*sizeof(int));
+    cudaMalloc((void**)&gpu_res, n*sizeof(int));
+
+    cudaMemcpy(gpu_arr1, arr1, n*sizeof(int), cudaMemcpyHostToDevice);
+    cudaMemcpy(gpu_arr2, arr2, n*sizeof(int), cudaMemcpyHostToDevice);
+
+    sumArrays<<<n,1>>>(gpu_arr1, gpu_arr2, gpu_res, n);
+
+    cudaMemcpy(res, gpu_res, n*sizeof(int), cudaMemcpyDeviceToHost);
+
+    cudaFree(gpu_arr1);
+    cudaFree(gpu_arr2);
+    cudaFree(gpu_res);
+}

CUDA_Hadr04/src/CUDA_G4NeutronHPVector.h

@@ -0,0 +1 @@
+void CUDA_sumArrays(int* arr1, int* arr2, int* res, int n);

CUDA_Hadr04/src/Hadr04/.README

@@ -0,0 +1,127 @@
+//$Id$
+
+///\file "hadronic/Hadr04/.README"
+///\brief Example Hadr04 README page
+
+/*! \page ExampleHadr04 Example Hadr04
+
+   This example is focused on neutronHP physics, especially neutron transport,
+   including thermal scattering.
+   See A.R. Garcia, E. Mendoza, D. Cano-Ott presentation at G4 Hadronic group
+   meeting (04/2013) and note on G4NeutronHP package
+
+\section Hadr04_s1 MATERIALS AND GEOMETRY DEFINITION
+
+   It is a single box representing a 'quasi infinite' homogeneous medium.
+   Two parameters define the geometry :
+ 	- the material of the box,
+	- the (full) size of the box.
+
+   The default geometry (1 m3 of pressurized water) is built in 
+   DetectorConstruction, but the above parameters can be changed interactively 
+   via commands defined in DetectorMessenger.
+
+   A function, and its associated UI command, allows to build a material
+   directly from a single isotope.
+
+   To be identified by the ThermalScattering module, the elements composing a
+   material must have a specific name (see G4NeutronHPThermalScatteringNames.cc)
+   Examples of such materials are build in DetectorConstruction.
+
+\section Hadr04_s2 PHYSICS LIST
+
+   Only processes of neutronHP package are registered : neutronElastic 
+   (including thermalScattering), neutronInelastic, nCapture; nFission.
+   See class NeutronHPphysics. No other hadronic nor electromagnetic processes.
+
+   A command allows to select or not ThermalScattering model.
+
+   Several hadronic physics options are controlled by environment variables.
+   To trigger them, an envHadronic.csh has been added in this example.
+   One must select the options wished, and do
+\verbatim   
+        source envHadronic.csh  (or sh)
+\endverbatim	 
+
+\section Hadr04_s3 AN EVENT : THE PRIMARY GENERATOR
+
+   The primary kinematic is a single particle randomly shooted at the 
+   centre of the box. The type of the particle and its energy are set in 
+   PrimaryGeneratorAction (neutron 2 MeV), and can be changed via the G4 
+   build-in commands of ParticleGun class (see the macros provided with 
+   this example).
+
+\section Hadr04_s4 PHYSICS
+
+   All secondaries are killed in StackingAction. Therefore an event consists of
+   the transport of the primary neutron. Then one survey the thermal and non
+   thermal part of this parcours.
+
+
+\section Hadr04_s5 HISTOGRAMS
+
+   The test contains 7 built-in 1D histograms, which are managed by
+   G4AnalysisManager and its Messenger. The histos can be individually 
+   activated with the command :
+   /analysis/h1/set id nbBins  valMin valMax unit 
+   where unit is the desired unit for the histo (MeV or keV, etc..)
+   (see the macros xxxx.mac).
+
+   - 1   "incident neutron: nb of collisions above 1 eV"
+   - 2   "incident neutron: total track length above 1 eV"
+   - 3   "incident neutron: time of flight above 1 eV"
+   - 4   "incident neutron: nb of collisions below 1 eV"
+   - 5   "incident neutron: total track length below 1*eV"
+   - 6   "incident neutron: time of flight below 1 eV"
+   - 7   "incident neutron: energy distribution below 1*eV"   
+
+   The histograms are managed by the HistoManager class and its Messenger. 
+   The histos can be individually activated with the command :
+\verbatim
+/analysis/h1/set id nbBins  valMin valMax unit 
+\endverbatim   
+   where unit is the desired unit for the histo (MeV or keV, deg or mrad, etc..)
+
+   One can control the name of the histograms file with the command:
+\verbatim
+/analysis/setFileName  name  (default Hadr04)
+\endverbatim   
+   It is possible to choose the format of the histogram file : root (default),
+   xml, csv, by using namespace in HistoManager.hh
+
+   It is also possible to print selected histograms on an ascii file:
+\verbatim
+/analysis/h1/setAscii id
+\endverbatim   
+   All selected histos will be written on a file name.ascii (default Hadr04) 
+
+\section Hadr04_s6 VISUALIZATION
+
+   The Visualization Manager is set in the main().
+   The initialisation of the drawing is done via the commands
+   /vis/... in the macro vis.mac. To get visualisation:
+\verbatim
+> /control/execute vis.mac
+\endverbatim
+
+   The detector has a default view which is a longitudinal view of the box.
+   The tracks are drawn at the end of event, and erased at the end of run.
+
+\section Hadr04_s7 HOW TO START ?
+
+   Execute Hadr04 in 'batch' mode from macro files :
+\verbatim   
+% Hadr04   run01.mac
+\endverbatim
+
+   Execute Hadr04 in 'interactive mode' with visualization :
+\verbatim   
+% Hadr04
+Idle> control/execute vis.mac
+ ....
+ Idle> type your commands
+ ....
+ Idle> exit
+\endverbatim
+
+*/

CUDA_Hadr04/src/Hadr04/CMakeLists.txt

@@ -0,0 +1,73 @@
+#----------------------------------------------------------------------------
+# Setup the project
+cmake_minimum_required(VERSION 2.6 FATAL_ERROR)
+project(Hadr04)
+
+#----------------------------------------------------------------------------
+# Find Geant4 package, activating all available UI and Vis drivers by default
+# You can set WITH_GEANT4_UIVIS to OFF via the command line or ccmake/cmake-gui
+# to build a batch mode only executable
+#
+option(WITH_GEANT4_UIVIS "Build example with Geant4 UI and Vis drivers" ON)
+if(WITH_GEANT4_UIVIS)
+  find_package(Geant4 REQUIRED ui_all vis_all)
+else()
+  find_package(Geant4 REQUIRED)
+endif()
+
+#----------------------------------------------------------------------------
+# Setup Geant4 include directories and compile definitions
+#
+include(${Geant4_USE_FILE})
+
+#----------------------------------------------------------------------------
+# Locate sources and headers for this project
+#
+include_directories(${PROJECT_SOURCE_DIR}/include 
+                    ${Geant4_INCLUDE_DIR}
+                    ../)
+file(GLOB sources ${PROJECT_SOURCE_DIR}/src/*.cc)
+file(GLOB headers ${PROJECT_SOURCE_DIR}/include/*.hh)
+
+#----------------------------------------------------------------------------
+# Add the executable, and link it to the Geant4 libraries
+#
+add_executable(Hadr04 Hadr04.cc ${sources} ${headers})
+if (USE_GPU)
+    target_link_libraries(Hadr04 ${Geant4_LIBRARIES} CUDA_G4NeutronHPVector /usr/local/cuda/lib/libcudart_static.a)
+    add_definitions(-DUSE_GPU=1)
+    MESSAGE (STATUS "GPU Computation with CUDA is: ON")
+else()
+    MESSAGE (STATUS "GPU Computation with CUDA is: OFF")
+    add_definitions(-DUSE_GPU=0)
+    target_link_libraries(Hadr04 ${Geant4_LIBRARIES})
+endif()
+
+#----------------------------------------------------------------------------
+# Copy all scripts to the build directory, i.e. the directory in which we
+# build Hadr04. This is so that we can run the executable directly because it
+# relies on these scripts being in the current working directory.
+#
+set(Hadr04_SCRIPTS
+    debug.mac
+    envHadronic.csh
+    envHadronic.sh 
+    graphite.mac 
+    hadr04.in 
+    run01.mac 
+    vis.mac
+  )
+
+foreach(_script ${Hadr04_SCRIPTS})
+  configure_file(
+    ${PROJECT_SOURCE_DIR}/${_script}
+    ${PROJECT_BINARY_DIR}/${_script}
+    COPYONLY
+    )
+endforeach()
+
+#----------------------------------------------------------------------------
+# Install the executable to 'bin' directory under CMAKE_INSTALL_PREFIX
+#
+install(TARGETS Hadr04 DESTINATION bin)
+

CUDA_Hadr04/src/Hadr04/GNUmakefile

@@ -0,0 +1,23 @@
+# $Id: GNUmakefile 67909 2013-03-12 18:51:09Z vnivanch $
+# --------------------------------------------------------------
+# GNUmakefile for examples module.  Gabriele Cosmo, 06/04/98.
+# --------------------------------------------------------------
+
+name := Hadr04
+G4TARGET := $(name)
+G4EXLIB := true
+
+ifndef G4INSTALL
+  G4INSTALL = ../../../..
+endif
+
+.PHONY: all
+all: lib bin
+
+include $(G4INSTALL)/config/architecture.gmk
+
+include $(G4INSTALL)/config/binmake.gmk
+
+visclean:
+	rm -f g4*.prim g4*.eps g4*.wrl
+	rm -f .DAWN_*