Initial commit of threaded CPU using hipSYCL

examples/standalone_cpp/eemumu/CMakeLists.txt

@@ -0,0 +1,40 @@
+cmake_minimum_required (VERSION 3.5)
+
+project(eemumu)
+
+set(CMAKE_EXPORT_COMPILE_COMMANDS 1)
+set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_CURRENT_LIST_DIR}/lib)
+set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_CURRENT_LIST_DIR}/SubProcesses/P1_Sigma_sm_epem_mupmum/)
+set(CONFIGURATION_NAME "Debug")
+set(CMAKE_CXX_STANDARD 17)
+set(CMAKE_CXX_STANDARD_REQUIRED ON)
+set(CMAKE_CXX_EXTENSIONS OFF)
+set(HIPSYCL_PLATFORM "cpu")
+
+set(HIPSYCL_INSTALL_LOCATION /opt/hipSYCL/)
+set(hipSYCL_DIR ${HIPSYCL_INSTALL_LOCATION}/lib/cmake)
+
+find_package(hipSYCL CONFIG REQUIRED)
+
+include_directories(
+    ${HIPSYCL_INSTALL_LOCATION}/include
+    ${HIPSYCL_INSTALL_LOCATION}/include/hipSYCL/contrib/
+    src 
+    SubProcesses/P1_Sigma_sm_epem_mupmum
+    )
+
+# Define the CXX sources
+set ( SOURCES
+    ${CMAKE_CURRENT_LIST_DIR}/src/read_slha.cc
+    ${CMAKE_CURRENT_LIST_DIR}/src/HelAmps_sm.cc
+    ${CMAKE_CURRENT_LIST_DIR}/src/Parameters_sm.cc
+    ${CMAKE_CURRENT_LIST_DIR}/src/rambo.cc
+    ${CMAKE_CURRENT_LIST_DIR}/SubProcesses/P1_Sigma_sm_epem_mupmum/CPPProcess.cc
+    ${CMAKE_CURRENT_LIST_DIR}/SubProcesses/P1_Sigma_sm_epem_mupmum/check_sa.cpp
+)
+
+# Compiler options
+add_definitions(-O0 -g )
+
+add_executable(check_sa.exe ${SOURCES})
+add_sycl_to_target(TARGET check_sa.exe SOURCES ${SOURCES})

examples/standalone_cpp/eemumu/SubProcesses/P1_Sigma_sm_epem_mupmum/CPPProcess.h

@@ -61,9 +61,10 @@ class CPPProcess {
 
   // print performance numbers
   void printPerformanceStats();
-
-private:
+
   std::vector<float> m_wavetimes;
+
+private:
   // Private functions to calculate the matrix element for all subprocesses
   // Calculate wavefunctions
   void calculate_wavefunctions(const int perm[], const int hel[]);

examples/standalone_cpp/eemumu/SubProcesses/check_sa.cpp

@@ -35,40 +35,52 @@ int main(int argc, char **argv) {
   if (verbose)
     std::cout << "num evts: " << numevts << std::endl;
 
-  cl::sycl::range<1> work_items{(unsigned long)numevts};
-
-  cl::sycl::queue q;
+  // Create a process object
+  CPPProcess process;
+
+  // Used to store timings
+  vector<float> t(numevts);
 
+  cl::sycl::range<1> work_items{(unsigned long)numevts};
+  cl::sycl::buffer<float> buff_t(t.data(), t.size());
+  cl::sycl::queue q; 
+
   q.submit([&](cl::sycl::handler& cgh){
 
-  // Create a process object
-  CPPProcess process;
+    auto access_t = buff_t.get_access<cl::sycl::access::mode::write>(cgh);
+
+    // Put here so that it captures the instance     
+    // Read param_card and set parameters
+    process.initProc("../../Cards/param_card.dat", verbose);
 
-  // Read param_card and set parameters
-  process.initProc("../../Cards/param_card.dat", verbose);
+    cgh.parallel_for<class my_kernel>(work_items,
+                                         [=] (cl::sycl::id<1> idx) {
+      // Create local (tread) copy (_lc) of the instance
+      CPPProcess process_lc = process; 
+      double energy = 1500;
+      double weight;
 
-  double energy = 1500;
-  double weight;
+      // Get phase space point
+      vector<double *> p = 
+      get_momenta(process_lc.ninitial, energy, process_lc.getMasses(), weight);
 
-  // Get phase space point
-  vector<double *> p =
-  get_momenta(process.ninitial, energy, process.getMasses(), weight);
+      // Set momenta for this event
+      process_lc.setMomenta(p);
 
-  // Set momenta for this event
-  process.setMomenta(p);
+      // Evaluate matrix element
+      process_lc.sigmaKin(false);
 
-  cgh.parallel_for<class vector_add>(work_items,
-                                         [=] (cl::sycl::id<1> tid) {
+      const double *matrix_elements = process_lc.getMatrixElements();
 
-    // Evaluate matrix element
-    process.sigmaKin(verbose);
+      // Store the timings back in the orginal instance
+      access_t[idx] = process_lc.m_wavetimes[0];
 
-    //const double *matrix_elements = process.getMatrixElements();
-
     });
-   process.printPerformanceStats();
+  });
+  q.wait();
+  process.m_wavetimes = t;
+  process.printPerformanceStats();
 
-    });
 /* 
    if (verbose) {
       cout << "Momenta:" << endl;

examples/standalone_cpp/eemumu/src/rambo.cc

@@ -19,9 +19,9 @@ double Random::ranmar() {
   ranu[iranmr] = uni;
   iranmr = iranmr - 1;
   jranmr = jranmr - 1;
-  if (iranmr == 0)
+  if (iranmr < 0) // fixed race condition
     iranmr = 97;
-  if (jranmr == 0)
+  if (jranmr < 0) // fixed race condition
     jranmr = 97;
   ranc = ranc - rancd;
   if (ranc < 0)