add missing zmat cmake file, add mmc-trinity cuda support in cmake

fangq · Sep 8, 2023 · d38e7e5 · d38e7e5
1 parent 74a3035
commit d38e7e5
Show file tree

Hide file tree

Showing 3 changed files with 131 additions and 9 deletions.
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
@@ -17,14 +17,13 @@ project(mmc)
 
 cmake_policy(SET CMP0074 NEW)  # to make -DOPENCL_ROOT= work with newer CMake versions as well
 
-find_package(CUDA QUIET REQUIRED)
-
 find_package(OpenMP REQUIRED)
 find_package(OpenCL REQUIRED)
 
 add_subdirectory(zmat)
 
 option(BUILD_MEX "Build mex" ON)
+option(BUILD_CUDA "Build cuda" OFF)
 
 if(BUILD_PYTHON)
     add_subdirectory(pybind11)
@@ -36,6 +35,10 @@ if(BUILD_MEX)
     find_package(Matlab)
 endif()
 
+if(BUILD_CUDA)
+  find_package(CUDA QUIET REQUIRED)
+endif()
+
 # C Options
 set(CMAKE_CXX_FLAGS "-Wall -g -DMCX_EMBED_CL -fno-strict-aliasing -m64 -DMMC_USE_SSE -DHAVE_SSE2 -msse -msse2 -msse3 -mssse3 -msse4.1 -O3 -DUSE_OS_TIMER -DUSE_OPENCL -DMMC_XORSHIFT -fPIC ${OpenMP_CXX_FLAGS}")
 set(CMAKE_C_FLAGS "${CMAKE_CXX_FLAGS}")
@@ -54,6 +57,15 @@ add_custom_target(
     VERBATIM
 )
 
+if(BUILD_CUDA)
+  # NVC# NVCC Options
+  set(
+      CUDA_NVCC_FLAGS
+      ${CUDA_NVCC_FLAGS};
+      -g -lineinfo -DUSE_ATOMIC -DMCX_SAVE_DETECTORS -DMCX_DO_REFLECTION -DUSE_DMMC -DUSE_BLBADOUEL -Xcompiler -fPIC -use_fast_math -Xptxas -O3,-v -Xcompiler -Wall -Xcompiler "${OpenMP_CXX_FLAGS}"
+      )
+endif()
+
 # Add include directories
 include_directories(cjson ubj zmat zmat/easylzma)
 
@@ -107,6 +119,29 @@ target_link_libraries(
 add_dependencies(mmc clheader)
 add_dependencies(mmc-exe clheader)
 
+if(BUILD_CUDA)
+  # Create mmc-trinity binary
+  cuda_add_executable(
+      mmc-trinity
+      mmc_cu_host.cu
+      mmc_cu_host.h
+      mmc.c
+      )
+  # Link options
+  target_link_libraries(
+      mmc-trinity
+      mmc OpenMP::OpenMP_CXX
+      zmat
+      m
+      OpenCL::OpenCL
+      )
+
+  add_dependencies(mmc-trinity mmc)
+
+  set_target_properties(mmc-trinity
+        PROPERTIES OUTPUT_NAME mmc-trinity)
+endif()
+
 # Build mex file
 if(BUILD_MEX AND Matlab_FOUND)
     # Create mcx-matlab library

diff --git a/src/mmc_core.cl b/src/mmc_core.cl
@@ -1074,10 +1074,10 @@ __device__ void launchphoton(__constant MCXParam* gcfg, ray* r, __global float3*
 #endif
         float rx = rand_uniform01(ran);
         float ry = rand_uniform01(ran);
-        ry *= gcfg->srcparam1.w * MCX_MATHFUN(rsqrt)(gcfg->srcparam1.x * gcfg->srcparam1.x + gcfg->srcparam1.y * gcfg->srcparam1.y + gcfg->srcparam1.z * gcfg->srcparam1.z);
-        r->p0.x = gcfg->srcpos.x + rx * gcfg->srcparam1.x + ry * (gcfg->srcdir.y * gcfg->srcparam1.z - gcfg->srcdir.z * gcfg->srcparam1.y);
-        r->p0.y = gcfg->srcpos.y + rx * gcfg->srcparam1.y + ry * (gcfg->srcdir.z * gcfg->srcparam1.x - gcfg->srcdir.x * gcfg->srcparam1.z);
-        r->p0.z = gcfg->srcpos.z + rx * gcfg->srcparam1.z + ry * (gcfg->srcdir.x * gcfg->srcparam1.y - gcfg->srcdir.y * gcfg->srcparam1.x);
+        float tmp = gcfg->srcparam1.w * MCX_MATHFUN(rsqrt)(gcfg->srcparam1.x * gcfg->srcparam1.x + gcfg->srcparam1.y * gcfg->srcparam1.y + gcfg->srcparam1.z * gcfg->srcparam1.z);
+        r->p0.x = gcfg->srcpos.x + rx * gcfg->srcparam1.x + ry * tmp * (gcfg->srcdir.y * gcfg->srcparam1.z - gcfg->srcdir.z * gcfg->srcparam1.y);
+        r->p0.y = gcfg->srcpos.y + rx * gcfg->srcparam1.y + ry * tmp * (gcfg->srcdir.z * gcfg->srcparam1.x - gcfg->srcdir.x * gcfg->srcparam1.z);
+        r->p0.z = gcfg->srcpos.z + rx * gcfg->srcparam1.z + ry * tmp * (gcfg->srcdir.x * gcfg->srcparam1.y - gcfg->srcdir.y * gcfg->srcparam1.x);
 #if defined(__NVCC__) || defined(MCX_SRC_FOURIERX2D)
 #ifdef __NVCC__
 
@@ -1093,9 +1093,9 @@ __device__ void launchphoton(__constant MCXParam* gcfg, ray* r, __global float3*
             r->weight = (MCX_MATHFUN(cos)((gcfg->srcparam2.x * rx + gcfg->srcparam2.y * ry + gcfg->srcparam2.z) * TWO_PI) * (1.f - gcfg->srcparam2.w) + 1.f) * 0.5f; //between 0 and 1
 
 #endif
-        origin.x += (gcfg->srcparam1.x + v2.x) * 0.5f;
-        origin.y += (gcfg->srcparam1.y + v2.y) * 0.5f;
-        origin.z += (gcfg->srcparam1.z + v2.z) * 0.5f;
+        origin.x += (gcfg->srcparam1.x + tmp * (gcfg->srcdir.y * gcfg->srcparam1.z - gcfg->srcdir.z * gcfg->srcparam1.y)) * 0.5f;
+        origin.y += (gcfg->srcparam1.y + tmp * (gcfg->srcdir.z * gcfg->srcparam1.x - gcfg->srcdir.x * gcfg->srcparam1.z)) * 0.5f;
+        origin.z += (gcfg->srcparam1.z + tmp * (gcfg->srcdir.x * gcfg->srcparam1.y - gcfg->srcdir.y * gcfg->srcparam1.x)) * 0.5f;
 #endif
 #if defined(__NVCC__) || defined(MCX_SRC_DISK) || defined(MCX_SRC_GAUSSIAN) // uniform disk distribution or Gaussian-beam
 #ifdef __NVCC__

diff --git a/src/zmat/CMakeLists.txt b/src/zmat/CMakeLists.txt
@@ -0,0 +1,87 @@
+#################################################################
+#  CMake configure file for ZMat
+#  Qianqian Fang <q.fang at neu.edu>
+#  2020/05/23
+#################################################################
+
+cmake_minimum_required(VERSION 3.3)
+
+project(zmat)
+
+find_package(ZLIB REQUIRED)
+#find_package(Matlab)
+
+option(STATIC_LIB "Build static library" ON)
+
+# C Options
+set(CMAKE_C_FLAGS "-g -Wall -O3 -fPIC -DNO_BLOSC2 -DNO_ZLIB -DNO_ZSTD -D_LARGEFILE64_SOURCE=1")
+set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/../)
+
+# Add include directories
+#include_directories(./)
+include_directories(lz4)
+include_directories(miniz)
+include_directories(easylzma)
+include_directories(easylzma/pavlov)
+
+# Add all project units
+
+if(STATIC_LIB)
+  add_library(zmat STATIC
+    zmatlib.c
+    lz4/lz4.c
+    lz4/lz4hc.c
+    miniz/miniz.c
+    miniz/miniz.h
+    easylzma/compress.c
+    easylzma/decompress.c
+    easylzma/lzma_header.c
+    easylzma/lzip_header.c
+    easylzma/common_internal.c
+    easylzma/pavlov/LzmaEnc.c
+    easylzma/pavlov/LzmaDec.c
+    easylzma/pavlov/LzmaLib.c
+    easylzma/pavlov/LzFind.c
+    easylzma/pavlov/Bra.c
+    easylzma/pavlov/BraIA64.c
+    easylzma/pavlov/Alloc.c
+    easylzma/pavlov/7zCrc.c
+    )
+else()
+# Add all project units
+  add_library(zmat SHARED
+    zmatlib.c
+    lz4/lz4.c
+    lz4/lz4hc.c
+    miniz/miniz.c
+    miniz/miniz.h
+    easylzma/compress.c
+    easylzma/decompress.c
+    easylzma/lzma_header.c
+    easylzma/lzip_header.c
+    easylzma/common_internal.c
+    easylzma/pavlov/LzmaEnc.c
+    easylzma/pavlov/LzmaDec.c
+    easylzma/pavlov/LzmaLib.c
+    easylzma/pavlov/LzFind.c
+    easylzma/pavlov/Bra.c
+    easylzma/pavlov/BraIA64.c
+    easylzma/pavlov/Alloc.c
+    easylzma/pavlov/7zCrc.c
+    )
+endif()
+
+# Link options
+target_link_libraries(
+    zmat
+    ZLIB::ZLIB
+    )
+
+if(Matlab_FOUND)
+    matlab_add_mex(
+      NAME zipmat
+      SRC zmat.cpp
+      LINK_TO mex mx zmat
+    )
+endif()
+