Unverified Commit f254b8e3 authored by Axel Kohlmeyer's avatar Axel Kohlmeyer Committed by GitHub
Browse files

Merge pull request #1226 from akohlmey/cmake-gpu-enhancements 

Enhancements for using CMake with the GPU package, improved compatibility with cmake 3.x versions, improved handling of shared library building.
parents 6e8c5375 ebacd5ca

cmake/CMakeLists.txt

+93 −52
Original line number Diff line number Diff line
@@ -11,6 +11,10 @@ get_filename_component(LAMMPS_LIB_SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/../lib 
get_filename_component(LAMMPS_LIB_BINARY_DIR ${CMAKE_BINARY_DIR}/lib ABSOLUTE)

get_filename_component(LAMMPS_DOC_DIR ${CMAKE_CURRENT_SOURCE_DIR}/../doc ABSOLUTE)



# by default, install into $HOME/.local (not /usr/local), so that no root access (and sudo!!) is needed

if (CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT)

  set(CMAKE_INSTALL_PREFIX "$ENV{HOME}/.local" CACHE PATH "default install path" FORCE )

endif()



# To avoid conflicts with the conventional Makefile build system, we build everything here

file(GLOB LIB_SOURCES ${LAMMPS_SOURCE_DIR}/[^.]*.cpp)
@@ -75,6 +79,7 @@ if(NOT CMAKE_BUILD_TYPE AND NOT CMAKE_CXX_FLAGS) 
  #release comes with -O3 by default

  set(CMAKE_BUILD_TYPE Release CACHE STRING "Choose the type of build, options are: None Debug Release RelWithDebInfo MinSizeRel." FORCE)

endif(NOT CMAKE_BUILD_TYPE AND NOT CMAKE_CXX_FLAGS)

string(TOUPPER "${CMAKE_BUILD_TYPE}" BTYPE)



# check for files auto-generated by make-based buildsystem

# this is fast, so check for it all the time
@@ -422,16 +427,30 @@ if(WITH_FFMPEG) 
  add_definitions(-DLAMMPS_FFMPEG)

endif()



if(BUILD_SHARED_LIBS)

  set(CONFIGURE_REQUEST_PIC "--with-pic")

  set(CMAKE_REQUEST_PIC "-DCMAKE_POSITION_INDEPENDENT_CODE=${CMAKE_POSITION_INDEPENDENT_CODE}")

  set(CUDA_REQUEST_PIC "-Xcompiler ${CMAKE_SHARED_LIBRARY_CXX_FLAGS}")

else()

  set(CONFIGURE_REQUEST_PIC)

  set(CMAKE_REQUEST_PIC)

  set(CUDA_REQUEST_PIC)

endif()





if(PKG_VORONOI)

  option(DOWNLOAD_VORO "Download voro++ (instead of using the system's one)" OFF)

  option(DOWNLOAD_VORO "Download and compile the Voro++ library instead of using an already installed one" OFF)

  if(DOWNLOAD_VORO)

    message(STATUS "Voro++ download requested - we will build our own")

    include(ExternalProject)



    if(BUILD_SHARED_LIBS)

        set(VORO_BUILD_OPTIONS "CFLAGS=-fPIC")

      set(VORO_BUILD_CFLAGS "${CMAKE_SHARED_LIBRARY_CXX_FLAGS} ${CMAKE_CXX_FLAGS} ${CMAKE_CXX_FLAGS_${BTYPE}}")

    else()

        set(VORO_BUILD_OPTIONS)

      set(VORO_BUILD_CFLAGS "${CMAKE_CXX_FLAGS} ${CMAKE_CXX_FLAGS_${BTYPE}}")

    endif()

    string(APPEND VORO_BUILD_CFLAGS ${CMAKE_CXX_FLAGS})

    set(VORO_BUILD_OPTIONS CXX=${CMAKE_CXX_COMPILER} CFLAGS=${VORO_BUILD_CFLAGS})



    ExternalProject_Add(voro_build

      URL http://math.lbl.gov/voro++/download/dir/voro++-0.4.6.tar.gz
@@ -445,7 +464,7 @@ if(PKG_VORONOI) 
  else()

    find_package(VORO)

    if(NOT VORO_FOUND)

      message(FATAL_ERROR "VORO not found, help CMake to find it by setting VORO_LIBRARY and VORO_INCLUDE_DIR, or set DOWNLOAD_VORO=ON to download it")

      message(FATAL_ERROR "Voro++ library not found. Help CMake to find it by setting VORO_LIBRARY and VORO_INCLUDE_DIR, or set DOWNLOAD_VORO=ON to download it")

    endif()

  endif()

  include_directories(${VORO_INCLUDE_DIRS})
@@ -453,18 +472,18 @@ if(PKG_VORONOI) 
endif()



if(PKG_LATTE)

  option(DOWNLOAD_LATTE "Download latte (instead of using the system's one)" OFF)

  option(DOWNLOAD_LATTE "Download the LATTE library instead of using an already installed one" OFF)

  if(DOWNLOAD_LATTE)

    if (CMAKE_VERSION VERSION_LESS "3.7") # due to SOURCE_SUBDIR

      message(FATAL_ERROR "For downlading LATTE you need at least cmake-3.7")

    endif()

    message(STATUS "LATTE not found - we will build our own")

    message(STATUS "LATTE download requested - we will build our own")

    include(ExternalProject)

    ExternalProject_Add(latte_build

      URL https://github.com/lanl/LATTE/archive/v1.2.1.tar.gz

      URL_MD5 85ac414fdada2d04619c8f936344df14

      SOURCE_SUBDIR cmake

      CMAKE_ARGS -DCMAKE_INSTALL_PREFIX=<INSTALL_DIR> -DCMAKE_POSITION_INDEPENDENT_CODE=${CMAKE_POSITION_INDEPENDENT_CODE}

      CMAKE_ARGS -DCMAKE_INSTALL_PREFIX=<INSTALL_DIR> ${CMAKE_REQUEST_PIC}

    )

    ExternalProject_get_property(latte_build INSTALL_DIR)

    set(LATTE_LIBRARIES ${INSTALL_DIR}/${CMAKE_INSTALL_LIBDIR}/liblatte.a)
@@ -472,7 +491,7 @@ if(PKG_LATTE) 
  else()

    find_package(LATTE)

    if(NOT LATTE_FOUND)

      message(FATAL_ERROR "LATTE not found, help CMake to find it by setting LATTE_LIBRARY, or set DOWNLOAD_LATTE=ON to download it")

      message(FATAL_ERROR "LATTE library not found, help CMake to find it by setting LATTE_LIBRARY, or set DOWNLOAD_LATTE=ON to download it")

    endif()

  endif()

  list(APPEND LAMMPS_LINK_LIBS ${LATTE_LIBRARIES} ${LAPACK_LIBRARIES})
@@ -480,19 +499,17 @@ endif() 


if(PKG_USER-SCAFACOS)

  find_package(GSL REQUIRED)

  option(DOWNLOAD_SCAFACOS "Download ScaFaCoS (instead of using the system's one)" OFF)

  option(DOWNLOAD_SCAFACOS "Download ScaFaCoS library instead of using an already installed one" OFF)

  if(DOWNLOAD_SCAFACOS)

    message(STATUS "ScaFaCoS download requested - we will build our own")

    include(ExternalProject)

    ExternalProject_Add(scafacos_build

      URL https://github.com/scafacos/scafacos/releases/download/v1.0.1/scafacos-1.0.1.tar.gz

      URL_MD5 bd46d74e3296bd8a444d731bb10c1738

      CONFIGURE_COMMAND <SOURCE_DIR>/configure --prefix=<INSTALL_DIR> 

                                               --disable-doc 

      CONFIGURE_COMMAND <SOURCE_DIR>/configure --prefix=<INSTALL_DIR> --disable-doc

                                               --enable-fcs-solvers=fmm,p2nfft,direct,ewald,p3m

                                               --with-internal-fftw

                                               --with-internal-pfft

                                               --with-internal-pnfft

                                               $<$<BOOL:${BUILD_SHARED_LIBS}>:--with-pic>                                               

                                               --with-internal-fftw --with-internal-pfft

                                               --with-internal-pnfft ${CONFIGURE_REQUEST_PIC}

                                               FC=${CMAKE_MPI_Fortran_COMPILER}

                                               CXX=${CMAKE_MPI_CXX_COMPILER}

                                               CC=${CMAKE_MPI_C_COMPILER}
@@ -537,15 +554,15 @@ if(PKG_USER-PLUMED) 
  validate_option(PLUMED_MODE PLUMED_MODE_VALUES)

  string(TOUPPER ${PLUMED_MODE} PLUMED_MODE)



  option(DOWNLOAD_PLUMED "Download Plumed (instead of using the system's one)" OFF)

  option(DOWNLOAD_PLUMED "Download Plumed package instead of using an already installed one" OFF)

  if(DOWNLOAD_PLUMED)

    message(STATUS "PLUMED download requested - we will build our own")

    include(ExternalProject)

    ExternalProject_Add(plumed_build

      URL https://github.com/plumed/plumed2/releases/download/v2.4.3/plumed-src-2.4.3.tgz

      URL_MD5 b1be7c48971627febc11c61b70767fc5

      BUILD_IN_SOURCE 1

      CONFIGURE_COMMAND <SOURCE_DIR>/configure --prefix=<INSTALL_DIR> 

                                               $<$<BOOL:${BUILD_SHARED_LIBS}>:--with-pic>                                                  )

      CONFIGURE_COMMAND <SOURCE_DIR>/configure --prefix=<INSTALL_DIR> ${CONFIGURE_REQUEST_PIC})

    ExternalProject_get_property(plumed_build INSTALL_DIR)

    set(PLUMED_INSTALL_DIR ${INSTALL_DIR})

    list(APPEND LAMMPS_DEPS plumed_build)
@@ -592,8 +609,9 @@ if(PKG_USER-NETCDF) 
endif()



if(PKG_USER-SMD)

  option(DOWNLOAD_EIGEN3 "Download Eigen3 (instead of using the system's one)" OFF)

  option(DOWNLOAD_EIGEN3 "Download Eigen3 instead of using an already installed one)" OFF)

  if(DOWNLOAD_EIGEN3)

    message(STATUS "Eigen3 download requested - we will build our own")

    include(ExternalProject)

    ExternalProject_Add(Eigen3_build

      URL http://bitbucket.org/eigen/eigen/get/3.3.4.tar.gz
@@ -633,8 +651,9 @@ if(PKG_USER-VTK) 
endif()



if(PKG_KIM)

  option(DOWNLOAD_KIM "Download kim-api (instead of using the system's one)" OFF)

  option(DOWNLOAD_KIM "Download KIM-API v1 from OpenKIM instead of using an already installed one)" OFF)

  if(DOWNLOAD_KIM)

    message(STATUS "KIM-API v1 download requested - we will build our own")

    include(ExternalProject)

    ExternalProject_Add(kim_build

      URL https://github.com/openkim/kim-api/archive/v1.9.5.tar.gz
@@ -649,7 +668,7 @@ if(PKG_KIM) 
  else()

    find_package(KIM)

    if(NOT KIM_FOUND)

      message(FATAL_ERROR "KIM not found, help CMake to find it by setting KIM_LIBRARY and KIM_INCLUDE_DIR, or set DOWNLOAD_KIM=ON to download it")

      message(FATAL_ERROR "KIM-API v1 not found, help CMake to find it by setting KIM_LIBRARY and KIM_INCLUDE_DIR, or set DOWNLOAD_KIM=ON to download it")

    endif()

  endif()

  list(APPEND LAMMPS_LINK_LIBS ${KIM_LIBRARIES})
@@ -662,12 +681,7 @@ if(PKG_MESSAGE) 
      ${LAMMPS_LIB_SOURCE_DIR}/message/cslib/[^.]*.c

      ${LAMMPS_LIB_SOURCE_DIR}/message/cslib/[^.]*.cpp)



  if(BUILD_SHARED_LIBS)

      add_library(cslib SHARED ${cslib_SOURCES})

  else()

  add_library(cslib STATIC ${cslib_SOURCES})

  endif()



  if(BUILD_MPI)

    target_compile_definitions(cslib PRIVATE -DMPI_YES)

    set_target_properties(cslib PROPERTIES OUTPUT_NAME "csmpi")
@@ -692,10 +706,10 @@ endif() 


if(PKG_MSCG)

  find_package(GSL REQUIRED)

  option(DOWNLOAD_MSCG "Download latte (instead of using the system's one)" OFF)

  option(DOWNLOAD_MSCG "Download MSCG library instead of using an already installed one)" OFF)

  if(DOWNLOAD_MSCG)

    if (CMAKE_VERSION VERSION_LESS "3.7") # due to SOURCE_SUBDIR

      message(FATAL_ERROR "For downlading LATTE you need at least cmake-3.7")

      message(FATAL_ERROR "For downlading MSCG you need at least cmake-3.7")

    endif()

    include(ExternalProject)

    if(NOT LAPACK_FOUND)
@@ -705,7 +719,7 @@ if(PKG_MSCG) 
      URL https://github.com/uchicago-voth/MSCG-release/archive/1.7.3.1.tar.gz

      URL_MD5 8c45e269ee13f60b303edd7823866a91

      SOURCE_SUBDIR src/CMake

      CMAKE_ARGS -DCMAKE_INSTALL_PREFIX=<INSTALL_DIR> -DCMAKE_POSITION_INDEPENDENT_CODE=${CMAKE_POSITION_INDEPENDENT_CODE} ${EXTRA_MSCG_OPTS}

      CMAKE_ARGS -DCMAKE_INSTALL_PREFIX=<INSTALL_DIR> ${CMAKE_REQUEST_PIC} ${EXTRA_MSCG_OPTS}

      BUILD_COMMAND make mscg INSTALL_COMMAND ""

      )

    ExternalProject_get_property(mscg_build BINARY_DIR)
@@ -753,7 +767,7 @@ set(MATH_LIBRARIES "m" CACHE STRING "math library") 
mark_as_advanced( MATH_LIBRARIES )

include(CheckLibraryExists)

if (CMAKE_VERSION VERSION_LESS "3.4")

  enable_language(C) # check_library_exists isn't supported without a c compiler before v3.4

  enable_language(C) # check_library_exists isn't supported without a C compiler before v3.4

endif()

# RB: disabled this check because it breaks with KOKKOS CUDA enabled

#foreach(FUNC sin cos)
@@ -1026,7 +1040,7 @@ if(PKG_USER-INTEL) 
    endif()



    if(CMAKE_CXX_COMPILER_VERSION VERSION_LESS 16)

        message(FATAL_ERROR "USER-INTEL is needed at least 2016 intel compiler, found ${CMAKE_CXX_COMPILER_VERSION}")

        message(FATAL_ERROR "USER-INTEL needs at least a 2016 intel compiler, found ${CMAKE_CXX_COMPILER_VERSION}")

    endif()



    if(NOT BUILD_OMP)
@@ -1126,11 +1140,11 @@ if(PKG_GPU) 
      find_package(CUDA REQUIRED)

      find_program(BIN2C bin2c)

      if(NOT BIN2C)

        message(FATAL_ERROR "Couldn't find bin2c, use -DBIN2C helping cmake to find it.")

        message(FATAL_ERROR "Could not find bin2c, use -DBIN2C=/path/to/bin2c to help cmake finding it.")

      endif()

      option(CUDPP_OPT "Enable CUDPP_OPT" ON)



      set(GPU_ARCH "sm_30" CACHE STRING "LAMMPS GPU CUDA SM architecture (e.g. sm_60)")

      set(GPU_ARCH "sm_30" CACHE STRING "LAMMPS GPU CUDA SM primary architecture (e.g. sm_60)")



      file(GLOB GPU_LIB_CU ${LAMMPS_LIB_SOURCE_DIR}/gpu/[^.]*.cu ${CMAKE_CURRENT_SOURCE_DIR}/gpu/[^.]*.cu)

      list(REMOVE_ITEM GPU_LIB_CU ${LAMMPS_LIB_SOURCE_DIR}/gpu/lal_pppm.cu)
@@ -1143,11 +1157,39 @@ if(PKG_GPU) 
        file(GLOB GPU_LIB_CUDPP_CU ${LAMMPS_LIB_SOURCE_DIR}/gpu/cudpp_mini/[^.]*.cu)

      endif()



      cuda_compile_cubin(GPU_GEN_OBJS ${GPU_LIB_CU} OPTIONS

                   -DUNIX -O3 -Xptxas -v --use_fast_math -DNV_KERNEL -DUCL_CUDADR -arch=${GPU_ARCH} -D_${GPU_PREC_SETTING})

      # build arch/gencode commands for nvcc based on CUDA toolkit version and use choice

      # --arch translates directly instead of JIT, so this should be for the preferred or most common architecture

      set(GPU_CUDA_GENCODE "-arch=${GPU_ARCH} ")

      # Fermi (GPU Arch 2.x) is supported by CUDA 3.2 to CUDA 8.0

      if((CUDA_VERSION VERSION_GREATER "3.1") AND (CUDA_VERSION VERSION_LESS "9.0"))

        string(APPEND GPU_CUDA_GENCODE "-gencode arch=compute_20,code=[sm_20,compute_20] ")

      endif()

      # Kepler (GPU Arch 3.x) is supported by CUDA 5 and later

      if(CUDA_VERSION VERSION_GREATER "4.9")

        string(APPEND GPU_CUDA_GENCODE "-gencode arch=compute_30,code=[sm_30,compute_30] -gencode arch=compute_35,code=[sm_35,compute_35] ")

      endif()

      # Maxwell (GPU Arch 5.x) is supported by CUDA 6 and later

      if(CUDA_VERSION VERSION_GREATER "5.9")

        string(APPEND GPU_CUDA_GENCODE "-gencode arch=compute_50,code=[sm_50,compute_50] -gencode arch=compute_52,code=[sm_52,compute_52] ")

      endif()

      # Pascal (GPU Arch 6.x) is supported by CUDA 8 and later

      if(CUDA_VERSION VERSION_GREATER "7.9")

        string(APPEND GPU_CUDA_GENCODE "-gencode arch=compute_60,code=[sm_60,compute_60] -gencode arch=compute_61,code=[sm_61,compute_61] ")

      endif()

      # Volta (GPU Arch 7.0) is supported by CUDA 9 and later

      if(CUDA_VERSION VERSION_GREATER "8.9")

        string(APPEND GPU_CUDA_GENCODE "-gencode arch=compute_70,code=[sm_70,compute_70] ")

      endif()

      # Turing (GPU Arch 7.5) is supported by CUDA 10 and later

      if(CUDA_VERSION VERSION_GREATER "9.9")

        string(APPEND GPU_CUDA_GENCODE "-gencode arch=compute_75,code=[sm_75,compute_75] ")

      endif()



      cuda_compile(GPU_OBJS ${GPU_LIB_CUDPP_CU} OPTIONS $<$<BOOL:${BUILD_SHARED_LIBS}>:-Xcompiler=-fPIC>

                   -DUNIX -O3 -Xptxas -v --use_fast_math -DUCL_CUDADR -arch=${GPU_ARCH} -D_${GPU_PREC_SETTING})

      cuda_compile_fatbin(GPU_GEN_OBJS ${GPU_LIB_CU} OPTIONS

              -DUNIX -O3 --use_fast_math -Wno-deprecated-gpu-targets -DNV_KERNEL -DUCL_CUDADR ${GPU_CUDA_GENCODE} -D_${GPU_PREC_SETTING})



      cuda_compile(GPU_OBJS ${GPU_LIB_CUDPP_CU} OPTIONS ${CUDA_REQUEST_PIC}

              -DUNIX -O3 --use_fast_math -Wno-deprecated-gpu-targets -DUCL_CUDADR ${GPU_CUDA_GENCODE} -D_${GPU_PREC_SETTING})



      foreach(CU_OBJ ${GPU_GEN_OBJS})

        get_filename_component(CU_NAME ${CU_OBJ} NAME_WE)
@@ -1416,7 +1458,6 @@ foreach(PKG ${DEFAULT_PACKAGES} ${ACCEL_PACKAGES} ${OTHER_PACKAGES}) 
  endif()

endforeach()



string(TOUPPER "${CMAKE_BUILD_TYPE}" BTYPE)

get_directory_property(CPPFLAGS DIRECTORY ${CMAKE_SOURCE_DIR} COMPILE_DEFINITIONS)

include(FeatureSummary)

feature_summary(DESCRIPTION "The following packages have been found:" WHAT PACKAGES_FOUND)

cmake/README.md

+1 −0
Original line number Diff line number Diff line
@@ -195,6 +195,7 @@ cmake -C ../cmake/presets/std_nolib.cmake -D PKG_GPU=on ../cmake 
  <td><code>CMAKE_INSTALL_PREFIX</code></td>

  <td>Install location where LAMMPS files will be copied to. In the Unix/Linux case with Makefiles this controls what `make install` will do.</td>

  <td>

   Default setting is <code>$HOME/.local</code>.

  </td>

</tr>

<tr>

doc/src/Build_cmake.txt

+1 −1
Original line number Diff line number Diff line
@@ -44,7 +44,7 @@ LAMMPS or need to re-compile LAMMPS repeatedly, installation of the 
ccache (= Compiler Cache) software may speed up compilation even more.



After compilation, you can optionally copy the LAMMPS executable and

library into your system folders (by default under /usr/local) with:

library into your system folders (by default under $HOME/.local) with:



make install    # optional, copy LAMMPS executable & library elsewhere :pre

doc/src/Build_extras.txt

+23 −9
Original line number Diff line number Diff line
@@ -87,22 +87,30 @@ which GPU hardware to build for. 
                      # value = double or mixed (default) or single

-D OCL_TUNE=value     # hardware choice for GPU_API=opencl

                      # generic (default) or intel (Intel CPU) or fermi, kepler, cypress (NVIDIA)

-D GPU_ARCH=value     # hardware choice for GPU_API=cuda

-D GPU_ARCH=value     # primary GPU hardware choice for GPU_API=cuda

                      # value = sm_XX, see below

                      # default is Cuda-compiler dependent, but typically sm_20

-D CUDPP_OPT=value    # optimization setting for GPU_API=cudea

-D CUDPP_OPT=value    # optimization setting for GPU_API=cuda

                      # enables CUDA Performance Primitives Optimizations

                      # yes (default) or no :pre



GPU_ARCH settings for different GPU hardware is as follows:



sm_20 for Fermi (C2050/C2070, deprecated as of CUDA 8.0) or GeForce GTX 580 or similar

sm_30 for Kepler (K10)

sm_35 for Kepler (K40) or GeForce GTX Titan or similar

sm_37 for Kepler (dual K80)

sm_50 for Maxwell

sm_60 for Pascal (P100)

sm_70 for Volta :ul

sm_20 or sm_21 for Fermi (supported by CUDA 3.2 until CUDA 7.5)

sm_30 or sm_35 or sm_37 for Kepler (supported since CUDA 5)

sm_50 or sm_52 for Maxwell (supported since CUDA 6)

sm_60 or sm_61 for Pascal (supported since CUDA 8)

sm_70 for Volta (supported since CUDA 9)

sm_75 for Turing (supported since CUDA 10) :ul



A more detailed list can be found, for example,

at "Wikipedia's CUDA article"_https://en.wikipedia.org/wiki/CUDA#GPUs_supported



CMake can detect which version of the CUDA toolkit is used and thus can

include support for [all] major GPU architectures supported by this toolkit.

Thus the GPU_ARCH setting is merely an optimization, to have code for

the preferred GPU architecture directly included rather than having to wait

for the JIT compiler of the CUDA driver to translate it.



[Traditional make]:
@@ -137,6 +145,11 @@ CUDA_ARCH = sm_XX, what GPU hardware you have, same as CMake GPU_ARCH above 
CUDA_PRECISION = precision (double, mixed, single)

EXTRAMAKE = which Makefile.lammps.* file to copy to Makefile.lammps :ul



The file Makefile.linux_multi is set up to include support for multiple

GPU architectures as supported by the CUDA toolkit in use. This is done

through using the "--gencode " flag, which can be used multiple times and

thus support all GPU architectures supported by your CUDA compiler.



If the library build is successful, 3 files should be created:

lib/gpu/libgpu.a, lib/gpu/nvc_get_devices, and

lib/gpu/Makefile.lammps.  The latter has settings that enable LAMMPS
@@ -150,6 +163,7 @@ re-build LAMMPS. This is because the compilation of files in the GPU 
package uses the library settings from the lib/gpu/Makefile.machine

used to build the GPU library.





:line



KIM package :h4,link(kim)

lib/gpu/lal_device.cu

+2 −1
Original line number Diff line number Diff line
@@ -11,7 +11,7 @@ 
//

//    begin                :

//    email                : brownw@ornl.gov

// ***************************************************************************/

// ***************************************************************************



#ifdef NV_KERNEL

#include "lal_preprocessor.h"
@@ -42,3 +42,4 @@ __kernel void kernel_info(__global int *info) { 
  info[13]=THREADS_PER_CHARGE;

  info[14]=BLOCK_ELLIPSE;

}

CRA Git | Maintained and supported by SUSTech CRA and CCSE