diff --git a/cmake/CMakeLists.txt b/cmake/CMakeLists.txt index 7ce7ca23e16423be98be34edf51e4b4321931817..4df2ddd0752bae38a0bce6d7cd797dc0c1b18fb1 100644 --- a/cmake/CMakeLists.txt +++ b/cmake/CMakeLists.txt @@ -2,9 +2,9 @@ # CMake build system # This file is part of LAMMPS # Created by Christoph Junghans and Richard Berger -cmake_minimum_required(VERSION 3.1) +cmake_minimum_required(VERSION 2.8.12) -project(lammps LANGUAGES CXX) +project(lammps CXX) set(SOVERSION 0) set(LAMMPS_SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/../src) set(LAMMPS_LIB_SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/../lib) @@ -129,7 +129,7 @@ foreach(PKG ${ACCEL_PACKAGES} ${OTHER_PACKAGES}) endforeach() macro(pkg_depends PKG1 PKG2) - if(PKG_${PKG1} AND NOT PKG_${PKG2}) + if(PKG_${PKG1} AND NOT (PKG_${PKG2} OR BUILD_${PKG2})) message(FATAL_ERROR "${PKG1} package needs LAMMPS to be build with ${PKG2}") endif() endmacro() @@ -137,23 +137,27 @@ endmacro() pkg_depends(MPIIO MPI) pkg_depends(QEQ MANYBODY) pkg_depends(USER-ATC MANYBODY) -pkg_depends(USER-H5MD MPI) pkg_depends(USER-LB MPI) pkg_depends(USER-MISC MANYBODY) pkg_depends(USER-PHONON KSPACE) +pkg_depends(CORESHELL KSPACE) ###################################################### # packages with special compiler needs or external libs ###################################################### if(PKG_REAX OR PKG_MEAM OR PKG_USER-QUIP OR PKG_USER-QMMM OR PKG_LATTE) enable_language(Fortran) + list(APPEND LAMMPS_LINK_LIBS ${CMAKE_Fortran_IMPLICIT_LINK_LIBRARIES}) endif() -if(PKG_MEAM) +if(PKG_MEAM OR PKG_USER-H5MD OR PKG_USER-QMMM) enable_language(C) endif() -if(PKG_KOKKOS OR PKG_MSCG) +if(PKG_MSCG) + if (CMAKE_VERSION VERSION_LESS "3.1") + message(FATAL_ERROR "For the MSCG package you need at least cmake-3.1") + endif() # starting with CMake 3.1 this is all you have to do to enforce C++11 set(CMAKE_CXX_STANDARD 11) # C++11... set(CMAKE_CXX_STANDARD_REQUIRED ON) #...is required... @@ -312,7 +316,7 @@ if(PKG_LATTE) message(FATAL_ERROR "LATTE 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} ${CMAKE_Fortran_IMPLICIT_LINK_LIBRARIES}) + list(APPEND LAMMPS_LINK_LIBS ${LATTE_LIBRARIES} ${LAPACK_LIBRARIES}) endif() if(PKG_USER-MOLFILE) @@ -355,14 +359,14 @@ endif() if(PKG_USER-QUIP) find_package(QUIP REQUIRED) - list(APPEND LAMMPS_LINK_LIBS ${QUIP_LIBRARIES} ${LAPACK_LIBRARIES} ${CMAKE_Fortran_IMPLICIT_LINK_LIBRARIES}) + list(APPEND LAMMPS_LINK_LIBS ${QUIP_LIBRARIES} ${LAPACK_LIBRARIES}) endif() if(PKG_USER-QMMM) message(WARNING "Building QMMM with CMake is still experimental") find_package(QE REQUIRED) include_directories(${QE_INCLUDE_DIRS}) - list(APPEND LAMMPS_LINK_LIBS ${QE_LIBRARIES} ${CMAKE_Fortran_IMPLICIT_LINK_LIBRARIES}) + list(APPEND LAMMPS_LINK_LIBS ${QE_LIBRARIES}) endif() if(PKG_USER-VTK) @@ -657,6 +661,9 @@ if(PKG_USER-INTEL) endif() if(PKG_GPU) + if (CMAKE_VERSION VERSION_LESS "3.1") + message(FATAL_ERROR "For the GPU package you need at least cmake-3.1") + endif() set(GPU_SOURCES_DIR ${LAMMPS_SOURCE_DIR}/GPU) set(GPU_SOURCES ${GPU_SOURCES_DIR}/gpu_extra.h ${GPU_SOURCES_DIR}/fix_gpu.h @@ -753,7 +760,7 @@ if(PKG_GPU) add_library(gpu STATIC ${GPU_LIB_SOURCES}) target_link_libraries(gpu ${OpenCL_LIBRARIES}) target_include_directories(gpu PRIVATE ${CMAKE_CURRENT_BINARY_DIR}/gpu ${OpenCL_INCLUDE_DIRS}) - target_compile_definitions(gpu PRIVATE -D_${GPU_PREC} -DMPI_GERYON -DUCL_NO_EXIT) + target_compile_definitions(gpu PRIVATE -D_${GPU_PREC} -D${OCL_TUNE}_OCL -DMPI_GERYON -DUCL_NO_EXIT) target_compile_definitions(gpu PRIVATE -DUSE_OPENCL) list(APPEND LAMMPS_LINK_LIBS gpu) diff --git a/cmake/README.md b/cmake/README.md index cc67cceb52ec7b40327af88b52709eda7a6c1416..67d5edc9f5024887f9db8d1ecdc35b656363c57a 100644 --- a/cmake/README.md +++ b/cmake/README.md @@ -9,7 +9,7 @@ cmake /path/to/lammps/source/cmake To enable package, e.g. GPU do ``` -cmake /path/to/lammps/source/cmake -DENABLE_GPU=ON +cmake /path/to/lammps/source/cmake -DPKG_GPU=ON ``` cmake has many many options, do get an overview use the curses-based cmake interface, ccmake: diff --git a/lib/gpu/Makefile.linux_multi b/lib/gpu/Makefile.linux_multi new file mode 100644 index 0000000000000000000000000000000000000000..02b405ac35114402cbd5609b635748fd2df37e55 --- /dev/null +++ b/lib/gpu/Makefile.linux_multi @@ -0,0 +1,53 @@ +# /* ---------------------------------------------------------------------- +# Generic Linux Makefile for CUDA +# - Change CUDA_ARCH for your GPU +# ------------------------------------------------------------------------- */ + +# which file will be copied to Makefile.lammps + +EXTRAMAKE = Makefile.lammps.standard + +ifeq ($(CUDA_HOME),) +CUDA_HOME = /usr/local/cuda +endif + +NVCC = nvcc + +# Kepler CUDA +#CUDA_ARCH = -arch=sm_35 +# newer CUDA +#CUDA_ARCH = -arch=sm_13 +# older CUDA +#CUDA_ARCH = -arch=sm_10 -DCUDA_PRE_THREE + +CUDA_ARCH = -gencode arch=compute_60,code=sm_60 -gencode arch=compute_61,code=sm_61 + +# this setting should match LAMMPS Makefile +# one of LAMMPS_SMALLBIG (default), LAMMPS_BIGBIG and LAMMPS_SMALLSMALL + +LMP_INC = -DLAMMPS_SMALLBIG + +# precision for GPU calculations +# -D_SINGLE_SINGLE # Single precision for all calculations +# -D_DOUBLE_DOUBLE # Double precision for all calculations +# -D_SINGLE_DOUBLE # Accumulation of forces, etc. in double + +CUDA_PRECISION = -D_SINGLE_DOUBLE + +CUDA_INCLUDE = -I$(CUDA_HOME)/include +CUDA_LIB = -L$(CUDA_HOME)/lib64 +CUDA_OPTS = -DUNIX -O3 -Xptxas -v --use_fast_math $(LMP_INC) -Xcompiler "-fPIC -std=c++98" + +CUDR_CPP = mpicxx -DMPI_GERYON -DUCL_NO_EXIT -DMPICH_IGNORE_CXX_SEEK -DOMPI_SKIP_MPICXX=1 -fPIC +CUDR_OPTS = -O2 $(LMP_INC) # -xHost -no-prec-div -ansi-alias + +BIN_DIR = ./ +OBJ_DIR = ./ +LIB_DIR = ./ +AR = ar +BSH = /bin/sh + +CUDPP_OPT = -DUSE_CUDPP -Icudpp_mini + +include Nvidia.makefile_multi + diff --git a/lib/gpu/Nvidia.makefile b/lib/gpu/Nvidia.makefile index 5f692cf66c8bbe9834ca36f81786fef1ad2e9ae5..21663ad85cb161c5acf97ee0463f295beaf58c17 100644 --- a/lib/gpu/Nvidia.makefile +++ b/lib/gpu/Nvidia.makefile @@ -77,7 +77,9 @@ OBJS = $(OBJ_DIR)/lal_atom.o $(OBJ_DIR)/lal_ans.o \ $(OBJ_DIR)/lal_coul_debye.o $(OBJ_DIR)/lal_coul_debye_ext.o \ $(OBJ_DIR)/lal_zbl.o $(OBJ_DIR)/lal_zbl_ext.o \ $(OBJ_DIR)/lal_lj_cubic.o $(OBJ_DIR)/lal_lj_cubic_ext.o \ - $(OBJ_DIR)/lal_ufm.o $(OBJ_DIR)/lal_ufm_ext.o + $(OBJ_DIR)/lal_ufm.o $(OBJ_DIR)/lal_ufm_ext.o \ + $(OBJ_DIR)/lal_dipole_long_lj.o $(OBJ_DIR)/lal_dipole_long_lj_ext.o \ + $(OBJ_DIR)/lal_lj_expand_coul_long.o $(OBJ_DIR)/lal_lj_expand_coul_long_ext.o CBNS = $(OBJ_DIR)/device.cubin $(OBJ_DIR)/device_cubin.h \ $(OBJ_DIR)/atom.cubin $(OBJ_DIR)/atom_cubin.h \ @@ -133,7 +135,9 @@ CBNS = $(OBJ_DIR)/device.cubin $(OBJ_DIR)/device_cubin.h \ $(OBJ_DIR)/coul_debye.cubin $(OBJ_DIR)/coul_debye_cubin.h \ $(OBJ_DIR)/zbl.cubin $(OBJ_DIR)/zbl_cubin.h \ $(OBJ_DIR)/lj_cubic.cubin $(OBJ_DIR)/lj_cubic_cubin.h \ - $(OBJ_DIR)/ufm.cubin $(OBJ_DIR)/ufm_cubin.h + $(OBJ_DIR)/ufm.cubin $(OBJ_DIR)/ufm_cubin.h \ + $(OBJ_DIR)/dipole_long_lj.cubin $(OBJ_DIR)/dipole_long_lj_cubin.h \ + $(OBJ_DIR)/lj_expand_coul_long.cubin $(OBJ_DIR)/lj_expand_coul_long_cubin.h all: $(OBJ_DIR) $(GPU_LIB) $(EXECS) @@ -809,6 +813,30 @@ $(OBJ_DIR)/lal_lj_cubic.o: $(ALL_H) lal_lj_cubic.h lal_lj_cubic.cpp $(OBJ_DIR)/l $(OBJ_DIR)/lal_lj_cubic_ext.o: $(ALL_H) lal_lj_cubic.h lal_lj_cubic_ext.cpp lal_base_atomic.h $(CUDR) -o $@ -c lal_lj_cubic_ext.cpp -I$(OBJ_DIR) +$(OBJ_DIR)/dipole_long_lj.cubin: lal_dipole_long_lj.cu lal_precision.h lal_preprocessor.h + $(CUDA) --cubin -DNV_KERNEL -o $@ lal_dipole_long_lj.cu + +$(OBJ_DIR)/dipole_long_lj_cubin.h: $(OBJ_DIR)/dipole_long_lj.cubin $(OBJ_DIR)/dipole_long_lj.cubin + $(BIN2C) -c -n dipole_long_lj $(OBJ_DIR)/dipole_long_lj.cubin > $(OBJ_DIR)/dipole_long_lj_cubin.h + +$(OBJ_DIR)/lal_dipole_long_lj.o: $(ALL_H) lal_dipole_long_lj.h lal_dipole_long_lj.cpp $(OBJ_DIR)/dipole_long_lj_cubin.h $(OBJ_DIR)/lal_base_dipole.o + $(CUDR) -o $@ -c lal_dipole_long_lj.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_dipole_long_lj_ext.o: $(ALL_H) lal_dipole_long_lj.h lal_dipole_long_lj_ext.cpp lal_base_dipole.h + $(CUDR) -o $@ -c lal_dipole_long_lj_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lj_expand_coul_long.cubin: lal_lj_expand_coul_long.cu lal_precision.h lal_preprocessor.h + $(CUDA) --cubin -DNV_KERNEL -o $@ lal_lj_expand_coul_long.cu + +$(OBJ_DIR)/lj_expand_coul_long_cubin.h: $(OBJ_DIR)/lj_expand_coul_long.cubin $(OBJ_DIR)/lj_expand_coul_long.cubin + $(BIN2C) -c -n lj_expand_coul_long $(OBJ_DIR)/lj_expand_coul_long.cubin > $(OBJ_DIR)/lj_expand_coul_long_cubin.h + +$(OBJ_DIR)/lal_lj_expand_coul_long.o: $(ALL_H) lal_lj_expand_coul_long.h lal_lj_expand_coul_long.cpp $(OBJ_DIR)/lj_expand_coul_long_cubin.h $(OBJ_DIR)/lal_base_charge.o + $(CUDR) -o $@ -c lal_lj_expand_coul_long.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_lj_expand_coul_long_ext.o: $(ALL_H) lal_lj_expand_coul_long.h lal_lj_expand_coul_long_ext.cpp lal_base_charge.h + $(CUDR) -o $@ -c lal_lj_expand_coul_long_ext.cpp -I$(OBJ_DIR) + $(BIN_DIR)/nvc_get_devices: ./geryon/ucl_get_devices.cpp $(NVD_H) $(CUDR) -o $@ ./geryon/ucl_get_devices.cpp -DUCL_CUDADR $(CUDA_LIB) -lcuda diff --git a/lib/gpu/Nvidia.makefile_multi b/lib/gpu/Nvidia.makefile_multi new file mode 100644 index 0000000000000000000000000000000000000000..5fb35cce3c30cdb55f5ca5ff76dcb0569c8d3920 --- /dev/null +++ b/lib/gpu/Nvidia.makefile_multi @@ -0,0 +1,854 @@ +CUDA = $(NVCC) $(CUDA_INCLUDE) $(CUDA_OPTS) -Icudpp_mini $(CUDA_ARCH) \ + $(CUDA_PRECISION) +CUDR = $(CUDR_CPP) $(CUDR_OPTS) $(CUDA_PRECISION) $(CUDA_INCLUDE) \ + $(CUDPP_OPT) +CUDA_LINK = $(CUDA_LIB) -lcudart +BIN2C = $(CUDA_HOME)/bin/bin2c + +GPU_LIB = $(LIB_DIR)/libgpu.a + +# Headers for Geryon +UCL_H = $(wildcard ./geryon/ucl*.h) +NVC_H = $(wildcard ./geryon/nvc*.h) $(UCL_H) +NVD_H = $(wildcard ./geryon/nvd*.h) $(UCL_H) lal_preprocessor.h +# Headers for Pair Stuff +PAIR_H = lal_atom.h lal_answer.h lal_neighbor_shared.h \ + lal_neighbor.h lal_precision.h lal_device.h \ + lal_balance.h lal_pppm.h + +ALL_H = $(NVD_H) $(PAIR_H) + +EXECS = $(BIN_DIR)/nvc_get_devices +ifdef CUDPP_OPT +CUDPP = $(OBJ_DIR)/cudpp.o $(OBJ_DIR)/cudpp_plan.o \ + $(OBJ_DIR)/cudpp_maximal_launch.o $(OBJ_DIR)/cudpp_plan_manager.o \ + $(OBJ_DIR)/radixsort_app.cu_o $(OBJ_DIR)/scan_app.cu_o +endif +OBJS = $(OBJ_DIR)/lal_atom.o $(OBJ_DIR)/lal_ans.o \ + $(OBJ_DIR)/lal_neighbor.o $(OBJ_DIR)/lal_neighbor_shared.o \ + $(OBJ_DIR)/lal_device.o $(OBJ_DIR)/lal_base_atomic.o \ + $(OBJ_DIR)/lal_base_charge.o $(OBJ_DIR)/lal_base_ellipsoid.o \ + $(OBJ_DIR)/lal_base_dipole.o $(OBJ_DIR)/lal_base_three.o \ + $(OBJ_DIR)/lal_base_dpd.o \ + $(OBJ_DIR)/lal_pppm.o $(OBJ_DIR)/lal_pppm_ext.o \ + $(OBJ_DIR)/lal_gayberne.o $(OBJ_DIR)/lal_gayberne_ext.o \ + $(OBJ_DIR)/lal_re_squared.o $(OBJ_DIR)/lal_re_squared_ext.o \ + $(OBJ_DIR)/lal_lj.o $(OBJ_DIR)/lal_lj_ext.o \ + $(OBJ_DIR)/lal_lj96.o $(OBJ_DIR)/lal_lj96_ext.o \ + $(OBJ_DIR)/lal_lj_expand.o $(OBJ_DIR)/lal_lj_expand_ext.o \ + $(OBJ_DIR)/lal_lj_coul.o $(OBJ_DIR)/lal_lj_coul_ext.o \ + $(OBJ_DIR)/lal_lj_coul_long.o $(OBJ_DIR)/lal_lj_coul_long_ext.o \ + $(OBJ_DIR)/lal_lj_dsf.o $(OBJ_DIR)/lal_lj_dsf_ext.o \ + $(OBJ_DIR)/lal_lj_class2_long.o $(OBJ_DIR)/lal_lj_class2_long_ext.o \ + $(OBJ_DIR)/lal_coul_long.o $(OBJ_DIR)/lal_coul_long_ext.o \ + $(OBJ_DIR)/lal_morse.o $(OBJ_DIR)/lal_morse_ext.o \ + $(OBJ_DIR)/lal_charmm_long.o $(OBJ_DIR)/lal_charmm_long_ext.o \ + $(OBJ_DIR)/lal_lj_sdk.o $(OBJ_DIR)/lal_lj_sdk_ext.o \ + $(OBJ_DIR)/lal_lj_sdk_long.o $(OBJ_DIR)/lal_lj_sdk_long_ext.o \ + $(OBJ_DIR)/lal_eam.o $(OBJ_DIR)/lal_eam_ext.o \ + $(OBJ_DIR)/lal_eam_fs_ext.o $(OBJ_DIR)/lal_eam_alloy_ext.o \ + $(OBJ_DIR)/lal_buck.o $(OBJ_DIR)/lal_buck_ext.o \ + $(OBJ_DIR)/lal_buck_coul.o $(OBJ_DIR)/lal_buck_coul_ext.o \ + $(OBJ_DIR)/lal_buck_coul_long.o $(OBJ_DIR)/lal_buck_coul_long_ext.o \ + $(OBJ_DIR)/lal_table.o $(OBJ_DIR)/lal_table_ext.o \ + $(OBJ_DIR)/lal_yukawa.o $(OBJ_DIR)/lal_yukawa_ext.o \ + $(OBJ_DIR)/lal_born.o $(OBJ_DIR)/lal_born_ext.o \ + $(OBJ_DIR)/lal_born_coul_wolf.o $(OBJ_DIR)/lal_born_coul_wolf_ext.o \ + $(OBJ_DIR)/lal_born_coul_long.o $(OBJ_DIR)/lal_born_coul_long_ext.o \ + $(OBJ_DIR)/lal_dipole_lj.o $(OBJ_DIR)/lal_dipole_lj_ext.o \ + $(OBJ_DIR)/lal_dipole_lj_sf.o $(OBJ_DIR)/lal_dipole_lj_sf_ext.o \ + $(OBJ_DIR)/lal_colloid.o $(OBJ_DIR)/lal_colloid_ext.o \ + $(OBJ_DIR)/lal_gauss.o $(OBJ_DIR)/lal_gauss_ext.o \ + $(OBJ_DIR)/lal_yukawa_colloid.o $(OBJ_DIR)/lal_yukawa_colloid_ext.o \ + $(OBJ_DIR)/lal_lj_coul_debye.o $(OBJ_DIR)/lal_lj_coul_debye_ext.o \ + $(OBJ_DIR)/lal_coul_dsf.o $(OBJ_DIR)/lal_coul_dsf_ext.o \ + $(OBJ_DIR)/lal_sw.o $(OBJ_DIR)/lal_sw_ext.o \ + $(OBJ_DIR)/lal_vashishta.o $(OBJ_DIR)/lal_vashishta_ext.o \ + $(OBJ_DIR)/lal_beck.o $(OBJ_DIR)/lal_beck_ext.o \ + $(OBJ_DIR)/lal_mie.o $(OBJ_DIR)/lal_mie_ext.o \ + $(OBJ_DIR)/lal_soft.o $(OBJ_DIR)/lal_soft_ext.o \ + $(OBJ_DIR)/lal_lj_coul_msm.o $(OBJ_DIR)/lal_lj_coul_msm_ext.o \ + $(OBJ_DIR)/lal_lj_gromacs.o $(OBJ_DIR)/lal_lj_gromacs_ext.o \ + $(OBJ_DIR)/lal_dpd.o $(OBJ_DIR)/lal_dpd_ext.o \ + $(OBJ_DIR)/lal_tersoff.o $(OBJ_DIR)/lal_tersoff_ext.o \ + $(OBJ_DIR)/lal_tersoff_zbl.o $(OBJ_DIR)/lal_tersoff_zbl_ext.o \ + $(OBJ_DIR)/lal_tersoff_mod.o $(OBJ_DIR)/lal_tersoff_mod_ext.o \ + $(OBJ_DIR)/lal_coul.o $(OBJ_DIR)/lal_coul_ext.o \ + $(OBJ_DIR)/lal_coul_debye.o $(OBJ_DIR)/lal_coul_debye_ext.o \ + $(OBJ_DIR)/lal_zbl.o $(OBJ_DIR)/lal_zbl_ext.o \ + $(OBJ_DIR)/lal_lj_cubic.o $(OBJ_DIR)/lal_lj_cubic_ext.o \ + $(OBJ_DIR)/lal_ufm.o $(OBJ_DIR)/lal_ufm_ext.o \ + $(OBJ_DIR)/lal_dipole_long_lj.o $(OBJ_DIR)/lal_dipole_long_lj_ext.o \ + $(OBJ_DIR)/lal_lj_expand_coul_long.o $(OBJ_DIR)/lal_lj_expand_coul_long_ext.o + +CBNS = $(OBJ_DIR)/device.cubin $(OBJ_DIR)/device_cubin.h \ + $(OBJ_DIR)/atom.cubin $(OBJ_DIR)/atom_cubin.h \ + $(OBJ_DIR)/neighbor_cpu.cubin $(OBJ_DIR)/neighbor_cpu_cubin.h \ + $(OBJ_DIR)/neighbor_gpu.cubin $(OBJ_DIR)/neighbor_gpu_cubin.h \ + $(OBJ_DIR)/pppm_f.cubin $(OBJ_DIR)/pppm_f_cubin.h \ + $(OBJ_DIR)/pppm_d.cubin $(OBJ_DIR)/pppm_d_cubin.h \ + $(OBJ_DIR)/ellipsoid_nbor.cubin $(OBJ_DIR)/ellipsoid_nbor_cubin.h \ + $(OBJ_DIR)/gayberne.cubin $(OBJ_DIR)/gayberne_lj.cubin \ + $(OBJ_DIR)/gayberne_cubin.h $(OBJ_DIR)/gayberne_lj_cubin.h \ + $(OBJ_DIR)/re_squared.cubin $(OBJ_DIR)/re_squared_lj.cubin \ + $(OBJ_DIR)/re_squared_cubin.h $(OBJ_DIR)/re_squared_lj_cubin.h \ + $(OBJ_DIR)/lj.cubin $(OBJ_DIR)/lj_cubin.h \ + $(OBJ_DIR)/lj96.cubin $(OBJ_DIR)/lj96_cubin.h \ + $(OBJ_DIR)/lj_expand.cubin $(OBJ_DIR)/lj_expand_cubin.h \ + $(OBJ_DIR)/lj_coul.cubin $(OBJ_DIR)/lj_coul_cubin.h \ + $(OBJ_DIR)/lj_coul_long.cubin $(OBJ_DIR)/lj_coul_long_cubin.h \ + $(OBJ_DIR)/lj_dsf.cubin $(OBJ_DIR)/lj_dsf_cubin.h \ + $(OBJ_DIR)/lj_class2_long.cubin $(OBJ_DIR)/lj_class2_long_cubin.h \ + $(OBJ_DIR)/coul_long.cubin $(OBJ_DIR)/coul_long_cubin.h \ + $(OBJ_DIR)/morse.cubin $(OBJ_DIR)/morse_cubin.h \ + $(OBJ_DIR)/charmm_long.cubin $(OBJ_DIR)/charmm_long_cubin.h \ + $(OBJ_DIR)/lj_sdk.cubin $(OBJ_DIR)/lj_sdk_cubin.h \ + $(OBJ_DIR)/lj_sdk_long.cubin $(OBJ_DIR)/lj_sdk_long_cubin.h \ + $(OBJ_DIR)/eam.cubin $(OBJ_DIR)/eam_cubin.h \ + $(OBJ_DIR)/buck.cubin $(OBJ_DIR)/buck_cubin.h \ + $(OBJ_DIR)/buck_coul_long.cubin $(OBJ_DIR)/buck_coul_long_cubin.h \ + $(OBJ_DIR)/buck_coul.cubin $(OBJ_DIR)/buck_coul_cubin.h \ + $(OBJ_DIR)/table.cubin $(OBJ_DIR)/table_cubin.h \ + $(OBJ_DIR)/yukawa.cubin $(OBJ_DIR)/yukawa_cubin.h \ + $(OBJ_DIR)/born.cubin $(OBJ_DIR)/born_cubin.h \ + $(OBJ_DIR)/born_coul_wolf.cubin $(OBJ_DIR)/born_coul_wolf_cubin.h \ + $(OBJ_DIR)/born_coul_long.cubin $(OBJ_DIR)/born_coul_long_cubin.h \ + $(OBJ_DIR)/dipole_lj.cubin $(OBJ_DIR)/dipole_lj_cubin.h \ + $(OBJ_DIR)/dipole_lj_sf.cubin $(OBJ_DIR)/dipole_lj_sf_cubin.h \ + $(OBJ_DIR)/colloid.cubin $(OBJ_DIR)/colloid_cubin.h \ + $(OBJ_DIR)/gauss.cubin $(OBJ_DIR)/gauss_cubin.h \ + $(OBJ_DIR)/yukawa_colloid.cubin $(OBJ_DIR)/yukawa_colloid_cubin.h \ + $(OBJ_DIR)/lj_coul_debye.cubin $(OBJ_DIR)/lj_coul_debye_cubin.h \ + $(OBJ_DIR)/coul_dsf.cubin $(OBJ_DIR)/coul_dsf_cubin.h \ + $(OBJ_DIR)/sw.cubin $(OBJ_DIR)/sw_cubin.h \ + $(OBJ_DIR)/vashishta.cubin $(OBJ_DIR)/vashishta_cubin.h \ + $(OBJ_DIR)/beck.cubin $(OBJ_DIR)/beck_cubin.h \ + $(OBJ_DIR)/mie.cubin $(OBJ_DIR)/mie_cubin.h \ + $(OBJ_DIR)/soft.cubin $(OBJ_DIR)/soft_cubin.h \ + $(OBJ_DIR)/lj_coul_msm.cubin $(OBJ_DIR)/lj_coul_msm_cubin.h \ + $(OBJ_DIR)/lj_gromacs.cubin $(OBJ_DIR)/lj_gromacs_cubin.h \ + $(OBJ_DIR)/dpd.cubin $(OBJ_DIR)/dpd_cubin.h \ + $(OBJ_DIR)/tersoff.cubin $(OBJ_DIR)/tersoff_cubin.h \ + $(OBJ_DIR)/tersoff_zbl.cubin $(OBJ_DIR)/tersoff_zbl_cubin.h \ + $(OBJ_DIR)/tersoff_mod.cubin $(OBJ_DIR)/tersoff_mod_cubin.h \ + $(OBJ_DIR)/coul.cubin $(OBJ_DIR)/coul_cubin.h \ + $(OBJ_DIR)/coul_debye.cubin $(OBJ_DIR)/coul_debye_cubin.h \ + $(OBJ_DIR)/zbl.cubin $(OBJ_DIR)/zbl_cubin.h \ + $(OBJ_DIR)/lj_cubic.cubin $(OBJ_DIR)/lj_cubic_cubin.h \ + $(OBJ_DIR)/ufm.cubin $(OBJ_DIR)/ufm_cubin.h \ + $(OBJ_DIR)/dipole_long_lj.cubin $(OBJ_DIR)/dipole_long_lj_cubin.h \ + $(OBJ_DIR)/lj_expand_coul_long.cubin $(OBJ_DIR)/lj_expand_coul_long_cubin.h + +all: $(OBJ_DIR) $(GPU_LIB) $(EXECS) + +$(OBJ_DIR): + mkdir -p $@ + +$(OBJ_DIR)/cudpp.o: cudpp_mini/cudpp.cpp + $(CUDR) -o $@ -c cudpp_mini/cudpp.cpp -Icudpp_mini + +$(OBJ_DIR)/cudpp_plan.o: cudpp_mini/cudpp_plan.cpp + $(CUDR) -o $@ -c cudpp_mini/cudpp_plan.cpp -Icudpp_mini + +$(OBJ_DIR)/cudpp_maximal_launch.o: cudpp_mini/cudpp_maximal_launch.cpp + $(CUDR) -o $@ -c cudpp_mini/cudpp_maximal_launch.cpp -Icudpp_mini + +$(OBJ_DIR)/cudpp_plan_manager.o: cudpp_mini/cudpp_plan_manager.cpp + $(CUDR) -o $@ -c cudpp_mini/cudpp_plan_manager.cpp -Icudpp_mini + +$(OBJ_DIR)/radixsort_app.cu_o: cudpp_mini/radixsort_app.cu + $(CUDA) -o $@ -c cudpp_mini/radixsort_app.cu + +$(OBJ_DIR)/scan_app.cu_o: cudpp_mini/scan_app.cu + $(CUDA) -o $@ -c cudpp_mini/scan_app.cu + +$(OBJ_DIR)/atom.cubin: lal_atom.cu lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_atom.cu + +$(OBJ_DIR)/atom_cubin.h: $(OBJ_DIR)/atom.cubin + $(BIN2C) -c -n atom $(OBJ_DIR)/atom.cubin > $(OBJ_DIR)/atom_cubin.h + +$(OBJ_DIR)/lal_atom.o: lal_atom.cpp lal_atom.h $(NVD_H) $(OBJ_DIR)/atom_cubin.h + $(CUDR) -o $@ -c lal_atom.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_ans.o: lal_answer.cpp lal_answer.h $(NVD_H) + $(CUDR) -o $@ -c lal_answer.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/neighbor_cpu.cubin: lal_neighbor_cpu.cu lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_neighbor_cpu.cu + +$(OBJ_DIR)/neighbor_cpu_cubin.h: $(OBJ_DIR)/neighbor_cpu.cubin + $(BIN2C) -c -n neighbor_cpu $(OBJ_DIR)/neighbor_cpu.cubin > $(OBJ_DIR)/neighbor_cpu_cubin.h + +$(OBJ_DIR)/neighbor_gpu.cubin: lal_neighbor_gpu.cu lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_neighbor_gpu.cu + +$(OBJ_DIR)/neighbor_gpu_cubin.h: $(OBJ_DIR)/neighbor_gpu.cubin + $(BIN2C) -c -n neighbor_gpu $(OBJ_DIR)/neighbor_gpu.cubin > $(OBJ_DIR)/neighbor_gpu_cubin.h + +$(OBJ_DIR)/lal_neighbor_shared.o: lal_neighbor_shared.cpp lal_neighbor_shared.h $(OBJ_DIR)/neighbor_cpu_cubin.h $(OBJ_DIR)/neighbor_gpu_cubin.h $(NVD_H) + $(CUDR) -o $@ -c lal_neighbor_shared.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_neighbor.o: lal_neighbor.cpp lal_neighbor.h lal_neighbor_shared.h $(NVD_H) + $(CUDR) -o $@ -c lal_neighbor.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/device.cubin: lal_device.cu lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_device.cu + +$(OBJ_DIR)/device_cubin.h: $(OBJ_DIR)/device.cubin + $(BIN2C) -c -n device $(OBJ_DIR)/device.cubin > $(OBJ_DIR)/device_cubin.h + +$(OBJ_DIR)/lal_device.o: lal_device.cpp lal_device.h $(ALL_H) $(OBJ_DIR)/device_cubin.h + $(CUDR) -o $@ -c lal_device.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_base_atomic.o: $(ALL_H) lal_base_atomic.h lal_base_atomic.cpp + $(CUDR) -o $@ -c lal_base_atomic.cpp + +$(OBJ_DIR)/lal_base_charge.o: $(ALL_H) lal_base_charge.h lal_base_charge.cpp + $(CUDR) -o $@ -c lal_base_charge.cpp + +$(OBJ_DIR)/lal_base_ellipsoid.o: $(ALL_H) lal_base_ellipsoid.h lal_base_ellipsoid.cpp $(OBJ_DIR)/ellipsoid_nbor_cubin.h + $(CUDR) -o $@ -c lal_base_ellipsoid.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_base_dipole.o: $(ALL_H) lal_base_dipole.h lal_base_dipole.cpp + $(CUDR) -o $@ -c lal_base_dipole.cpp + +$(OBJ_DIR)/lal_base_three.o: $(ALL_H) lal_base_three.h lal_base_three.cpp + $(CUDR) -o $@ -c lal_base_three.cpp + +$(OBJ_DIR)/lal_base_dpd.o: $(ALL_H) lal_base_dpd.h lal_base_dpd.cpp + $(CUDR) -o $@ -c lal_base_dpd.cpp + +$(OBJ_DIR)/pppm_f.cubin: lal_pppm.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -Dgrdtyp=float -Dgrdtyp4=float4 -o $@ lal_pppm.cu + +$(OBJ_DIR)/pppm_f_cubin.h: $(OBJ_DIR)/pppm_f.cubin + $(BIN2C) -c -n pppm_f $(OBJ_DIR)/pppm_f.cubin > $(OBJ_DIR)/pppm_f_cubin.h + +$(OBJ_DIR)/pppm_d.cubin: lal_pppm.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -Dgrdtyp=double -Dgrdtyp4=double4 -o $@ lal_pppm.cu + +$(OBJ_DIR)/pppm_d_cubin.h: $(OBJ_DIR)/pppm_d.cubin + $(BIN2C) -c -n pppm_d $(OBJ_DIR)/pppm_d.cubin > $(OBJ_DIR)/pppm_d_cubin.h + +$(OBJ_DIR)/lal_pppm.o: $(ALL_H) lal_pppm.h lal_pppm.cpp $(OBJ_DIR)/pppm_f_cubin.h $(OBJ_DIR)/pppm_d_cubin.h + $(CUDR) -o $@ -c lal_pppm.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_pppm_ext.o: $(ALL_H) lal_pppm.h lal_pppm_ext.cpp + $(CUDR) -o $@ -c lal_pppm_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/ellipsoid_nbor.cubin: lal_ellipsoid_nbor.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_ellipsoid_nbor.cu + +$(OBJ_DIR)/ellipsoid_nbor_cubin.h: $(OBJ_DIR)/ellipsoid_nbor.cubin + $(BIN2C) -c -n ellipsoid_nbor $(OBJ_DIR)/ellipsoid_nbor.cubin > $(OBJ_DIR)/ellipsoid_nbor_cubin.h + +$(OBJ_DIR)/gayberne.cubin: lal_gayberne.cu lal_precision.h lal_ellipsoid_extra.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_gayberne.cu + +$(OBJ_DIR)/gayberne_lj.cubin: lal_gayberne_lj.cu lal_precision.h lal_ellipsoid_extra.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_gayberne_lj.cu + +$(OBJ_DIR)/gayberne_cubin.h: $(OBJ_DIR)/gayberne.cubin + $(BIN2C) -c -n gayberne $(OBJ_DIR)/gayberne.cubin > $(OBJ_DIR)/gayberne_cubin.h + +$(OBJ_DIR)/gayberne_lj_cubin.h: $(OBJ_DIR)/gayberne_lj.cubin + $(BIN2C) -c -n gayberne_lj $(OBJ_DIR)/gayberne_lj.cubin > $(OBJ_DIR)/gayberne_lj_cubin.h + +$(OBJ_DIR)/lal_gayberne.o: $(ALL_H) lal_gayberne.h lal_gayberne.cpp $(OBJ_DIR)/gayberne_cubin.h $(OBJ_DIR)/gayberne_lj_cubin.h $(OBJ_DIR)/lal_base_ellipsoid.o + $(CUDR) -o $@ -c lal_gayberne.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_gayberne_ext.o: $(ALL_H) $(OBJ_DIR)/lal_gayberne.o lal_gayberne_ext.cpp + $(CUDR) -o $@ -c lal_gayberne_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/re_squared.cubin: lal_re_squared.cu lal_precision.h lal_ellipsoid_extra.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_re_squared.cu + +$(OBJ_DIR)/re_squared_lj.cubin: lal_re_squared_lj.cu lal_precision.h lal_ellipsoid_extra.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_re_squared_lj.cu + +$(OBJ_DIR)/re_squared_cubin.h: $(OBJ_DIR)/re_squared.cubin + $(BIN2C) -c -n re_squared $(OBJ_DIR)/re_squared.cubin > $(OBJ_DIR)/re_squared_cubin.h + +$(OBJ_DIR)/re_squared_lj_cubin.h: $(OBJ_DIR)/re_squared_lj.cubin + $(BIN2C) -c -n re_squared_lj $(OBJ_DIR)/re_squared_lj.cubin > $(OBJ_DIR)/re_squared_lj_cubin.h + +$(OBJ_DIR)/lal_re_squared.o: $(ALL_H) lal_re_squared.h lal_re_squared.cpp $(OBJ_DIR)/re_squared_cubin.h $(OBJ_DIR)/re_squared_lj_cubin.h $(OBJ_DIR)/lal_base_ellipsoid.o + $(CUDR) -o $@ -c lal_re_squared.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_re_squared_ext.o: $(ALL_H) $(OBJ_DIR)/lal_re_squared.o lal_re_squared_ext.cpp + $(CUDR) -o $@ -c lal_re_squared_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lj.cubin: lal_lj.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_lj.cu + +$(OBJ_DIR)/lj_cubin.h: $(OBJ_DIR)/lj.cubin $(OBJ_DIR)/lj.cubin + $(BIN2C) -c -n lj $(OBJ_DIR)/lj.cubin > $(OBJ_DIR)/lj_cubin.h + +$(OBJ_DIR)/lal_lj.o: $(ALL_H) lal_lj.h lal_lj.cpp $(OBJ_DIR)/lj_cubin.h $(OBJ_DIR)/lal_base_atomic.o + $(CUDR) -o $@ -c lal_lj.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_lj_ext.o: $(ALL_H) lal_lj.h lal_lj_ext.cpp lal_base_atomic.h + $(CUDR) -o $@ -c lal_lj_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lj_coul.cubin: lal_lj_coul.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_lj_coul.cu + +$(OBJ_DIR)/lj_coul_cubin.h: $(OBJ_DIR)/lj_coul.cubin $(OBJ_DIR)/lj_coul.cubin + $(BIN2C) -c -n lj_coul $(OBJ_DIR)/lj_coul.cubin > $(OBJ_DIR)/lj_coul_cubin.h + +$(OBJ_DIR)/lal_lj_coul.o: $(ALL_H) lal_lj_coul.h lal_lj_coul.cpp $(OBJ_DIR)/lj_coul_cubin.h $(OBJ_DIR)/lal_base_charge.o + $(CUDR) -o $@ -c lal_lj_coul.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_lj_coul_ext.o: $(ALL_H) lal_lj_coul.h lal_lj_coul_ext.cpp lal_base_charge.h + $(CUDR) -o $@ -c lal_lj_coul_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lj_class2_long.cubin: lal_lj_class2_long.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_lj_class2_long.cu + +$(OBJ_DIR)/lj_class2_long_cubin.h: $(OBJ_DIR)/lj_class2_long.cubin $(OBJ_DIR)/lj_class2_long.cubin + $(BIN2C) -c -n lj_class2_long $(OBJ_DIR)/lj_class2_long.cubin > $(OBJ_DIR)/lj_class2_long_cubin.h + +$(OBJ_DIR)/lal_lj_class2_long.o: $(ALL_H) lal_lj_class2_long.h lal_lj_class2_long.cpp $(OBJ_DIR)/lj_class2_long_cubin.h $(OBJ_DIR)/lal_base_charge.o + $(CUDR) -o $@ -c lal_lj_class2_long.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_lj_class2_long_ext.o: $(ALL_H) lal_lj_class2_long.h lal_lj_class2_long_ext.cpp lal_base_charge.h + $(CUDR) -o $@ -c lal_lj_class2_long_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/coul_long.cubin: lal_coul_long.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_coul_long.cu + +$(OBJ_DIR)/coul_long_cubin.h: $(OBJ_DIR)/coul_long.cubin $(OBJ_DIR)/coul_long.cubin + $(BIN2C) -c -n coul_long $(OBJ_DIR)/coul_long.cubin > $(OBJ_DIR)/coul_long_cubin.h + +$(OBJ_DIR)/lal_coul_long.o: $(ALL_H) lal_coul_long.h lal_coul_long.cpp $(OBJ_DIR)/coul_long_cubin.h $(OBJ_DIR)/lal_base_charge.o + $(CUDR) -o $@ -c lal_coul_long.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_coul_long_ext.o: $(ALL_H) lal_coul_long.h lal_coul_long_ext.cpp lal_base_charge.h + $(CUDR) -o $@ -c lal_coul_long_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lj_coul_long.cubin: lal_lj_coul_long.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_lj_coul_long.cu + +$(OBJ_DIR)/lj_coul_long_cubin.h: $(OBJ_DIR)/lj_coul_long.cubin $(OBJ_DIR)/lj_coul_long.cubin + $(BIN2C) -c -n lj_coul_long $(OBJ_DIR)/lj_coul_long.cubin > $(OBJ_DIR)/lj_coul_long_cubin.h + +$(OBJ_DIR)/lal_lj_coul_long.o: $(ALL_H) lal_lj_coul_long.h lal_lj_coul_long.cpp $(OBJ_DIR)/lj_coul_long_cubin.h $(OBJ_DIR)/lal_base_charge.o + $(CUDR) -o $@ -c lal_lj_coul_long.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_lj_coul_long_ext.o: $(ALL_H) lal_lj_coul_long.h lal_lj_coul_long_ext.cpp lal_base_charge.h + $(CUDR) -o $@ -c lal_lj_coul_long_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lj_dsf.cubin: lal_lj_dsf.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_lj_dsf.cu + +$(OBJ_DIR)/lj_dsf_cubin.h: $(OBJ_DIR)/lj_dsf.cubin $(OBJ_DIR)/lj_dsf.cubin + $(BIN2C) -c -n lj_dsf $(OBJ_DIR)/lj_dsf.cubin > $(OBJ_DIR)/lj_dsf_cubin.h + +$(OBJ_DIR)/lal_lj_dsf.o: $(ALL_H) lal_lj_dsf.h lal_lj_dsf.cpp $(OBJ_DIR)/lj_dsf_cubin.h $(OBJ_DIR)/lal_base_charge.o + $(CUDR) -o $@ -c lal_lj_dsf.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_lj_dsf_ext.o: $(ALL_H) lal_lj_dsf.h lal_lj_dsf_ext.cpp lal_base_charge.h + $(CUDR) -o $@ -c lal_lj_dsf_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/morse.cubin: lal_morse.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_morse.cu + +$(OBJ_DIR)/morse_cubin.h: $(OBJ_DIR)/morse.cubin $(OBJ_DIR)/morse.cubin + $(BIN2C) -c -n morse $(OBJ_DIR)/morse.cubin > $(OBJ_DIR)/morse_cubin.h + +$(OBJ_DIR)/lal_morse.o: $(ALL_H) lal_morse.h lal_morse.cpp $(OBJ_DIR)/morse_cubin.h $(OBJ_DIR)/lal_base_atomic.o + $(CUDR) -o $@ -c lal_morse.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_morse_ext.o: $(ALL_H) lal_morse.h lal_morse_ext.cpp lal_base_atomic.h + $(CUDR) -o $@ -c lal_morse_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/charmm_long.cubin: lal_charmm_long.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_charmm_long.cu + +$(OBJ_DIR)/charmm_long_cubin.h: $(OBJ_DIR)/charmm_long.cubin $(OBJ_DIR)/charmm_long.cubin + $(BIN2C) -c -n charmm_long $(OBJ_DIR)/charmm_long.cubin > $(OBJ_DIR)/charmm_long_cubin.h + +$(OBJ_DIR)/lal_charmm_long.o: $(ALL_H) lal_charmm_long.h lal_charmm_long.cpp $(OBJ_DIR)/charmm_long_cubin.h $(OBJ_DIR)/lal_base_charge.o + $(CUDR) -o $@ -c lal_charmm_long.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_charmm_long_ext.o: $(ALL_H) lal_charmm_long.h lal_charmm_long_ext.cpp lal_base_charge.h + $(CUDR) -o $@ -c lal_charmm_long_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lj96.cubin: lal_lj96.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_lj96.cu + +$(OBJ_DIR)/lj96_cubin.h: $(OBJ_DIR)/lj96.cubin $(OBJ_DIR)/lj96.cubin + $(BIN2C) -c -n lj96 $(OBJ_DIR)/lj96.cubin > $(OBJ_DIR)/lj96_cubin.h + +$(OBJ_DIR)/lal_lj96.o: $(ALL_H) lal_lj96.h lal_lj96.cpp $(OBJ_DIR)/lj96_cubin.h $(OBJ_DIR)/lal_base_atomic.o + $(CUDR) -o $@ -c lal_lj96.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_lj96_ext.o: $(ALL_H) lal_lj96.h lal_lj96_ext.cpp lal_base_atomic.h + $(CUDR) -o $@ -c lal_lj96_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lj_expand.cubin: lal_lj_expand.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_lj_expand.cu + +$(OBJ_DIR)/lj_expand_cubin.h: $(OBJ_DIR)/lj_expand.cubin $(OBJ_DIR)/lj_expand.cubin + $(BIN2C) -c -n lj_expand $(OBJ_DIR)/lj_expand.cubin > $(OBJ_DIR)/lj_expand_cubin.h + +$(OBJ_DIR)/lal_lj_expand.o: $(ALL_H) lal_lj_expand.h lal_lj_expand.cpp $(OBJ_DIR)/lj_expand_cubin.h $(OBJ_DIR)/lal_base_atomic.o + $(CUDR) -o $@ -c lal_lj_expand.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_lj_expand_ext.o: $(ALL_H) lal_lj_expand.h lal_lj_expand_ext.cpp lal_base_atomic.h + $(CUDR) -o $@ -c lal_lj_expand_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lj_sdk.cubin: lal_lj_sdk.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_lj_sdk.cu + +$(OBJ_DIR)/lj_sdk_cubin.h: $(OBJ_DIR)/lj_sdk.cubin $(OBJ_DIR)/lj_sdk.cubin + $(BIN2C) -c -n lj_sdk $(OBJ_DIR)/lj_sdk.cubin > $(OBJ_DIR)/lj_sdk_cubin.h + +$(OBJ_DIR)/lal_lj_sdk.o: $(ALL_H) lal_lj_sdk.h lal_lj_sdk.cpp $(OBJ_DIR)/lj_sdk_cubin.h $(OBJ_DIR)/lal_base_atomic.o + $(CUDR) -o $@ -c lal_lj_sdk.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_lj_sdk_ext.o: $(ALL_H) lal_lj_sdk.h lal_lj_sdk_ext.cpp lal_base_atomic.h + $(CUDR) -o $@ -c lal_lj_sdk_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lj_sdk_long.cubin: lal_lj_sdk_long.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_lj_sdk_long.cu + +$(OBJ_DIR)/lj_sdk_long_cubin.h: $(OBJ_DIR)/lj_sdk_long.cubin $(OBJ_DIR)/lj_sdk_long.cubin + $(BIN2C) -c -n lj_sdk_long $(OBJ_DIR)/lj_sdk_long.cubin > $(OBJ_DIR)/lj_sdk_long_cubin.h + +$(OBJ_DIR)/lal_lj_sdk_long.o: $(ALL_H) lal_lj_sdk_long.h lal_lj_sdk_long.cpp $(OBJ_DIR)/lj_sdk_long_cubin.h $(OBJ_DIR)/lal_base_atomic.o + $(CUDR) -o $@ -c lal_lj_sdk_long.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_lj_sdk_long_ext.o: $(ALL_H) lal_lj_sdk_long.h lal_lj_sdk_long_ext.cpp lal_base_charge.h + $(CUDR) -o $@ -c lal_lj_sdk_long_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/eam.cubin: lal_eam.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_eam.cu + +$(OBJ_DIR)/eam_cubin.h: $(OBJ_DIR)/eam.cubin $(OBJ_DIR)/eam.cubin + $(BIN2C) -c -n eam $(OBJ_DIR)/eam.cubin > $(OBJ_DIR)/eam_cubin.h + +$(OBJ_DIR)/lal_eam.o: $(ALL_H) lal_eam.h lal_eam.cpp $(OBJ_DIR)/eam_cubin.h $(OBJ_DIR)/lal_base_atomic.o + $(CUDR) -o $@ -c lal_eam.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_eam_ext.o: $(ALL_H) lal_eam.h lal_eam_ext.cpp lal_base_atomic.h + $(CUDR) -o $@ -c lal_eam_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_eam_fs_ext.o: $(ALL_H) lal_eam.h lal_eam_fs_ext.cpp lal_base_atomic.h + $(CUDR) -o $@ -c lal_eam_fs_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_eam_alloy_ext.o: $(ALL_H) lal_eam.h lal_eam_alloy_ext.cpp lal_base_atomic.h + $(CUDR) -o $@ -c lal_eam_alloy_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/buck.cubin: lal_buck.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_buck.cu + +$(OBJ_DIR)/buck_cubin.h: $(OBJ_DIR)/buck.cubin $(OBJ_DIR)/buck.cubin + $(BIN2C) -c -n buck $(OBJ_DIR)/buck.cubin > $(OBJ_DIR)/buck_cubin.h + +$(OBJ_DIR)/lal_buck.o: $(ALL_H) lal_buck.h lal_buck.cpp $(OBJ_DIR)/buck_cubin.h $(OBJ_DIR)/lal_base_atomic.o + $(CUDR) -o $@ -c lal_buck.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_buck_ext.o: $(ALL_H) lal_buck.h lal_buck_ext.cpp lal_base_atomic.h + $(CUDR) -o $@ -c lal_buck_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/buck_coul.cubin: lal_buck_coul.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_buck_coul.cu + +$(OBJ_DIR)/buck_coul_cubin.h: $(OBJ_DIR)/buck_coul.cubin $(OBJ_DIR)/buck_coul.cubin + $(BIN2C) -c -n buck_coul $(OBJ_DIR)/buck_coul.cubin > $(OBJ_DIR)/buck_coul_cubin.h + +$(OBJ_DIR)/lal_buck_coul.o: $(ALL_H) lal_buck_coul.h lal_buck_coul.cpp $(OBJ_DIR)/buck_coul_cubin.h $(OBJ_DIR)/lal_base_charge.o + $(CUDR) -o $@ -c lal_buck_coul.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_buck_coul_ext.o: $(ALL_H) lal_buck_coul.h lal_buck_coul_ext.cpp lal_base_charge.h + $(CUDR) -o $@ -c lal_buck_coul_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/buck_coul_long.cubin: lal_buck_coul_long.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_buck_coul_long.cu + +$(OBJ_DIR)/buck_coul_long_cubin.h: $(OBJ_DIR)/buck_coul_long.cubin $(OBJ_DIR)/buck_coul_long.cubin + $(BIN2C) -c -n buck_coul_long $(OBJ_DIR)/buck_coul_long.cubin > $(OBJ_DIR)/buck_coul_long_cubin.h + +$(OBJ_DIR)/lal_buck_coul_long.o: $(ALL_H) lal_buck_coul_long.h lal_buck_coul_long.cpp $(OBJ_DIR)/buck_coul_long_cubin.h $(OBJ_DIR)/lal_base_charge.o + $(CUDR) -o $@ -c lal_buck_coul_long.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_buck_coul_long_ext.o: $(ALL_H) lal_buck_coul_long.h lal_buck_coul_long_ext.cpp lal_base_charge.h + $(CUDR) -o $@ -c lal_buck_coul_long_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/table.cubin: lal_table.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_table.cu + +$(OBJ_DIR)/table_cubin.h: $(OBJ_DIR)/table.cubin $(OBJ_DIR)/table.cubin + $(BIN2C) -c -n table $(OBJ_DIR)/table.cubin > $(OBJ_DIR)/table_cubin.h + +$(OBJ_DIR)/lal_table.o: $(ALL_H) lal_table.h lal_table.cpp $(OBJ_DIR)/table_cubin.h $(OBJ_DIR)/lal_base_atomic.o + $(CUDR) -o $@ -c lal_table.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_table_ext.o: $(ALL_H) lal_table.h lal_table_ext.cpp lal_base_atomic.h + $(CUDR) -o $@ -c lal_table_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/yukawa.cubin: lal_yukawa.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_yukawa.cu + +$(OBJ_DIR)/yukawa_cubin.h: $(OBJ_DIR)/yukawa.cubin $(OBJ_DIR)/yukawa.cubin + $(BIN2C) -c -n yukawa $(OBJ_DIR)/yukawa.cubin > $(OBJ_DIR)/yukawa_cubin.h + +$(OBJ_DIR)/lal_yukawa.o: $(ALL_H) lal_yukawa.h lal_yukawa.cpp $(OBJ_DIR)/yukawa_cubin.h $(OBJ_DIR)/lal_base_atomic.o + $(CUDR) -o $@ -c lal_yukawa.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_yukawa_ext.o: $(ALL_H) lal_yukawa.h lal_yukawa_ext.cpp lal_base_atomic.h + $(CUDR) -o $@ -c lal_yukawa_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/born.cubin: lal_born.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_born.cu + +$(OBJ_DIR)/born_cubin.h: $(OBJ_DIR)/born.cubin $(OBJ_DIR)/born.cubin + $(BIN2C) -c -n born $(OBJ_DIR)/born.cubin > $(OBJ_DIR)/born_cubin.h + +$(OBJ_DIR)/lal_born.o: $(ALL_H) lal_born.h lal_born.cpp $(OBJ_DIR)/born_cubin.h $(OBJ_DIR)/lal_base_atomic.o + $(CUDR) -o $@ -c lal_born.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_born_ext.o: $(ALL_H) lal_born.h lal_born_ext.cpp lal_base_atomic.h + $(CUDR) -o $@ -c lal_born_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/born_coul_wolf.cubin: lal_born_coul_wolf.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_born_coul_wolf.cu + +$(OBJ_DIR)/born_coul_wolf_cubin.h: $(OBJ_DIR)/born_coul_wolf.cubin $(OBJ_DIR)/born_coul_wolf.cubin + $(BIN2C) -c -n born_coul_wolf $(OBJ_DIR)/born_coul_wolf.cubin > $(OBJ_DIR)/born_coul_wolf_cubin.h + +$(OBJ_DIR)/lal_born_coul_wolf.o: $(ALL_H) lal_born_coul_wolf.h lal_born_coul_wolf.cpp $(OBJ_DIR)/born_coul_wolf_cubin.h $(OBJ_DIR)/lal_base_charge.o + $(CUDR) -o $@ -c lal_born_coul_wolf.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_born_coul_wolf_ext.o: $(ALL_H) lal_born_coul_wolf.h lal_born_coul_wolf_ext.cpp lal_base_charge.h + $(CUDR) -o $@ -c lal_born_coul_wolf_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/born_coul_long.cubin: lal_born_coul_long.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_born_coul_long.cu + +$(OBJ_DIR)/born_coul_long_cubin.h: $(OBJ_DIR)/born_coul_long.cubin $(OBJ_DIR)/born_coul_long.cubin + $(BIN2C) -c -n born_coul_long $(OBJ_DIR)/born_coul_long.cubin > $(OBJ_DIR)/born_coul_long_cubin.h + +$(OBJ_DIR)/lal_born_coul_long.o: $(ALL_H) lal_born_coul_long.h lal_born_coul_long.cpp $(OBJ_DIR)/born_coul_long_cubin.h $(OBJ_DIR)/lal_base_charge.o + $(CUDR) -o $@ -c lal_born_coul_long.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_born_coul_long_ext.o: $(ALL_H) lal_born_coul_long.h lal_born_coul_long_ext.cpp lal_base_charge.h + $(CUDR) -o $@ -c lal_born_coul_long_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/dipole_lj.cubin: lal_dipole_lj.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_dipole_lj.cu + +$(OBJ_DIR)/dipole_lj_cubin.h: $(OBJ_DIR)/dipole_lj.cubin $(OBJ_DIR)/dipole_lj.cubin + $(BIN2C) -c -n dipole_lj $(OBJ_DIR)/dipole_lj.cubin > $(OBJ_DIR)/dipole_lj_cubin.h + +$(OBJ_DIR)/lal_dipole_lj.o: $(ALL_H) lal_dipole_lj.h lal_dipole_lj.cpp $(OBJ_DIR)/dipole_lj_cubin.h $(OBJ_DIR)/lal_base_dipole.o + $(CUDR) -o $@ -c lal_dipole_lj.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_dipole_lj_ext.o: $(ALL_H) lal_dipole_lj.h lal_dipole_lj_ext.cpp lal_base_dipole.h + $(CUDR) -o $@ -c lal_dipole_lj_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/dipole_lj_sf.cubin: lal_dipole_lj_sf.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_dipole_lj_sf.cu + +$(OBJ_DIR)/dipole_lj_sf_cubin.h: $(OBJ_DIR)/dipole_lj_sf.cubin $(OBJ_DIR)/dipole_lj_sf.cubin + $(BIN2C) -c -n dipole_lj_sf $(OBJ_DIR)/dipole_lj_sf.cubin > $(OBJ_DIR)/dipole_lj_sf_cubin.h + +$(OBJ_DIR)/lal_dipole_lj_sf.o: $(ALL_H) lal_dipole_lj_sf.h lal_dipole_lj_sf.cpp $(OBJ_DIR)/dipole_lj_sf_cubin.h $(OBJ_DIR)/lal_base_dipole.o + $(CUDR) -o $@ -c lal_dipole_lj_sf.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_dipole_lj_sf_ext.o: $(ALL_H) lal_dipole_lj_sf.h lal_dipole_lj_sf_ext.cpp lal_base_dipole.h + $(CUDR) -o $@ -c lal_dipole_lj_sf_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/colloid.cubin: lal_colloid.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_colloid.cu + +$(OBJ_DIR)/colloid_cubin.h: $(OBJ_DIR)/colloid.cubin $(OBJ_DIR)/colloid.cubin + $(BIN2C) -c -n colloid $(OBJ_DIR)/colloid.cubin > $(OBJ_DIR)/colloid_cubin.h + +$(OBJ_DIR)/lal_colloid.o: $(ALL_H) lal_colloid.h lal_colloid.cpp $(OBJ_DIR)/colloid_cubin.h $(OBJ_DIR)/lal_base_atomic.o + $(CUDR) -o $@ -c lal_colloid.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_colloid_ext.o: $(ALL_H) lal_colloid.h lal_colloid_ext.cpp lal_base_atomic.h + $(CUDR) -o $@ -c lal_colloid_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/gauss.cubin: lal_gauss.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_gauss.cu + +$(OBJ_DIR)/gauss_cubin.h: $(OBJ_DIR)/gauss.cubin $(OBJ_DIR)/gauss.cubin + $(BIN2C) -c -n gauss $(OBJ_DIR)/gauss.cubin > $(OBJ_DIR)/gauss_cubin.h + +$(OBJ_DIR)/lal_gauss.o: $(ALL_H) lal_gauss.h lal_gauss.cpp $(OBJ_DIR)/gauss_cubin.h $(OBJ_DIR)/lal_base_atomic.o + $(CUDR) -o $@ -c lal_gauss.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_gauss_ext.o: $(ALL_H) lal_gauss.h lal_gauss_ext.cpp lal_base_atomic.h + $(CUDR) -o $@ -c lal_gauss_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/yukawa_colloid.cubin: lal_yukawa_colloid.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_yukawa_colloid.cu + +$(OBJ_DIR)/yukawa_colloid_cubin.h: $(OBJ_DIR)/yukawa_colloid.cubin $(OBJ_DIR)/yukawa_colloid.cubin + $(BIN2C) -c -n yukawa_colloid $(OBJ_DIR)/yukawa_colloid.cubin > $(OBJ_DIR)/yukawa_colloid_cubin.h + +$(OBJ_DIR)/lal_yukawa_colloid.o: $(ALL_H) lal_yukawa_colloid.h lal_yukawa_colloid.cpp $(OBJ_DIR)/yukawa_colloid_cubin.h $(OBJ_DIR)/lal_base_atomic.o + $(CUDR) -o $@ -c lal_yukawa_colloid.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_yukawa_colloid_ext.o: $(ALL_H) lal_yukawa_colloid.h lal_yukawa_colloid_ext.cpp lal_base_atomic.h + $(CUDR) -o $@ -c lal_yukawa_colloid_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lj_coul_debye.cubin: lal_lj_coul_debye.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_lj_coul_debye.cu + +$(OBJ_DIR)/lj_coul_debye_cubin.h: $(OBJ_DIR)/lj_coul_debye.cubin $(OBJ_DIR)/lj_coul_debye.cubin + $(BIN2C) -c -n lj_coul_debye $(OBJ_DIR)/lj_coul_debye.cubin > $(OBJ_DIR)/lj_coul_debye_cubin.h + +$(OBJ_DIR)/lal_lj_coul_debye.o: $(ALL_H) lal_lj_coul_debye.h lal_lj_coul_debye.cpp $(OBJ_DIR)/lj_coul_debye_cubin.h $(OBJ_DIR)/lal_base_charge.o + $(CUDR) -o $@ -c lal_lj_coul_debye.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_lj_coul_debye_ext.o: $(ALL_H) lal_lj_coul_debye.h lal_lj_coul_debye_ext.cpp lal_base_charge.h + $(CUDR) -o $@ -c lal_lj_coul_debye_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/coul_dsf.cubin: lal_coul_dsf.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_coul_dsf.cu + +$(OBJ_DIR)/coul_dsf_cubin.h: $(OBJ_DIR)/coul_dsf.cubin $(OBJ_DIR)/coul_dsf.cubin + $(BIN2C) -c -n coul_dsf $(OBJ_DIR)/coul_dsf.cubin > $(OBJ_DIR)/coul_dsf_cubin.h + +$(OBJ_DIR)/lal_coul_dsf.o: $(ALL_H) lal_coul_dsf.h lal_coul_dsf.cpp $(OBJ_DIR)/coul_dsf_cubin.h $(OBJ_DIR)/lal_base_charge.o + $(CUDR) -o $@ -c lal_coul_dsf.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_coul_dsf_ext.o: $(ALL_H) lal_coul_dsf.h lal_coul_dsf_ext.cpp lal_base_charge.h + $(CUDR) -o $@ -c lal_coul_dsf_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/vashishta.cubin: lal_vashishta.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_vashishta.cu + +$(OBJ_DIR)/vashishta_cubin.h: $(OBJ_DIR)/vashishta.cubin $(OBJ_DIR)/vashishta.cubin + $(BIN2C) -c -n vashishta $(OBJ_DIR)/vashishta.cubin > $(OBJ_DIR)/vashishta_cubin.h + +$(OBJ_DIR)/lal_vashishta.o: $(ALL_H) lal_vashishta.h lal_vashishta.cpp $(OBJ_DIR)/vashishta_cubin.h $(OBJ_DIR)/lal_base_three.o + $(CUDR) -o $@ -c lal_vashishta.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_vashishta_ext.o: $(ALL_H) lal_vashishta.h lal_vashishta_ext.cpp lal_base_three.h + $(CUDR) -o $@ -c lal_vashishta_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/sw.cubin: lal_sw.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_sw.cu + +$(OBJ_DIR)/sw_cubin.h: $(OBJ_DIR)/sw.cubin $(OBJ_DIR)/sw.cubin + $(BIN2C) -c -n sw $(OBJ_DIR)/sw.cubin > $(OBJ_DIR)/sw_cubin.h + +$(OBJ_DIR)/lal_sw.o: $(ALL_H) lal_sw.h lal_sw.cpp $(OBJ_DIR)/sw_cubin.h $(OBJ_DIR)/lal_base_three.o + $(CUDR) -o $@ -c lal_sw.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_sw_ext.o: $(ALL_H) lal_sw.h lal_sw_ext.cpp lal_base_three.h + $(CUDR) -o $@ -c lal_sw_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/beck.cubin: lal_beck.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_beck.cu + +$(OBJ_DIR)/beck_cubin.h: $(OBJ_DIR)/beck.cubin $(OBJ_DIR)/beck.cubin + $(BIN2C) -c -n beck $(OBJ_DIR)/beck.cubin > $(OBJ_DIR)/beck_cubin.h + +$(OBJ_DIR)/lal_beck.o: $(ALL_H) lal_beck.h lal_beck.cpp $(OBJ_DIR)/beck_cubin.h $(OBJ_DIR)/lal_base_atomic.o + $(CUDR) -o $@ -c lal_beck.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_beck_ext.o: $(ALL_H) lal_beck.h lal_beck_ext.cpp lal_base_atomic.h + $(CUDR) -o $@ -c lal_beck_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/mie.cubin: lal_mie.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_mie.cu + +$(OBJ_DIR)/mie_cubin.h: $(OBJ_DIR)/mie.cubin $(OBJ_DIR)/mie.cubin + $(BIN2C) -c -n mie $(OBJ_DIR)/mie.cubin > $(OBJ_DIR)/mie_cubin.h + +$(OBJ_DIR)/lal_mie.o: $(ALL_H) lal_mie.h lal_mie.cpp $(OBJ_DIR)/mie_cubin.h $(OBJ_DIR)/lal_base_atomic.o + $(CUDR) -o $@ -c lal_mie.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_mie_ext.o: $(ALL_H) lal_mie.h lal_mie_ext.cpp lal_base_atomic.h + $(CUDR) -o $@ -c lal_mie_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/soft.cubin: lal_soft.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_soft.cu + +$(OBJ_DIR)/soft_cubin.h: $(OBJ_DIR)/soft.cubin $(OBJ_DIR)/soft.cubin + $(BIN2C) -c -n soft $(OBJ_DIR)/soft.cubin > $(OBJ_DIR)/soft_cubin.h + +$(OBJ_DIR)/lal_soft.o: $(ALL_H) lal_soft.h lal_soft.cpp $(OBJ_DIR)/soft_cubin.h $(OBJ_DIR)/lal_base_atomic.o + $(CUDR) -o $@ -c lal_soft.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_soft_ext.o: $(ALL_H) lal_soft.h lal_soft_ext.cpp lal_base_atomic.h + $(CUDR) -o $@ -c lal_soft_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lj_coul_msm.cubin: lal_lj_coul_msm.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_lj_coul_msm.cu + +$(OBJ_DIR)/lj_coul_msm_cubin.h: $(OBJ_DIR)/lj_coul_msm.cubin $(OBJ_DIR)/lj_coul_msm.cubin + $(BIN2C) -c -n lj_coul_msm $(OBJ_DIR)/lj_coul_msm.cubin > $(OBJ_DIR)/lj_coul_msm_cubin.h + +$(OBJ_DIR)/lal_lj_coul_msm.o: $(ALL_H) lal_lj_coul_msm.h lal_lj_coul_msm.cpp $(OBJ_DIR)/lj_coul_msm_cubin.h $(OBJ_DIR)/lal_base_charge.o + $(CUDR) -o $@ -c lal_lj_coul_msm.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_lj_coul_msm_ext.o: $(ALL_H) lal_lj_coul_msm.h lal_lj_coul_msm_ext.cpp lal_base_charge.h + $(CUDR) -o $@ -c lal_lj_coul_msm_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lj_gromacs.cubin: lal_lj_gromacs.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_lj_gromacs.cu + +$(OBJ_DIR)/lj_gromacs_cubin.h: $(OBJ_DIR)/lj_gromacs.cubin $(OBJ_DIR)/lj_gromacs.cubin + $(BIN2C) -c -n lj_gromacs $(OBJ_DIR)/lj_gromacs.cubin > $(OBJ_DIR)/lj_gromacs_cubin.h + +$(OBJ_DIR)/lal_lj_gromacs.o: $(ALL_H) lal_lj_gromacs.h lal_lj_gromacs.cpp $(OBJ_DIR)/lj_gromacs_cubin.h $(OBJ_DIR)/lal_base_atomic.o + $(CUDR) -o $@ -c lal_lj_gromacs.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_lj_gromacs_ext.o: $(ALL_H) lal_lj_gromacs.h lal_lj_gromacs_ext.cpp lal_base_atomic.h + $(CUDR) -o $@ -c lal_lj_gromacs_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/dpd.cubin: lal_dpd.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_dpd.cu + +$(OBJ_DIR)/dpd_cubin.h: $(OBJ_DIR)/dpd.cubin $(OBJ_DIR)/dpd.cubin + $(BIN2C) -c -n dpd $(OBJ_DIR)/dpd.cubin > $(OBJ_DIR)/dpd_cubin.h + +$(OBJ_DIR)/ufm.cubin: lal_ufm.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_ufm.cu + +$(OBJ_DIR)/ufm_cubin.h: $(OBJ_DIR)/ufm.cubin $(OBJ_DIR)/ufm.cubin + $(BIN2C) -c -n ufm $(OBJ_DIR)/ufm.cubin > $(OBJ_DIR)/ufm_cubin.h + +$(OBJ_DIR)/lal_ufm.o: $(ALL_H) lal_ufm.h lal_ufm.cpp $(OBJ_DIR)/ufm_cubin.h $(OBJ_DIR)/lal_base_atomic.o + $(CUDR) -o $@ -c lal_ufm.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_ufm_ext.o: $(ALL_H) lal_ufm.h lal_ufm_ext.cpp lal_base_atomic.h + $(CUDR) -o $@ -c lal_ufm_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_dpd.o: $(ALL_H) lal_dpd.h lal_dpd.cpp $(OBJ_DIR)/dpd_cubin.h $(OBJ_DIR)/lal_base_dpd.o + $(CUDR) -o $@ -c lal_dpd.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_dpd_ext.o: $(ALL_H) lal_dpd.h lal_dpd_ext.cpp lal_base_dpd.h + $(CUDR) -o $@ -c lal_dpd_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/tersoff.cubin: lal_tersoff.cu lal_precision.h lal_tersoff_extra.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_tersoff.cu + +$(OBJ_DIR)/tersoff_cubin.h: $(OBJ_DIR)/tersoff.cubin $(OBJ_DIR)/tersoff.cubin + $(BIN2C) -c -n tersoff $(OBJ_DIR)/tersoff.cubin > $(OBJ_DIR)/tersoff_cubin.h + +$(OBJ_DIR)/lal_tersoff.o: $(ALL_H) lal_tersoff.h lal_tersoff.cpp $(OBJ_DIR)/tersoff_cubin.h $(OBJ_DIR)/lal_base_three.o + $(CUDR) -o $@ -c lal_tersoff.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_tersoff_ext.o: $(ALL_H) lal_tersoff.h lal_tersoff_ext.cpp lal_base_three.h + $(CUDR) -o $@ -c lal_tersoff_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/tersoff_zbl.cubin: lal_tersoff_zbl.cu lal_precision.h lal_tersoff_zbl_extra.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_tersoff_zbl.cu + +$(OBJ_DIR)/tersoff_zbl_cubin.h: $(OBJ_DIR)/tersoff_zbl.cubin $(OBJ_DIR)/tersoff_zbl.cubin + $(BIN2C) -c -n tersoff_zbl $(OBJ_DIR)/tersoff_zbl.cubin > $(OBJ_DIR)/tersoff_zbl_cubin.h + +$(OBJ_DIR)/lal_tersoff_zbl.o: $(ALL_H) lal_tersoff_zbl.h lal_tersoff_zbl.cpp $(OBJ_DIR)/tersoff_zbl_cubin.h $(OBJ_DIR)/lal_base_three.o + $(CUDR) -o $@ -c lal_tersoff_zbl.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_tersoff_zbl_ext.o: $(ALL_H) lal_tersoff_zbl.h lal_tersoff_zbl_ext.cpp lal_base_three.h + $(CUDR) -o $@ -c lal_tersoff_zbl_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/tersoff_mod.cubin: lal_tersoff_mod.cu lal_precision.h lal_tersoff_mod_extra.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_tersoff_mod.cu + +$(OBJ_DIR)/tersoff_mod_cubin.h: $(OBJ_DIR)/tersoff_mod.cubin $(OBJ_DIR)/tersoff_mod.cubin + $(BIN2C) -c -n tersoff_mod $(OBJ_DIR)/tersoff_mod.cubin > $(OBJ_DIR)/tersoff_mod_cubin.h + +$(OBJ_DIR)/lal_tersoff_mod.o: $(ALL_H) lal_tersoff_mod.h lal_tersoff_mod.cpp $(OBJ_DIR)/tersoff_mod_cubin.h $(OBJ_DIR)/lal_base_three.o + $(CUDR) -o $@ -c lal_tersoff_mod.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_tersoff_mod_ext.o: $(ALL_H) lal_tersoff_mod.h lal_tersoff_mod_ext.cpp lal_base_three.h + $(CUDR) -o $@ -c lal_tersoff_mod_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/coul.cubin: lal_coul.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_coul.cu + +$(OBJ_DIR)/coul_cubin.h: $(OBJ_DIR)/coul.cubin $(OBJ_DIR)/coul.cubin + $(BIN2C) -c -n coul $(OBJ_DIR)/coul.cubin > $(OBJ_DIR)/coul_cubin.h + +$(OBJ_DIR)/lal_coul.o: $(ALL_H) lal_coul.h lal_coul.cpp $(OBJ_DIR)/coul_cubin.h $(OBJ_DIR)/lal_base_charge.o + $(CUDR) -o $@ -c lal_coul.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_coul_ext.o: $(ALL_H) lal_coul.h lal_coul_ext.cpp lal_base_charge.h + $(CUDR) -o $@ -c lal_coul_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/coul_debye.cubin: lal_coul_debye.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_coul_debye.cu + +$(OBJ_DIR)/coul_debye_cubin.h: $(OBJ_DIR)/coul_debye.cubin $(OBJ_DIR)/coul_debye.cubin + $(BIN2C) -c -n coul_debye $(OBJ_DIR)/coul_debye.cubin > $(OBJ_DIR)/coul_debye_cubin.h + +$(OBJ_DIR)/lal_coul_debye.o: $(ALL_H) lal_coul_debye.h lal_coul_debye.cpp $(OBJ_DIR)/coul_debye_cubin.h $(OBJ_DIR)/lal_base_charge.o + $(CUDR) -o $@ -c lal_coul_debye.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_coul_debye_ext.o: $(ALL_H) lal_coul_debye.h lal_coul_debye_ext.cpp lal_base_charge.h + $(CUDR) -o $@ -c lal_coul_debye_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/zbl.cubin: lal_zbl.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_zbl.cu + +$(OBJ_DIR)/zbl_cubin.h: $(OBJ_DIR)/zbl.cubin $(OBJ_DIR)/zbl.cubin + $(BIN2C) -c -n zbl $(OBJ_DIR)/zbl.cubin > $(OBJ_DIR)/zbl_cubin.h + +$(OBJ_DIR)/lal_zbl.o: $(ALL_H) lal_zbl.h lal_zbl.cpp $(OBJ_DIR)/zbl_cubin.h $(OBJ_DIR)/lal_base_atomic.o + $(CUDR) -o $@ -c lal_zbl.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_zbl_ext.o: $(ALL_H) lal_zbl.h lal_zbl_ext.cpp lal_base_atomic.h + $(CUDR) -o $@ -c lal_zbl_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lj_cubic.cubin: lal_lj_cubic.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_lj_cubic.cu + +$(OBJ_DIR)/lj_cubic_cubin.h: $(OBJ_DIR)/lj_cubic.cubin $(OBJ_DIR)/lj_cubic.cubin + $(BIN2C) -c -n lj_cubic $(OBJ_DIR)/lj_cubic.cubin > $(OBJ_DIR)/lj_cubic_cubin.h + +$(OBJ_DIR)/lal_lj_cubic.o: $(ALL_H) lal_lj_cubic.h lal_lj_cubic.cpp $(OBJ_DIR)/lj_cubic_cubin.h $(OBJ_DIR)/lal_base_atomic.o + $(CUDR) -o $@ -c lal_lj_cubic.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_lj_cubic_ext.o: $(ALL_H) lal_lj_cubic.h lal_lj_cubic_ext.cpp lal_base_atomic.h + $(CUDR) -o $@ -c lal_lj_cubic_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/dipole_long_lj.cubin: lal_dipole_long_lj.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_dipole_long_lj.cu + +$(OBJ_DIR)/dipole_long_lj_cubin.h: $(OBJ_DIR)/dipole_long_lj.cubin $(OBJ_DIR)/dipole_long_lj.cubin + $(BIN2C) -c -n dipole_long_lj $(OBJ_DIR)/dipole_long_lj.cubin > $(OBJ_DIR)/dipole_long_lj_cubin.h + +$(OBJ_DIR)/lal_dipole_long_lj.o: $(ALL_H) lal_dipole_long_lj.h lal_dipole_long_lj.cpp $(OBJ_DIR)/dipole_long_lj_cubin.h $(OBJ_DIR)/lal_base_dipole.o + $(CUDR) -o $@ -c lal_dipole_long_lj.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_dipole_long_lj_ext.o: $(ALL_H) lal_dipole_long_lj.h lal_dipole_long_lj_ext.cpp lal_base_dipole.h + $(CUDR) -o $@ -c lal_dipole_long_lj_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lj_expand_coul_long.cubin: lal_lj_expand_coul_long.cu lal_precision.h lal_preprocessor.h + $(CUDA) --fatbin -DNV_KERNEL -o $@ lal_lj_expand_coul_long.cu + +$(OBJ_DIR)/lj_expand_coul_long_cubin.h: $(OBJ_DIR)/lj_expand_coul_long.cubin $(OBJ_DIR)/lj_expand_coul_long.cubin + $(BIN2C) -c -n lj_expand_coul_long $(OBJ_DIR)/lj_expand_coul_long.cubin > $(OBJ_DIR)/lj_expand_coul_long_cubin.h + +$(OBJ_DIR)/lal_lj_expand_coul_long.o: $(ALL_H) lal_lj_expand_coul_long.h lal_lj_expand_coul_long.cpp $(OBJ_DIR)/lj_expand_coul_long_cubin.h $(OBJ_DIR)/lal_base_charge.o + $(CUDR) -o $@ -c lal_lj_expand_coul_long.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_lj_expand_coul_long_ext.o: $(ALL_H) lal_lj_expand_coul_long.h lal_lj_expand_coul_long_ext.cpp lal_base_charge.h + $(CUDR) -o $@ -c lal_lj_expand_coul_long_ext.cpp -I$(OBJ_DIR) + +$(BIN_DIR)/nvc_get_devices: ./geryon/ucl_get_devices.cpp $(NVD_H) + $(CUDR) -o $@ ./geryon/ucl_get_devices.cpp -DUCL_CUDADR $(CUDA_LIB) -lcuda + +$(GPU_LIB): $(OBJS) $(CUDPP) + $(AR) -crusv $(GPU_LIB) $(OBJS) $(CUDPP) + @cp $(EXTRAMAKE) Makefile.lammps + +clean: + -rm -f $(EXECS) $(GPU_LIB) $(OBJS) $(CUDPP) $(CBNS) *.linkinfo + +veryclean: clean + -rm -rf *~ *.linkinfo + +cleanlib: + -rm -f $(EXECS) $(GPU_LIB) $(OBJS) $(CBNS) *.linkinfo diff --git a/lib/gpu/Opencl.makefile b/lib/gpu/Opencl.makefile index bb071514474655a9fb36a73fab05635fa9ffdd99..3e73e6256c7b998c449a60e32e2a3974f636b227 100644 --- a/lib/gpu/Opencl.makefile +++ b/lib/gpu/Opencl.makefile @@ -66,7 +66,9 @@ OBJS = $(OBJ_DIR)/lal_atom.o $(OBJ_DIR)/lal_answer.o \ $(OBJ_DIR)/lal_coul_debye.o $(OBJ_DIR)/lal_coul_debye_ext.o \ $(OBJ_DIR)/lal_zbl.o $(OBJ_DIR)/lal_zbl_ext.o \ $(OBJ_DIR)/lal_lj_cubic.o $(OBJ_DIR)/lal_lj_cubic_ext.o \ - $(OBJ_DIR)/lal_ufm.o $(OBJ_DIR)/lal_ufm_ext.o + $(OBJ_DIR)/lal_ufm.o $(OBJ_DIR)/lal_ufm_ext.o \ + $(OBJ_DIR)/lal_dipole_long_lj.o $(OBJ_DIR)/lal_dipole_long_lj_ext.o \ + $(OBJ_DIR)/lal_lj_expand_coul_long.o $(OBJ_DIR)/lal_lj_expand_coul_long_ext.o KERS = $(OBJ_DIR)/device_cl.h $(OBJ_DIR)/atom_cl.h \ $(OBJ_DIR)/neighbor_cpu_cl.h $(OBJ_DIR)/pppm_cl.h \ @@ -95,7 +97,8 @@ KERS = $(OBJ_DIR)/device_cl.h $(OBJ_DIR)/atom_cl.h \ $(OBJ_DIR)/tersoff_mod_cl.h $(OBJ_DIR)/coul_cl.h \ $(OBJ_DIR)/coul_debye_cl.h $(OBJ_DIR)/zbl_cl.h \ $(OBJ_DIR)/lj_cubic_cl.h $(OBJ_DIR)/vashishta_cl.h \ - $(OBJ_DIR)/ufm_cl.h + $(OBJ_DIR)/ufm_cl.h $(OBJ_DIR)/dipole_long_lj_cl.h \ + $(OBJ_DIR)/lj_expand_coul_long_cl.h OCL_EXECS = $(BIN_DIR)/ocl_get_devices @@ -588,7 +591,25 @@ $(OBJ_DIR)/lal_ufm.o: $(ALL_H) lal_ufm.h lal_ufm.cpp $(OBJ_DIR)/ufm_cl.h $(OBJ_ $(OBJ_DIR)/lal_ufm_ext.o: $(ALL_H) lal_ufm.h lal_ufm_ext.cpp lal_base_atomic.h $(OCL) -o $@ -c lal_ufm_ext.cpp -I$(OBJ_DIR) -$(BIN_DIR)/ocl_get_devices: ./geryon/ucl_get_devices.cpp +$(OBJ_DIR)/dipole_long_lj_cl.h: lal_dipole_long_lj.cu $(PRE1_H) + $(BSH) ./geryon/file_to_cstr.sh dipole_long_lj $(PRE1_H) lal_dipole_long_lj.cu $(OBJ_DIR)/dipole_long_lj_cl.h; + +$(OBJ_DIR)/lal_dipole_long_lj.o: $(ALL_H) lal_dipole_long_lj.h lal_dipole_long_lj.cpp $(OBJ_DIR)/dipole_long_lj_cl.h $(OBJ_DIR)/lj_expand_coul_long_cl.h $(OBJ_DIR)/lal_base_charge.o + $(OCL) -o $@ -c lal_dipole_long_lj.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_dipole_long_lj_ext.o: $(ALL_H) lal_dipole_long_lj.h lal_dipole_long_lj_ext.cpp lal_base_dipole.h + $(OCL) -o $@ -c lal_dipole_long_lj_ext.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lj_expand_coul_long_cl.h: lal_lj_expand_coul_long.cu $(PRE1_H) + $(BSH) ./geryon/file_to_cstr.sh lj_expand_coul_long $(PRE1_H) lal_lj_expand_coul_long.cu $(OBJ_DIR)/lj_expand_coul_long_cl.h; + +$(OBJ_DIR)/lal_lj_expand_coul_long.o: $(ALL_H) lal_lj_expand_coul_long.h lal_lj_expand_coul_long.cpp $(OBJ_DIR)/lj_expand_coul_long_cl.h $(OBJ_DIR)/lj_expand_coul_long_cl.h $(OBJ_DIR)/lal_base_charge.o + $(OCL) -o $@ -c lal_lj_expand_coul_long.cpp -I$(OBJ_DIR) + +$(OBJ_DIR)/lal_lj_expand_coul_long_ext.o: $(ALL_H) lal_lj_expand_coul_long.h lal_lj_expand_coul_long_ext.cpp lal_base_charge.h + $(OCL) -o $@ -c lal_lj_expand_coul_long_ext.cpp -I$(OBJ_DIR) + +$(BIN_DIR)/ocl_get_devices: ./geryon/ucl_get_devices.cpp $(OCL_H) $(OCL) -o $@ ./geryon/ucl_get_devices.cpp -DUCL_OPENCL $(OCL_LINK) $(OCL_LIB): $(OBJS) $(PTXS) diff --git a/lib/gpu/geryon/ocl_device.h b/lib/gpu/geryon/ocl_device.h index 584d04e616521056d11372e4b2e48291e2df04ab..2b2367545e837427527cad44843a5cb31fd09d30 100644 --- a/lib/gpu/geryon/ocl_device.h +++ b/lib/gpu/geryon/ocl_device.h @@ -280,6 +280,9 @@ class UCL_Device { /// Return the OpenCL type for the device inline cl_device_id & cl_device() { return _cl_device; } + /// Select the platform that has accelerators + inline void set_platform_accelerator(int pid=-1); + private: int _num_platforms; // Number of platforms int _platform; // UCL_Device ID for current platform @@ -311,8 +314,8 @@ UCL_Device::UCL_Device() { return; } else _num_platforms=static_cast(nplatforms); - - set_platform(0); + // note that platform 0 may not necessarily be associated with accelerators + set_platform_accelerator(); } UCL_Device::~UCL_Device() { @@ -320,6 +323,7 @@ UCL_Device::~UCL_Device() { } void UCL_Device::clear() { + _properties.clear(); if (_device>-1) { for (size_t i=0; i<_cq.size(); i++) { CL_DESTRUCT_CALL(clReleaseCommandQueue(_cq.back())); @@ -529,75 +533,105 @@ int UCL_Device::set(int num) { return create_context(); } -// List all devices along with all properties +// List all devices from all platforms along with all properties void UCL_Device::print_all(std::ostream &out) { - if (num_devices() == 0) - out << "There is no device supporting OpenCL\n"; - for (int i=0; i +using namespace LAMMPS_AL; +#define DipoleLongLJT DipoleLongLJ + +extern Device device; + +template +DipoleLongLJT::DipoleLongLJ() : BaseDipole(), + _allocated(false) { +} + +template +DipoleLongLJT::~DipoleLongLJ() { + clear(); +} + +template +int DipoleLongLJT::bytes_per_atom(const int max_nbors) const { + return this->bytes_per_atom_atomic(max_nbors); +} + +template +int DipoleLongLJT::init(const int ntypes, + double **host_cutsq, double **host_lj1, + double **host_lj2, double **host_lj3, + double **host_lj4, double **host_offset, + double *host_special_lj, const int nlocal, + const int nall, const int max_nbors, + const int maxspecial, const double cell_size, + const double gpu_split, FILE *_screen, + double **host_cut_ljsq, const double host_cut_coulsq, + double *host_special_coul, const double qqrd2e, + const double g_ewald) { + int success; + success=this->init_atomic(nlocal,nall,max_nbors,maxspecial,cell_size,gpu_split, + _screen,dipole_long_lj,"k_dipole_long_lj"); + if (success!=0) + return success; + + // If atom type constants fit in shared memory use fast kernel + int lj_types=ntypes; + shared_types=false; + int max_shared_types=this->device->max_shared_types(); + if (lj_types<=max_shared_types && this->_block_size>=max_shared_types) { + lj_types=max_shared_types; + shared_types=true; + } + _lj_types=lj_types; + + // Allocate a host write buffer for data initialization + UCL_H_Vec host_write(lj_types*lj_types*32,*(this->ucl_device), + UCL_WRITE_ONLY); + + for (int i=0; iucl_device),UCL_READ_ONLY); + this->atom->type_pack4(ntypes,lj_types,lj1,host_write,host_lj1,host_lj2, + host_cut_ljsq); + + lj3.alloc(lj_types*lj_types,*(this->ucl_device),UCL_READ_ONLY); + this->atom->type_pack4(ntypes,lj_types,lj3,host_write,host_lj3,host_lj4, + host_offset); + + cutsq.alloc(lj_types*lj_types,*(this->ucl_device),UCL_READ_ONLY); + this->atom->type_pack1(ntypes,lj_types,cutsq,host_write,host_cutsq); + + sp_lj.alloc(8,*(this->ucl_device),UCL_READ_ONLY); + for (int i=0; i<4; i++) { + host_write[i]=host_special_lj[i]; + host_write[i+4]=host_special_coul[i]; + } + ucl_copy(sp_lj,host_write,8,false); + + _cut_coulsq=host_cut_coulsq; + _qqrd2e=qqrd2e; + _g_ewald=g_ewald; + + _allocated=true; + this->_max_bytes=lj1.row_bytes()+lj3.row_bytes()+cutsq.row_bytes()+ + sp_lj.row_bytes(); + return 0; +} + +template +void DipoleLongLJT::clear() { + if (!_allocated) + return; + _allocated=false; + + lj1.clear(); + lj3.clear(); + cutsq.clear(); + sp_lj.clear(); + this->clear_atomic(); +} + +template +double DipoleLongLJT::host_memory_usage() const { + return this->host_memory_usage_atomic()+sizeof(DipoleLongLJ); +} + +// --------------------------------------------------------------------------- +// Calculate energies, forces, and torques +// --------------------------------------------------------------------------- +template +void DipoleLongLJT::loop(const bool _eflag, const bool _vflag) { + // Compute the block size and grid size to keep all cores busy + const int BX=this->block_size(); + int eflag, vflag; + if (_eflag) + eflag=1; + else + eflag=0; + + if (_vflag) + vflag=1; + else + vflag=0; + + int GX=static_cast(ceil(static_cast(this->ans->inum())/ + (BX/this->_threads_per_atom))); + + int ainum=this->ans->inum(); + int nbor_pitch=this->nbor->nbor_pitch(); + this->time_pair.start(); + if (shared_types) { + this->k_pair_fast.set_size(GX,BX); + this->k_pair_fast.run(&this->atom->x, &lj1, &lj3, &sp_lj, + &this->nbor->dev_nbor, + &this->_nbor_data->begin(), + &this->ans->force, &this->ans->engv, &eflag, &vflag, + &ainum, &nbor_pitch, &this->atom->q, + &this->atom->quat, &cutsq, &_cut_coulsq, + &_qqrd2e, &_g_ewald, &this->_threads_per_atom); + } else { + this->k_pair.set_size(GX,BX); + this->k_pair.run(&this->atom->x, &lj1, &lj3, + &_lj_types, &sp_lj, &this->nbor->dev_nbor, + &this->_nbor_data->begin(), &this->ans->force, + &this->ans->engv, &eflag, &vflag, &ainum, + &nbor_pitch, &this->atom->q, + &this->atom->quat, &cutsq, &_cut_coulsq, + &_qqrd2e, &_g_ewald, &this->_threads_per_atom); + } + this->time_pair.stop(); +} + +template class DipoleLongLJ; diff --git a/lib/gpu/lal_dipole_long_lj.cu b/lib/gpu/lal_dipole_long_lj.cu new file mode 100644 index 0000000000000000000000000000000000000000..f888dece9b513001b4af9309ebc76e7b30f060d2 --- /dev/null +++ b/lib/gpu/lal_dipole_long_lj.cu @@ -0,0 +1,640 @@ +// ************************************************************************** +// dipole_lj.cu +// ------------------- +// Trung Dac Nguyen (ORNL) +// +// Device code for acceleration of the dipole/cut pair style +// +// __________________________________________________________________________ +// This file is part of the LAMMPS Accelerator Library (LAMMPS_AL) +// __________________________________________________________________________ +// +// begin : +// email : nguyentd@ornl.gov +// ***************************************************************************/ + +#ifdef NV_KERNEL +#include "lal_aux_fun1.h" +#ifndef _DOUBLE_DOUBLE +texture pos_tex; +texture q_tex; +texture mu_tex; +#else +texture pos_tex; +texture q_tex; +texture mu_tex; +#endif + +#else +#define pos_tex x_ +#define q_tex q_ +#define mu_tex mu_ +#endif + +#if (ARCH < 300) + +#define store_answers_tq(f, tor, energy, ecoul, virial, ii, inum, tid, \ + t_per_atom, offset, eflag, vflag, ans, engv) \ + if (t_per_atom>1) { \ + __local acctyp red_acc[8][BLOCK_PAIR]; \ + red_acc[0][tid]=f.x; \ + red_acc[1][tid]=f.y; \ + red_acc[2][tid]=f.z; \ + red_acc[3][tid]=tor.x; \ + red_acc[4][tid]=tor.y; \ + red_acc[5][tid]=tor.z; \ + for (unsigned int s=t_per_atom/2; s>0; s>>=1) { \ + if (offset < s) { \ + for (int r=0; r<6; r++) \ + red_acc[r][tid] += red_acc[r][tid+s]; \ + } \ + } \ + f.x=red_acc[0][tid]; \ + f.y=red_acc[1][tid]; \ + f.z=red_acc[2][tid]; \ + tor.x=red_acc[3][tid]; \ + tor.y=red_acc[4][tid]; \ + tor.z=red_acc[5][tid]; \ + if (eflag>0 || vflag>0) { \ + for (int r=0; r<6; r++) \ + red_acc[r][tid]=virial[r]; \ + red_acc[6][tid]=energy; \ + red_acc[7][tid]=ecoul; \ + for (unsigned int s=t_per_atom/2; s>0; s>>=1) { \ + if (offset < s) { \ + for (int r=0; r<8; r++) \ + red_acc[r][tid] += red_acc[r][tid+s]; \ + } \ + } \ + for (int r=0; r<6; r++) \ + virial[r]=red_acc[r][tid]; \ + energy=red_acc[6][tid]; \ + ecoul=red_acc[7][tid]; \ + } \ + } \ + if (offset==0) { \ + int ei=ii; \ + if (eflag>0) { \ + engv[ei]=energy*(acctyp)0.5; \ + ei+=inum; \ + engv[ei]=e_coul*(acctyp)0.5; \ + ei+=inum; \ + } \ + if (vflag>0) { \ + for (int i=0; i<6; i++) { \ + engv[ei]=virial[i]*(acctyp)0.5; \ + ei+=inum; \ + } \ + } \ + ans[ii]=f; \ + ans[ii+inum]=tor; \ + } + +#else + +#define store_answers_tq(f, tor, energy, e_coul, virial, ii, inum, tid, \ + t_per_atom, offset, eflag, vflag, ans, engv) \ + if (t_per_atom>1) { \ + for (unsigned int s=t_per_atom/2; s>0; s>>=1) { \ + f.x += shfl_xor(f.x, s, t_per_atom); \ + f.y += shfl_xor(f.y, s, t_per_atom); \ + f.z += shfl_xor(f.z, s, t_per_atom); \ + tor.x += shfl_xor(tor.x, s, t_per_atom); \ + tor.y += shfl_xor(tor.y, s, t_per_atom); \ + tor.z += shfl_xor(tor.z, s, t_per_atom); \ + energy += shfl_xor(energy, s, t_per_atom); \ + e_coul += shfl_xor(e_coul, s, t_per_atom); \ + } \ + if (vflag>0) { \ + for (unsigned int s=t_per_atom/2; s>0; s>>=1) { \ + for (int r=0; r<6; r++) \ + virial[r] += shfl_xor(virial[r], s, t_per_atom); \ + } \ + } \ + } \ + if (offset==0) { \ + int ei=ii; \ + if (eflag>0) { \ + engv[ei]=energy*(acctyp)0.5; \ + ei+=inum; \ + engv[ei]=e_coul*(acctyp)0.5; \ + ei+=inum; \ + } \ + if (vflag>0) { \ + for (int i=0; i<6; i++) { \ + engv[ei]=virial[i]*(acctyp)0.5; \ + ei+=inum; \ + } \ + } \ + ans[ii]=f; \ + ans[ii+inum]=tor; \ + } + +#endif + +#define MY_PIS (acctyp)1.77245385090551602729 + +__kernel void k_dipole_long_lj(const __global numtyp4 *restrict x_, + const __global numtyp4 *restrict lj1, + const __global numtyp4 *restrict lj3, + const int lj_types, + const __global numtyp *restrict sp_lj_in, + const __global int *dev_nbor, + const __global int *dev_packed, + __global acctyp4 *restrict ans, + __global acctyp *restrict engv, + const int eflag, const int vflag, const int inum, + const int nbor_pitch, + const __global numtyp *restrict q_, + const __global numtyp4 *restrict mu_, + const __global numtyp *restrict cutsq, + const numtyp cut_coulsq, const numtyp qqrd2e, + const numtyp g_ewald, const int t_per_atom) { + int tid, ii, offset; + atom_info(t_per_atom,ii,tid,offset); + + __local numtyp sp_lj[8]; + sp_lj[0]=sp_lj_in[0]; + sp_lj[1]=sp_lj_in[1]; + sp_lj[2]=sp_lj_in[2]; + sp_lj[3]=sp_lj_in[3]; + sp_lj[4]=sp_lj_in[4]; + sp_lj[5]=sp_lj_in[5]; + sp_lj[6]=sp_lj_in[6]; + sp_lj[7]=sp_lj_in[7]; + + acctyp energy=(acctyp)0; + acctyp e_coul=(acctyp)0; + acctyp4 f; + f.x=(acctyp)0; f.y=(acctyp)0; f.z=(acctyp)0; + acctyp4 tor; + tor.x=(acctyp)0; + tor.y=(acctyp)0; + tor.z=(acctyp)0; + acctyp virial[6]; + for (int i=0; i<6; i++) + virial[i]=(acctyp)0; + + numtyp pre1 = numtyp(2.0) * g_ewald / MY_PIS; + numtyp pre2 = numtyp(4.0) * (g_ewald*g_ewald*g_ewald) / MY_PIS; + numtyp pre3 = numtyp(8.0) * (g_ewald*g_ewald*g_ewald*g_ewald*g_ewald) / MY_PIS; + + if (ii (numtyp)0.0) { + b0 = _erfc * rinv; + b1 = (b0 + pre1*expm2) * r2inv; + b2 = ((numtyp)3.0*b1 + pre2*expm2) * r2inv; + b3 = ((numtyp)5.0*b2 + pre3*expm2) * r2inv; + + g0b1_g1b2_g2b3 = g0*b1 + g1*b2 + g2*b3; + fdx = delx * g0b1_g1b2_g2b3 - + b1 * (qtmp*muj.x - qj*mui.x) + + b2 * (pjdotr*mui.x + pidotr*muj.x); + fdy = dely * g0b1_g1b2_g2b3 - + b1 * (qtmp*muj.y - qj*mui.y) + + b2 * (pjdotr*mui.y + pidotr*muj.y); + fdz = delz * g0b1_g1b2_g2b3 - + b1 * (qtmp*muj.z - qj*mui.z) + + b2 * (pjdotr*mui.z + pidotr*muj.z); + + zdix = delx * (qj*b1 + b2*pjdotr) - b1*muj.x; + zdiy = dely * (qj*b1 + b2*pjdotr) - b1*muj.y; + zdiz = delz * (qj*b1 + b2*pjdotr) - b1*muj.z; + zdjx = delx * (-qtmp*b1 + b2*pidotr) - b1*mui.x; + zdjy = dely * (-qtmp*b1 + b2*pidotr) - b1*mui.y; + zdjz = delz * (-qtmp*b1 + b2*pidotr) - b1*mui.z; + + if (factor_coul < (numtyp)1.0) { + fdx *= factor_coul; + fdy *= factor_coul; + fdz *= factor_coul; + zdix *= factor_coul; + zdiy *= factor_coul; + zdiz *= factor_coul; + zdjx *= factor_coul; + zdjy *= factor_coul; + zdjz *= factor_coul; + } + } else { + fdx = fdy = fdz = (numtyp)0.0; + zdix = zdiy = zdiz = (numtyp)0.0; + zdjx = zdjy = zdjz = (numtyp)0.0; + } + + if (factor_coul < (numtyp)1.0) { + d0 = (_erfc - (numtyp)1.0) * rinv; + d1 = (d0 + pre1*expm2) * r2inv; + d2 = ((numtyp)3.0*d1 + pre2*expm2) * r2inv; + d3 = ((numtyp)5.0*d2 + pre3*expm2) * r2inv; + + g0d1_g1d2_g2d3 = g0*d1 + g1*d2 + g2*d3; + fax = delx * g0d1_g1d2_g2d3 - + d1 * (qtmp*muj.x - qj*mui.x) + + d2 * (pjdotr*mui.x + pidotr*muj.x); + fay = dely * g0d1_g1d2_g2d3 - + d1 * (qtmp*muj.y - qj*mui.y) + + d2 * (pjdotr*mui.y + pidotr*muj.y); + faz = delz * g0d1_g1d2_g2d3 - + d1 * (qtmp*muj.z - qj*mui.z) + + d2 * (pjdotr*mui.z + pidotr*muj.z); + + zaix = delx * (qj*d1 + d2*pjdotr) - d1*muj.x; + zaiy = dely * (qj*d1 + d2*pjdotr) - d1*muj.y; + zaiz = delz * (qj*d1 + d2*pjdotr) - d1*muj.z; + zajx = delx * (-qtmp*d1 + d2*pidotr) - d1*mui.x; + zajy = dely * (-qtmp*d1 + d2*pidotr) - d1*mui.y; + zajz = delz * (-qtmp*d1 + d2*pidotr) - d1*mui.z; + + if (factor_coul > (numtyp)0.0) { + facm1 = (numtyp)1.0 - factor_coul; + fax *= facm1; + fay *= facm1; + faz *= facm1; + zaix *= facm1; + zaiy *= facm1; + zaiz *= facm1; + zajx *= facm1; + zajy *= facm1; + zajz *= facm1; + } + } else { + fax = fay = faz = (numtyp)0.0; + zaix = zaiy = zaiz = (numtyp)0.0; + zajx = zajy = zajz = (numtyp)0.0; + } + + forcecoul.x = fdx + fax; + forcecoul.y = fdy + fay; + forcecoul.z = fdz + faz; + + ticoul.x = mui.y*(zdiz + zaiz) - mui.z*(zdiy + zaiy); + ticoul.y = mui.z*(zdix + zaix) - mui.x*(zdiz + zaiz); + ticoul.z = mui.x*(zdiy + zaiy) - mui.y*(zdix + zaix); + + } else { + forcecoul.x = forcecoul.y = forcecoul.z = (numtyp)0.0; + ticoul.x = ticoul.y = ticoul.z = (numtyp)0.0; + } + + force.x = qqrd2e*forcecoul.x + delx*force_lj; + force.y = qqrd2e*forcecoul.y + dely*force_lj; + force.z = qqrd2e*forcecoul.z + delz*force_lj; + f.x+=force.x; + f.y+=force.y; + f.z+=force.z; + tor.x+=qqrd2e*ticoul.x; + tor.y+=qqrd2e*ticoul.y; + tor.z+=qqrd2e*ticoul.z; + + if (eflag>0) { + acctyp e = (acctyp)0.0; + if (rsq < cut_coulsq && factor_coul > (numtyp)0.0) { + e = qqrd2e*(b0*g0 + b1*g1 + b2*g2); + if (factor_coul < (numtyp)1.0) { + e_coul *= factor_coul; + e_coul += ((numtyp)1.0-factor_coul) * qqrd2e * (d0*g0 + d1*g1 + d2*g2); + } + } else e = (acctyp)0.0; + e_coul += e; + + if (rsq < lj1[mtype].z) { + e=r6inv*(lj3[mtype].x*r6inv-lj3[mtype].y); + energy+=factor_lj*(e-lj3[mtype].z); + } + } + if (vflag>0) { + virial[0] += delx*force.x; + virial[1] += dely*force.y; + virial[2] += delz*force.z; + virial[3] += delx*force.y; + virial[4] += delx*force.z; + virial[5] += dely*force.z; + } + } + + } // for nbor + store_answers_tq(f,tor,energy,e_coul,virial,ii,inum,tid,t_per_atom,offset,eflag, + vflag,ans,engv); + } // if ii +} + +__kernel void k_dipole_long_lj_fast(const __global numtyp4 *restrict x_, + const __global numtyp4 *restrict lj1_in, + const __global numtyp4 *restrict lj3_in, + const __global numtyp *restrict sp_lj_in, + const __global int *dev_nbor, + const __global int *dev_packed, + __global acctyp4 *restrict ans, + __global acctyp *restrict engv, + const int eflag, const int vflag, const int inum, + const int nbor_pitch, + const __global numtyp *restrict q_, + const __global numtyp4 *restrict mu_, + const __global numtyp *restrict _cutsq, + const numtyp cut_coulsq, const numtyp qqrd2e, + const numtyp g_ewald, const int t_per_atom) { + int tid, ii, offset; + atom_info(t_per_atom,ii,tid,offset); + + __local numtyp4 lj1[MAX_SHARED_TYPES*MAX_SHARED_TYPES]; + __local numtyp4 lj3[MAX_SHARED_TYPES*MAX_SHARED_TYPES]; + __local numtyp cutsq[MAX_SHARED_TYPES*MAX_SHARED_TYPES]; + __local numtyp sp_lj[8]; + if (tid<8) + sp_lj[tid]=sp_lj_in[tid]; + if (tid0) + lj3[tid]=lj3_in[tid]; + } + + acctyp energy=(acctyp)0; + acctyp e_coul=(acctyp)0; + acctyp4 f; + f.x=(acctyp)0; f.y=(acctyp)0; f.z=(acctyp)0; + acctyp4 tor; + tor.x=(acctyp)0; + tor.y=(acctyp)0; + tor.z=(acctyp)0; + acctyp virial[6]; + for (int i=0; i<6; i++) + virial[i]=(acctyp)0; + + __syncthreads(); + + numtyp pre1 = numtyp(2.0) * g_ewald / MY_PIS; + numtyp pre2 = numtyp(4.0) * (g_ewald*g_ewald*g_ewald) / MY_PIS; + numtyp pre3 = numtyp(8.0) * (g_ewald*g_ewald*g_ewald*g_ewald*g_ewald) / MY_PIS; + + if (ii (numtyp)0.0) { + b0 = _erfc * rinv; + b1 = (b0 + pre1*expm2) * r2inv; + b2 = ((numtyp)3.0*b1 + pre2*expm2) * r2inv; + b3 = ((numtyp)5.0*b2 + pre3*expm2) * r2inv; + + g0b1_g1b2_g2b3 = g0*b1 + g1*b2 + g2*b3; + fdx = delx * g0b1_g1b2_g2b3 - + b1 * (qtmp*muj.x - qj*mui.x) + + b2 * (pjdotr*mui.x + pidotr*muj.x); + fdy = dely * g0b1_g1b2_g2b3 - + b1 * (qtmp*muj.y - qj*mui.y) + + b2 * (pjdotr*mui.y + pidotr*muj.y); + fdz = delz * g0b1_g1b2_g2b3 - + b1 * (qtmp*muj.z - qj*mui.z) + + b2 * (pjdotr*mui.z + pidotr*muj.z); + + zdix = delx * (qj*b1 + b2*pjdotr) - b1*muj.x; + zdiy = dely * (qj*b1 + b2*pjdotr) - b1*muj.y; + zdiz = delz * (qj*b1 + b2*pjdotr) - b1*muj.z; + zdjx = delx * (-qtmp*b1 + b2*pidotr) - b1*mui.x; + zdjy = dely * (-qtmp*b1 + b2*pidotr) - b1*mui.y; + zdjz = delz * (-qtmp*b1 + b2*pidotr) - b1*mui.z; + + if (factor_coul < (numtyp)1.0) { + fdx *= factor_coul; + fdy *= factor_coul; + fdz *= factor_coul; + zdix *= factor_coul; + zdiy *= factor_coul; + zdiz *= factor_coul; + zdjx *= factor_coul; + zdjy *= factor_coul; + zdjz *= factor_coul; + } + } else { + fdx = fdy = fdz = (numtyp)0.0; + zdix = zdiy = zdiz = (numtyp)0.0; + zdjx = zdjy = zdjz = (numtyp)0.0; + } + + if (factor_coul < (numtyp)1.0) { + d0 = (_erfc - (numtyp)1.0) * rinv; + d1 = (d0 + pre1*expm2) * r2inv; + d2 = ((numtyp)3.0*d1 + pre2*expm2) * r2inv; + d3 = ((numtyp)5.0*d2 + pre3*expm2) * r2inv; + + g0d1_g1d2_g2d3 = g0*d1 + g1*d2 + g2*d3; + fax = delx * g0d1_g1d2_g2d3 - + d1 * (qtmp*muj.x - qj*mui.x) + + d2 * (pjdotr*mui.x + pidotr*muj.x); + fay = dely * g0d1_g1d2_g2d3 - + d1 * (qtmp*muj.y - qj*mui.y) + + d2 * (pjdotr*mui.y + pidotr*muj.y); + faz = delz * g0d1_g1d2_g2d3 - + d1 * (qtmp*muj.z - qj*mui.z) + + d2 * (pjdotr*mui.z + pidotr*muj.z); + + zaix = delx * (qj*d1 + d2*pjdotr) - d1*muj.x; + zaiy = dely * (qj*d1 + d2*pjdotr) - d1*muj.y; + zaiz = delz * (qj*d1 + d2*pjdotr) - d1*muj.z; + zajx = delx * (-qtmp*d1 + d2*pidotr) - d1*mui.x; + zajy = dely * (-qtmp*d1 + d2*pidotr) - d1*mui.y; + zajz = delz * (-qtmp*d1 + d2*pidotr) - d1*mui.z; + + if (factor_coul > (numtyp)0.0) { + facm1 = (numtyp)1.0 - factor_coul; + fax *= facm1; + fay *= facm1; + faz *= facm1; + zaix *= facm1; + zaiy *= facm1; + zaiz *= facm1; + zajx *= facm1; + zajy *= facm1; + zajz *= facm1; + } + } else { + fax = fay = faz = (numtyp)0.0; + zaix = zaiy = zaiz = (numtyp)0.0; + zajx = zajy = zajz = (numtyp)0.0; + } + + forcecoul.x = fdx + fax; + forcecoul.y = fdy + fay; + forcecoul.z = fdz + faz; + + ticoul.x = mui.y*(zdiz + zaiz) - mui.z*(zdiy + zaiy); + ticoul.y = mui.z*(zdix + zaix) - mui.x*(zdiz + zaiz); + ticoul.z = mui.x*(zdiy + zaiy) - mui.y*(zdix + zaix); + + } else { + forcecoul.x = forcecoul.y = forcecoul.z = (numtyp)0.0; + ticoul.x = ticoul.y = ticoul.z = (numtyp)0.0; + } + + force.x = qqrd2e*forcecoul.x + delx*force_lj; + force.y = qqrd2e*forcecoul.y + dely*force_lj; + force.z = qqrd2e*forcecoul.z + delz*force_lj; + f.x+=force.x; + f.y+=force.y; + f.z+=force.z; + tor.x+=qqrd2e*ticoul.x; + tor.y+=qqrd2e*ticoul.y; + tor.z+=qqrd2e*ticoul.z; + + if (eflag>0) { + acctyp e = (acctyp)0.0; + if (rsq < cut_coulsq && factor_coul > (numtyp)0.0) { + e = qqrd2e*(b0*g0 + b1*g1 + b2*g2); + if (factor_coul < (numtyp)1.0) { + e_coul *= factor_coul; + e_coul += ((numtyp)1.0-factor_coul) * qqrd2e * (d0*g0 + d1*g1 + d2*g2); + } + } else e = (acctyp)0.0; + e_coul += e; + + if (rsq < lj1[mtype].z) { + e=r6inv*(lj3[mtype].x*r6inv-lj3[mtype].y); + energy+=factor_lj*(e-lj3[mtype].z); + } + } + if (vflag>0) { + virial[0] += delx*force.x; + virial[1] += dely*force.y; + virial[2] += delz*force.z; + virial[3] += delx*force.y; + virial[4] += delx*force.z; + virial[5] += dely*force.z; + } + } + + } // for nbor + store_answers_tq(f,tor,energy,e_coul,virial,ii,inum,tid,t_per_atom,offset,eflag, + vflag,ans,engv); + } // if ii +} + diff --git a/lib/gpu/lal_dipole_long_lj.h b/lib/gpu/lal_dipole_long_lj.h new file mode 100644 index 0000000000000000000000000000000000000000..1381e243267bcc548d7eef6da409b29f3da36931 --- /dev/null +++ b/lib/gpu/lal_dipole_long_lj.h @@ -0,0 +1,85 @@ +/*************************************************************************** + dipole_long_lj.h + ------------------- + Trung Dac Nguyen (Northwestern) + + Class for acceleration of the lj/cut/dipole/long pair style. + + __________________________________________________________________________ + This file is part of the LAMMPS Accelerator Library (LAMMPS_AL) + __________________________________________________________________________ + + begin : + email : ndactrung@gmail.com + ***************************************************************************/ + +#ifndef LAL_DIPOLE_LONG_LJ_H +#define LAL_DIPOLE_LONG_LJ_H + +#include "lal_base_dipole.h" + +namespace LAMMPS_AL { + +template +class DipoleLongLJ : public BaseDipole { + public: + DipoleLongLJ(); + ~DipoleLongLJ(); + + /// Clear any previous data and set up for a new LAMMPS run + /** \param max_nbors initial number of rows in the neighbor matrix + * \param cell_size cutoff + skin + * \param gpu_split fraction of particles handled by device + * + * Returns: + * - 0 if successfull + * - -1 if fix gpu not found + * - -3 if there is an out of memory error + * - -4 if the GPU library was not compiled for GPU + * - -5 Double precision is not supported on card **/ + int init(const int ntypes, double **host_cutsq, double **host_lj1, + double **host_lj2, double **host_lj3, double **host_lj4, + double **host_offset, double *host_special_lj, + const int nlocal, const int nall, const int max_nbors, + const int maxspecial, const double cell_size, + const double gpu_split, FILE *screen, double **host_cut_ljsq, + const double host_cut_coulsq, double *host_special_coul, + const double qqrd2e, const double g_ewald); + + /// Clear all host and device data + /** \note This is called at the beginning of the init() routine **/ + void clear(); + + /// Returns memory usage on device per atom + int bytes_per_atom(const int max_nbors) const; + + /// Total host memory used by library for pair style + double host_memory_usage() const; + + // --------------------------- TYPE DATA -------------------------- + + /// lj1.x = lj1, lj1.y = lj2, lj1.z = cutsq_vdw, lj1.w = cutsq_coul + UCL_D_Vec lj1; + /// lj3.x = lj3, lj3.y = lj4, lj3.z = offset + UCL_D_Vec lj3; + /// cutsq + UCL_D_Vec cutsq; + /// Special LJ values [0-3] and Special Coul values [4-7] + UCL_D_Vec sp_lj; + + /// If atom type constants fit in shared memory, use fast kernels + bool shared_types; + + /// Number of atom types + int _lj_types; + + numtyp _cut_coulsq, _qqrd2e, _g_ewald; + + private: + bool _allocated; + void loop(const bool _eflag, const bool _vflag); +}; + +} + +#endif diff --git a/lib/gpu/lal_dipole_long_lj_ext.cpp b/lib/gpu/lal_dipole_long_lj_ext.cpp new file mode 100644 index 0000000000000000000000000000000000000000..b2751e8a8221a1620a1ebb61211158eee8f455ad --- /dev/null +++ b/lib/gpu/lal_dipole_long_lj_ext.cpp @@ -0,0 +1,129 @@ +/*************************************************************************** + dipole_long_lj_ext.cpp + ------------------- + Trung Dac Nguyen (ORNL) + + Functions for LAMMPS access to dipole/cut acceleration routines. + + __________________________________________________________________________ + This file is part of the LAMMPS Accelerator Library (LAMMPS_AL) + __________________________________________________________________________ + + begin : + email : nguyentd@ornl.gov + ***************************************************************************/ + +#include +#include +#include + +#include "lal_dipole_long_lj.h" + +using namespace std; +using namespace LAMMPS_AL; + +static DipoleLongLJ DPLJMF; + +// --------------------------------------------------------------------------- +// Allocate memory on host and device and copy constants to device +// --------------------------------------------------------------------------- +int dplj_gpu_init(const int ntypes, double **cutsq, double **host_lj1, + double **host_lj2, double **host_lj3, double **host_lj4, + double **offset, double *special_lj, const int inum, + const int nall, const int max_nbors, const int maxspecial, + const double cell_size, int &gpu_mode, FILE *screen, + double **host_cut_ljsq, const double host_cut_coulsq, + double *host_special_coul, const double qqrd2e, + const double g_ewald) { + DPLJMF.clear(); + gpu_mode=DPLJMF.device->gpu_mode(); + double gpu_split=DPLJMF.device->particle_split(); + int first_gpu=DPLJMF.device->first_device(); + int last_gpu=DPLJMF.device->last_device(); + int world_me=DPLJMF.device->world_me(); + int gpu_rank=DPLJMF.device->gpu_rank(); + int procs_per_gpu=DPLJMF.device->procs_per_gpu(); + + DPLJMF.device->init_message(screen,"lj/cut/dipole/long",first_gpu,last_gpu); + + bool message=false; + if (DPLJMF.device->replica_me()==0 && screen) + message=true; + + if (message) { + fprintf(screen,"Initializing Device and compiling on process 0..."); + fflush(screen); + } + + int init_ok=0; + if (world_me==0) + init_ok=DPLJMF.init(ntypes, cutsq, host_lj1, host_lj2, host_lj3, + host_lj4, offset, special_lj, inum, nall, 300, + maxspecial, cell_size, gpu_split, screen, host_cut_ljsq, + host_cut_coulsq, host_special_coul, qqrd2e, g_ewald); + + DPLJMF.device->world_barrier(); + if (message) + fprintf(screen,"Done.\n"); + + for (int i=0; igpu_barrier(); + if (message) + fprintf(screen,"Done.\n"); + } + if (message) + fprintf(screen,"\n"); + + if (init_ok==0) + DPLJMF.estimate_gpu_overhead(); + return init_ok; +} + +void dplj_gpu_clear() { + DPLJMF.clear(); +} + +int** dplj_gpu_compute_n(const int ago, const int inum_full, + const int nall, double **host_x, int *host_type, + double *sublo, double *subhi, tagint *tag, int **nspecial, + tagint **special, const bool eflag, const bool vflag, + const bool eatom, const bool vatom, int &host_start, + int **ilist, int **jnum, const double cpu_time, + bool &success, double *host_q, double **host_mu, + double *boxlo, double *prd) { + return DPLJMF.compute(ago, inum_full, nall, host_x, host_type, sublo, + subhi, tag, nspecial, special, eflag, vflag, eatom, + vatom, host_start, ilist, jnum, cpu_time, success, + host_q, host_mu, boxlo, prd); +} + +void dplj_gpu_compute(const int ago, const int inum_full, const int nall, + double **host_x, int *host_type, int *ilist, int *numj, + int **firstneigh, const bool eflag, const bool vflag, + const bool eatom, const bool vatom, int &host_start, + const double cpu_time, bool &success, double *host_q, + double **host_mu, const int nlocal, double *boxlo, double *prd) { + DPLJMF.compute(ago,inum_full,nall,host_x,host_type,ilist,numj,firstneigh,eflag, + vflag,eatom,vatom,host_start,cpu_time,success,host_q,host_mu, + nlocal,boxlo,prd); +} + +double dplj_gpu_bytes() { + return DPLJMF.host_memory_usage(); +} + + diff --git a/lib/gpu/lal_lj_expand_coul_long.cpp b/lib/gpu/lal_lj_expand_coul_long.cpp new file mode 100644 index 0000000000000000000000000000000000000000..6a3da9b68604ed9e6d4f5ddf936c141b14b90aae --- /dev/null +++ b/lib/gpu/lal_lj_expand_coul_long.cpp @@ -0,0 +1,183 @@ +/*************************************************************************** + lj_expand_coul_long.cpp + -------------------------- + Trung Nguyen (Northwestern) + + Class for acceleration of the lj/expand/coul/long pair style. + + __________________________________________________________________________ + This file is part of the LAMMPS Accelerator Library (LAMMPS_AL) + __________________________________________________________________________ + + begin : + email : trung.nguyen@northwestern.edu + ***************************************************************************/ + +#if defined(USE_OPENCL) +#include "lj_expand_coul_long_cl.h" +#elif defined(USE_CUDART) +const char *lj_expand_coul_long=0; +#else +#include "lj_expand_coul_long_cubin.h" +#endif + +#include "lal_lj_expand_coul_long.h" +#include +using namespace LAMMPS_AL; +#define LJExpandCoulLongT LJExpandCoulLong + +extern Device device; + +template +LJExpandCoulLongT::LJExpandCoulLong() : BaseCharge(), + _allocated(false) { +} + +template +LJExpandCoulLongT::~LJExpandCoulLong() { + clear(); +} + +template +int LJExpandCoulLongT::bytes_per_atom(const int max_nbors) const { + return this->bytes_per_atom_atomic(max_nbors); +} + +template +int LJExpandCoulLongT::init(const int ntypes, + double **host_cutsq, double **host_lj1, + double **host_lj2, double **host_lj3, + double **host_lj4, double **host_offset, double **host_shift, + double *host_special_lj, const int nlocal, + const int nall, const int max_nbors, + const int maxspecial, const double cell_size, + const double gpu_split, FILE *_screen, + double **host_cut_ljsq, const double host_cut_coulsq, + double *host_special_coul, const double qqrd2e, + const double g_ewald) { + int success; + success=this->init_atomic(nlocal,nall,max_nbors,maxspecial,cell_size,gpu_split, + _screen,lj_expand_coul_long,"k_lj_expand_coul_long"); + if (success!=0) + return success; + + // If atom type constants fit in shared memory use fast kernel + int lj_types=ntypes; + shared_types=false; + int max_shared_types=this->device->max_shared_types(); + if (lj_types<=max_shared_types && this->_block_size>=max_shared_types) { + lj_types=max_shared_types; + shared_types=true; + } + _lj_types=lj_types; + + // Allocate a host write buffer for data initialization + UCL_H_Vec host_write(lj_types*lj_types*32,*(this->ucl_device), + UCL_WRITE_ONLY); + + for (int i=0; iucl_device),UCL_READ_ONLY); + this->atom->type_pack4(ntypes,lj_types,lj1,host_write,host_lj1,host_lj2, + host_cutsq, host_cut_ljsq); + + lj3.alloc(lj_types*lj_types,*(this->ucl_device),UCL_READ_ONLY); + this->atom->type_pack4(ntypes,lj_types,lj3,host_write,host_lj3,host_lj4, + host_offset,host_shift); + + sp_lj.alloc(8,*(this->ucl_device),UCL_READ_ONLY); + for (int i=0; i<4; i++) { + host_write[i]=host_special_lj[i]; + host_write[i+4]=host_special_coul[i]; + } + ucl_copy(sp_lj,host_write,8,false); + + _cut_coulsq=host_cut_coulsq; + _qqrd2e=qqrd2e; + _g_ewald=g_ewald; + + _allocated=true; + this->_max_bytes=lj1.row_bytes()+lj3.row_bytes()+sp_lj.row_bytes(); + return 0; +} + +template +void LJExpandCoulLongT::reinit(const int ntypes, double **host_cutsq, double **host_lj1, + double **host_lj2, double **host_lj3, double **host_lj4, + double **host_offset, double **host_shift, double **host_cut_ljsq) { + // Allocate a host write buffer for data initialization + UCL_H_Vec host_write(_lj_types*_lj_types*32,*(this->ucl_device), + UCL_WRITE_ONLY); + + for (int i=0; i<_lj_types*_lj_types; i++) + host_write[i]=0.0; + + this->atom->type_pack4(ntypes,_lj_types,lj1,host_write,host_lj1,host_lj2, + host_cutsq, host_cut_ljsq); + this->atom->type_pack4(ntypes,_lj_types,lj3,host_write,host_lj3,host_lj4, + host_offset,host_shift); +} + +template +void LJExpandCoulLongT::clear() { + if (!_allocated) + return; + _allocated=false; + + lj1.clear(); + lj3.clear(); + sp_lj.clear(); + this->clear_atomic(); +} + +template +double LJExpandCoulLongT::host_memory_usage() const { + return this->host_memory_usage_atomic()+sizeof(LJExpandCoulLong); +} + +// --------------------------------------------------------------------------- +// Calculate energies, forces, and torques +// --------------------------------------------------------------------------- +template +void LJExpandCoulLongT::loop(const bool _eflag, const bool _vflag) { + // Compute the block size and grid size to keep all cores busy + const int BX=this->block_size(); + int eflag, vflag; + if (_eflag) + eflag=1; + else + eflag=0; + + if (_vflag) + vflag=1; + else + vflag=0; + + int GX=static_cast(ceil(static_cast(this->ans->inum())/ + (BX/this->_threads_per_atom))); + + int ainum=this->ans->inum(); + int nbor_pitch=this->nbor->nbor_pitch(); + this->time_pair.start(); + if (shared_types) { + this->k_pair_fast.set_size(GX,BX); + this->k_pair_fast.run(&this->atom->x, &lj1, &lj3, &sp_lj, + &this->nbor->dev_nbor, &this->_nbor_data->begin(), + &this->ans->force, &this->ans->engv, &eflag, + &vflag, &ainum, &nbor_pitch, &this->atom->q, + &_cut_coulsq, &_qqrd2e, &_g_ewald, + &this->_threads_per_atom); + } else { + this->k_pair.set_size(GX,BX); + this->k_pair.run(&this->atom->x, &lj1, &lj3, + &_lj_types, &sp_lj, &this->nbor->dev_nbor, + &this->_nbor_data->begin(), &this->ans->force, + &this->ans->engv, &eflag, &vflag, &ainum, + &nbor_pitch, &this->atom->q, &_cut_coulsq, + &_qqrd2e, &_g_ewald, &this->_threads_per_atom); + } + this->time_pair.stop(); +} + +template class LJExpandCoulLong; diff --git a/lib/gpu/lal_lj_expand_coul_long.cu b/lib/gpu/lal_lj_expand_coul_long.cu new file mode 100644 index 0000000000000000000000000000000000000000..eaea5b0e12aeada9e58288d3f33b08e61c150b46 --- /dev/null +++ b/lib/gpu/lal_lj_expand_coul_long.cu @@ -0,0 +1,265 @@ +// ************************************************************************** +// lj_coul_long.cu +// ------------------- +// Trung Nguyen (Northwestern) +// +// Device code for acceleration of the lj/expand/coul/long pair style +// +// __________________________________________________________________________ +// This file is part of the LAMMPS Accelerator Library (LAMMPS_AL) +// __________________________________________________________________________ +// +// begin : +// email : trung.nguyen@northwestern.edu +// ***************************************************************************/ + +#ifdef NV_KERNEL + +#include "lal_aux_fun1.h" +#ifndef _DOUBLE_DOUBLE +texture pos_tex; +texture q_tex; +#else +texture pos_tex; +texture q_tex; +#endif + +#else +#define pos_tex x_ +#define q_tex q_ +#endif + +__kernel void k_lj_expand_coul_long(const __global numtyp4 *restrict x_, + const __global numtyp4 *restrict lj1, + const __global numtyp4 *restrict lj3, + const int lj_types, + const __global numtyp *restrict sp_lj_in, + const __global int *dev_nbor, + const __global int *dev_packed, + __global acctyp4 *restrict ans, + __global acctyp *restrict engv, + const int eflag, const int vflag, const int inum, + const int nbor_pitch, + const __global numtyp *restrict q_, + const numtyp cut_coulsq, const numtyp qqrd2e, + const numtyp g_ewald, const int t_per_atom) { + int tid, ii, offset; + atom_info(t_per_atom,ii,tid,offset); + + __local numtyp sp_lj[8]; + sp_lj[0]=sp_lj_in[0]; + sp_lj[1]=sp_lj_in[1]; + sp_lj[2]=sp_lj_in[2]; + sp_lj[3]=sp_lj_in[3]; + sp_lj[4]=sp_lj_in[4]; + sp_lj[5]=sp_lj_in[5]; + sp_lj[6]=sp_lj_in[6]; + sp_lj[7]=sp_lj_in[7]; + + acctyp energy=(acctyp)0; + acctyp e_coul=(acctyp)0; + acctyp4 f; + f.x=(acctyp)0; f.y=(acctyp)0; f.z=(acctyp)0; + acctyp virial[6]; + for (int i=0; i<6; i++) + virial[i]=(acctyp)0; + + if (ii0) { + if (rsq < cut_coulsq) + e_coul += prefactor*(_erfc-factor_coul); + if (rsq < lj1[mtype].w) { + numtyp e=r6inv*(lj3[mtype].x*r6inv-lj3[mtype].y); + energy+=factor_lj*(e-lj3[mtype].z); + } + } + if (vflag>0) { + virial[0] += delx*delx*force; + virial[1] += dely*dely*force; + virial[2] += delz*delz*force; + virial[3] += delx*dely*force; + virial[4] += delx*delz*force; + virial[5] += dely*delz*force; + } + } + + } // for nbor + store_answers_q(f,energy,e_coul,virial,ii,inum,tid,t_per_atom,offset,eflag, + vflag,ans,engv); + } // if ii +} + +__kernel void k_lj_expand_coul_long_fast(const __global numtyp4 *restrict x_, + const __global numtyp4 *restrict lj1_in, + const __global numtyp4 *restrict lj3_in, + const __global numtyp *restrict sp_lj_in, + const __global int *dev_nbor, + const __global int *dev_packed, + __global acctyp4 *restrict ans, + __global acctyp *restrict engv, + const int eflag, const int vflag, + const int inum, const int nbor_pitch, + const __global numtyp *restrict q_, + const numtyp cut_coulsq, const numtyp qqrd2e, + const numtyp g_ewald, const int t_per_atom) { + int tid, ii, offset; + atom_info(t_per_atom,ii,tid,offset); + + __local numtyp4 lj1[MAX_SHARED_TYPES*MAX_SHARED_TYPES]; + __local numtyp4 lj3[MAX_SHARED_TYPES*MAX_SHARED_TYPES]; + __local numtyp sp_lj[8]; + if (tid<8) + sp_lj[tid]=sp_lj_in[tid]; + if (tid0) + lj3[tid]=lj3_in[tid]; + } + + acctyp energy=(acctyp)0; + acctyp e_coul=(acctyp)0; + acctyp4 f; + f.x=(acctyp)0; f.y=(acctyp)0; f.z=(acctyp)0; + acctyp virial[6]; + for (int i=0; i<6; i++) + virial[i]=(acctyp)0; + + __syncthreads(); + + if (ii0) { + if (rsq < cut_coulsq) + e_coul += prefactor*(_erfc-factor_coul); + if (rsq < lj1[mtype].w) { + numtyp e=r6inv*(lj3[mtype].x*r6inv-lj3[mtype].y); + energy+=factor_lj*(e-lj3[mtype].z); + } + } + if (vflag>0) { + virial[0] += delx*delx*force; + virial[1] += dely*dely*force; + virial[2] += delz*delz*force; + virial[3] += delx*dely*force; + virial[4] += delx*delz*force; + virial[5] += dely*delz*force; + } + } + + } // for nbor + store_answers_q(f,energy,e_coul,virial,ii,inum,tid,t_per_atom,offset,eflag, + vflag,ans,engv); + } // if ii +} + diff --git a/lib/gpu/lal_lj_expand_coul_long.h b/lib/gpu/lal_lj_expand_coul_long.h new file mode 100644 index 0000000000000000000000000000000000000000..c5c7a468a98d68f979bb098f49ad7aa5a70a85b4 --- /dev/null +++ b/lib/gpu/lal_lj_expand_coul_long.h @@ -0,0 +1,88 @@ +/*************************************************************************** + lj_expand_coul_long.h + ------------------- + Trung Nguyen (Northwestern) + + Class for acceleration of the lj/expand/coul/long pair style. + + __________________________________________________________________________ + This file is part of the LAMMPS Accelerator Library (LAMMPS_AL) + __________________________________________________________________________ + + begin : + email : trung.nguyen@northwestern.edu + ***************************************************************************/ + +#ifndef LAL_LJ_EXPAND_COUL_LONG_H +#define LAL_LJ_EXPAND_COUL_LONG_H + +#include "lal_base_charge.h" + +namespace LAMMPS_AL { + +template +class LJExpandCoulLong : public BaseCharge { + public: + LJExpandCoulLong(); + ~LJExpandCoulLong(); + + /// Clear any previous data and set up for a new LAMMPS run + /** \param max_nbors initial number of rows in the neighbor matrix + * \param cell_size cutoff + skin + * \param gpu_split fraction of particles handled by device + * + * Returns: + * - 0 if successfull + * - -1 if fix gpu not found + * - -3 if there is an out of memory error + * - -4 if the GPU library was not compiled for GPU + * - -5 Double precision is not supported on card **/ + int init(const int ntypes, double **host_cutsq, + double **host_lj1, double **host_lj2, double **host_lj3, + double **host_lj4, double **host_offset, double **host_shift, double *host_special_lj, + const int nlocal, const int nall, const int max_nbors, + const int maxspecial, const double cell_size, + const double gpu_split, FILE *screen, double **host_cut_ljsq, + const double host_cut_coulsq, double *host_special_coul, + const double qqrd2e, const double g_ewald); + + /// Send updated coeffs from host to device (to be compatible with fix adapt) + void reinit(const int ntypes, double **host_cutsq, + double **host_lj1, double **host_lj2, double **host_lj3, + double **host_lj4, double **host_offset, double **host_shift, double **host_cut_ljsq); + + /// Clear all host and device data + /** \note This is called at the beginning of the init() routine **/ + void clear(); + + /// Returns memory usage on device per atom + int bytes_per_atom(const int max_nbors) const; + + /// Total host memory used by library for pair style + double host_memory_usage() const; + + // --------------------------- TYPE DATA -------------------------- + + /// lj1.x = lj1, lj1.y = lj2, lj1.z = cutsq, lj1.w = cutsq_vdw + UCL_D_Vec lj1; + /// lj3.x = lj3, lj3.y = lj4, lj3.z = offset, lj3.w = shift + UCL_D_Vec lj3; + /// Special LJ values [0-3] and Special Coul values [4-7] + UCL_D_Vec sp_lj; + + /// If atom type constants fit in shared memory, use fast kernels + bool shared_types; + + /// Number of atom types + int _lj_types; + + numtyp _cut_coulsq, _qqrd2e, _g_ewald; + + private: + bool _allocated; + void loop(const bool _eflag, const bool _vflag); +}; + +} + +#endif diff --git a/lib/gpu/lal_lj_expand_coul_long_ext.cpp b/lib/gpu/lal_lj_expand_coul_long_ext.cpp new file mode 100644 index 0000000000000000000000000000000000000000..e4b53bbf20710681c2049fc1dea520f049a6e4b2 --- /dev/null +++ b/lib/gpu/lal_lj_expand_coul_long_ext.cpp @@ -0,0 +1,152 @@ +/*************************************************************************** + lj_coul_long_ext.cpp + ------------------------ + Trung Nguyen (Northwestern) + + Functions for LAMMPS access to lj/expand/coul/long acceleration routines. + + __________________________________________________________________________ + This file is part of the LAMMPS Accelerator Library (LAMMPS_AL) + __________________________________________________________________________ + + begin : + email : trung.nguyen@northwestern.edu + ***************************************************************************/ + +#include +#include +#include + +#include "lal_lj_expand_coul_long.h" + +using namespace std; +using namespace LAMMPS_AL; + +static LJExpandCoulLong LJECLMF; + +// --------------------------------------------------------------------------- +// Allocate memory on host and device and copy constants to device +// --------------------------------------------------------------------------- +int ljecl_gpu_init(const int ntypes, double **cutsq, double **host_lj1, + double **host_lj2, double **host_lj3, double **host_lj4, + double **offset, double **shift, double *special_lj, const int inum, + const int nall, const int max_nbors, const int maxspecial, + const double cell_size, int &gpu_mode, FILE *screen, + double **host_cut_ljsq, double host_cut_coulsq, + double *host_special_coul, const double qqrd2e, + const double g_ewald) { + LJECLMF.clear(); + gpu_mode=LJECLMF.device->gpu_mode(); + double gpu_split=LJECLMF.device->particle_split(); + int first_gpu=LJECLMF.device->first_device(); + int last_gpu=LJECLMF.device->last_device(); + int world_me=LJECLMF.device->world_me(); + int gpu_rank=LJECLMF.device->gpu_rank(); + int procs_per_gpu=LJECLMF.device->procs_per_gpu(); + + LJECLMF.device->init_message(screen,"lj/expand/coul/long",first_gpu,last_gpu); + + bool message=false; + if (LJECLMF.device->replica_me()==0 && screen) + message=true; + + if (message) { + fprintf(screen,"Initializing Device and compiling on process 0..."); + fflush(screen); + } + + int init_ok=0; + if (world_me==0) + init_ok=LJECLMF.init(ntypes, cutsq, host_lj1, host_lj2, host_lj3, host_lj4, + offset, shift, special_lj, inum, nall, 300, maxspecial, + cell_size, gpu_split, screen, host_cut_ljsq, + host_cut_coulsq, host_special_coul, qqrd2e, g_ewald); + + LJECLMF.device->world_barrier(); + if (message) + fprintf(screen,"Done.\n"); + + for (int i=0; igpu_barrier(); + if (message) + fprintf(screen,"Done.\n"); + } + if (message) + fprintf(screen,"\n"); + + if (init_ok==0) + LJECLMF.estimate_gpu_overhead(); + return init_ok; +} + +// --------------------------------------------------------------------------- +// Copy updated coeffs from host to device +// --------------------------------------------------------------------------- +void ljecl_gpu_reinit(const int ntypes, double **cutsq, double **host_lj1, + double **host_lj2, double **host_lj3, double **host_lj4, + double **offset, double **shift, double **host_cut_ljsq) { + int world_me=LJECLMF.device->world_me(); + int gpu_rank=LJECLMF.device->gpu_rank(); + int procs_per_gpu=LJECLMF.device->procs_per_gpu(); + + if (world_me==0) + LJECLMF.reinit(ntypes, cutsq, host_lj1, host_lj2, host_lj3, host_lj4, + offset, shift, host_cut_ljsq); + LJECLMF.device->world_barrier(); + + for (int i=0; igpu_barrier(); + } +} + +void ljecl_gpu_clear() { + LJECLMF.clear(); +} + +int** ljecl_gpu_compute_n(const int ago, const int inum_full, + const int nall, double **host_x, int *host_type, + double *sublo, double *subhi, tagint *tag, int **nspecial, + tagint **special, const bool eflag, const bool vflag, + const bool eatom, const bool vatom, int &host_start, + int **ilist, int **jnum, const double cpu_time, + bool &success, double *host_q, double *boxlo, + double *prd) { + return LJECLMF.compute(ago, inum_full, nall, host_x, host_type, sublo, + subhi, tag, nspecial, special, eflag, vflag, eatom, + vatom, host_start, ilist, jnum, cpu_time, success, + host_q, boxlo, prd); +} + +void ljecl_gpu_compute(const int ago, const int inum_full, const int nall, + double **host_x, int *host_type, int *ilist, int *numj, + int **firstneigh, const bool eflag, const bool vflag, + const bool eatom, const bool vatom, int &host_start, + const double cpu_time, bool &success, double *host_q, + const int nlocal, double *boxlo, double *prd) { + LJECLMF.compute(ago,inum_full,nall,host_x,host_type,ilist,numj, + firstneigh,eflag,vflag,eatom,vatom,host_start,cpu_time,success, + host_q,nlocal,boxlo,prd); +} + +double ljecl_gpu_bytes() { + return LJECLMF.host_memory_usage(); +} + + diff --git a/lib/kokkos/CHANGELOG.md b/lib/kokkos/CHANGELOG.md index feb2bd547f034a5ef233667a3a9b5c3d0e02d3f8..145cc62706eea04fab0bb6e3fbcb676ee03b9ef8 100644 --- a/lib/kokkos/CHANGELOG.md +++ b/lib/kokkos/CHANGELOG.md @@ -1,5 +1,58 @@ # Change Log +## [2.7.00](https://github.com/kokkos/kokkos/tree/2.7.00) (2018-05-24) +[Full Changelog](https://github.com/kokkos/kokkos/compare/2.6.00...2.7.00) + +**Part of the Kokkos C++ Performance Portability Programming EcoSystem 2.7** + +**Implemented enhancements:** + +- Deprecate team\_size auto adjusting to maximal value possible [\#1618](https://github.com/kokkos/kokkos/issues/1618) +- DynamicView - remove restrictions to std::is\_trivial types and value\_type is power of two [\#1586](https://github.com/kokkos/kokkos/issues/1586) +- Kokkos::StaticCrsGraph does not propagate memory traits \(e.g., Unmanaged\) [\#1581](https://github.com/kokkos/kokkos/issues/1581) +- Adding ETI for DeepCopy / ViewFill etc. [\#1578](https://github.com/kokkos/kokkos/issues/1578) +- Deprecate all the left over KOKKOS\_HAVE\_ Macros and Kokkos\_OldMacros.hpp [\#1572](https://github.com/kokkos/kokkos/issues/1572) +- Error if Kokkos\_ARCH set in CMake [\#1555](https://github.com/kokkos/kokkos/issues/1555) +- Deprecate ExecSpace::initialize / ExecSpace::finalize [\#1532](https://github.com/kokkos/kokkos/issues/1532) +- New API for TeamPolicy property setting [\#1531](https://github.com/kokkos/kokkos/issues/1531) +- clang 6.0 + cuda debug out-of-memory test failure [\#1521](https://github.com/kokkos/kokkos/issues/1521) +- Cuda UniqueToken interface not consistent with other backends [\#1505](https://github.com/kokkos/kokkos/issues/1505) +- Move Reducers out of Experimental namespace [\#1494](https://github.com/kokkos/kokkos/issues/1494) +- Provide scope guard for initialize/finalize [\#1479](https://github.com/kokkos/kokkos/issues/1479) +- Check Kokkos::is\_initialized in SharedAllocationRecord dtor [\#1465](https://github.com/kokkos/kokkos/issues/1465) +- Remove static list of allocations [\#1464](https://github.com/kokkos/kokkos/issues/1464) +- Makefiles: Support single compile/link line use case [\#1402](https://github.com/kokkos/kokkos/issues/1402) +- ThreadVectorRange with a range [\#1400](https://github.com/kokkos/kokkos/issues/1400) +- Exclusive scan + last value API [\#1358](https://github.com/kokkos/kokkos/issues/1358) +- Install kokkos\_generated\_settings.cmake [\#1348](https://github.com/kokkos/kokkos/issues/1348) +- Kokkos arrays \(not views!\) don't do bounds checking in debug mode [\#1342](https://github.com/kokkos/kokkos/issues/1342) +- Expose round-robin GPU assignment outside of initialize\(int, char\*\*\) [\#1318](https://github.com/kokkos/kokkos/issues/1318) +- DynamicView misses use\_count and label function [\#1298](https://github.com/kokkos/kokkos/issues/1298) +- View constructor should check arguments [\#1286](https://github.com/kokkos/kokkos/issues/1286) +- False Positive on Oversubscription Warning [\#1207](https://github.com/kokkos/kokkos/issues/1207) +- Allow \(require\) execution space for 1st arg of VerifyExecutionCanAccessMemorySpace [\#1192](https://github.com/kokkos/kokkos/issues/1192) +- ROCm: Add ROCmHostPinnedSpace [\#958](https://github.com/kokkos/kokkos/issues/958) +- power of two functions [\#656](https://github.com/kokkos/kokkos/issues/656) +- CUDA 8 has 64bit \_\_shfl [\#361](https://github.com/kokkos/kokkos/issues/361) +- Add TriBITS/CMake configure information about node types [\#243](https://github.com/kokkos/kokkos/issues/243) + +**Fixed bugs:** + +- CUDA atomic\_fetch\_sub for doubles is hitting CAS instead of intrinsic [\#1624](https://github.com/kokkos/kokkos/issues/1624) +- Bug: use of ballot on Volta [\#1612](https://github.com/kokkos/kokkos/issues/1612) +- Kokkos::deep\_copy memory access failures [\#1583](https://github.com/kokkos/kokkos/issues/1583) +- g++ -std option doubly set for cmake project [\#1548](https://github.com/kokkos/kokkos/issues/1548) +- ViewFill for 1D Views of larger 32bit entries fails [\#1541](https://github.com/kokkos/kokkos/issues/1541) +- CUDA Volta another warpsync bug [\#1520](https://github.com/kokkos/kokkos/issues/1520) +- triple\_nested\_parallelism fails with KOKKOS\_DEBUG and CUDA [\#1513](https://github.com/kokkos/kokkos/issues/1513) +- Jenkins errors in Kokkos\_SharedAlloc.cpp with debug build [\#1511](https://github.com/kokkos/kokkos/issues/1511) +- Kokkos::Sort out-of-bounds with empty bins [\#1504](https://github.com/kokkos/kokkos/issues/1504) +- Get rid of deprecated functions inside Kokkos [\#1484](https://github.com/kokkos/kokkos/issues/1484) +- get\_work\_partition casts int64\_t to int, causing a seg fault [\#1481](https://github.com/kokkos/kokkos/issues/1481) +- NVCC bug with \_\_device\_\_ on defaulted function [\#1470](https://github.com/kokkos/kokkos/issues/1470) +- CMake example broken with CUDA backend [\#1468](https://github.com/kokkos/kokkos/issues/1468) + + ## [2.6.00](https://github.com/kokkos/kokkos/tree/2.6.00) (2018-03-07) [Full Changelog](https://github.com/kokkos/kokkos/compare/2.5.00...2.6.00) diff --git a/lib/kokkos/CMakeLists.txt b/lib/kokkos/CMakeLists.txt index cd1f4ea981339a5ff01095be9bca144cdb5d9064..9e5308f1c3245ea40a01079909cd9639cf22c82f 100644 --- a/lib/kokkos/CMakeLists.txt +++ b/lib/kokkos/CMakeLists.txt @@ -44,6 +44,7 @@ IF(NOT KOKKOS_HAS_TRILINOS) "${KOKKOS_SETTINGS} make -f ${KOKKOS_SRC_PATH}/cmake/Makefile.generate_cmake_settings CXX=${CMAKE_CXX_COMPILER} generate_build_settings") endif() include(${Kokkos_BINARY_DIR}/kokkos_generated_settings.cmake) + install(FILES ${Kokkos_BINARY_DIR}/kokkos_generated_settings.cmake DESTINATION lib/cmake/Kokkos) string(REPLACE " " ";" KOKKOS_TPL_INCLUDE_DIRS "${KOKKOS_GMAKE_TPL_INCLUDE_DIRS}") string(REPLACE " " ";" KOKKOS_TPL_LIBRARY_DIRS "${KOKKOS_GMAKE_TPL_LIBRARY_DIRS}") string(REPLACE " " ";" KOKKOS_TPL_LIBRARY_NAMES "${KOKKOS_GMAKE_TPL_LIBRARY_NAMES}") diff --git a/lib/kokkos/Makefile.kokkos b/lib/kokkos/Makefile.kokkos index a7bb63f190348405267c3ea50e0f55dcf20645a3..52cd3969f9f74d736fc9be99ad64114b8f5af6cf 100644 --- a/lib/kokkos/Makefile.kokkos +++ b/lib/kokkos/Makefile.kokkos @@ -1,7 +1,9 @@ # Default settings common options. #LAMMPS specific settings: -KOKKOS_PATH=../../lib/kokkos +ifndef KOKKOS_PATH + KOKKOS_PATH=../../lib/kokkos +endif CXXFLAGS=$(CCFLAGS) # Options: Cuda,ROCm,OpenMP,Pthreads,Qthreads,Serial @@ -21,8 +23,10 @@ KOKKOS_DEBUG ?= "no" KOKKOS_USE_TPLS ?= "" # Options: c++11,c++1z KOKKOS_CXX_STANDARD ?= "c++11" -# Options: aggressive_vectorization,disable_profiling,disable_deprecated_code +# Options: aggressive_vectorization,disable_profiling,disable_deprecated_code,enable_large_mem_tests KOKKOS_OPTIONS ?= "" +# Option for setting ETI path +KOKKOS_ETI_PATH ?= ${KOKKOS_PATH}/core/src/eti # Default settings specific options. # Options: force_uvm,use_ldg,rdc,enable_lambda @@ -51,10 +55,12 @@ KOKKOS_INTERNAL_DISABLE_PROFILING := $(call kokkos_has_string,$(KOKKOS_OPTIONS), KOKKOS_INTERNAL_DISABLE_DEPRECATED_CODE := $(call kokkos_has_string,$(KOKKOS_OPTIONS),disable_deprecated_code) KOKKOS_INTERNAL_DISABLE_DUALVIEW_MODIFY_CHECK := $(call kokkos_has_string,$(KOKKOS_OPTIONS),disable_dualview_modify_check) KOKKOS_INTERNAL_ENABLE_PROFILING_LOAD_PRINT := $(call kokkos_has_string,$(KOKKOS_OPTIONS),enable_profile_load_print) +KOKKOS_INTERNAL_ENABLE_LARGE_MEM_TESTS := $(call kokkos_has_string,$(KOKKOS_OPTIONS),enable_large_mem_tests) KOKKOS_INTERNAL_CUDA_USE_LDG := $(call kokkos_has_string,$(KOKKOS_CUDA_OPTIONS),use_ldg) KOKKOS_INTERNAL_CUDA_USE_UVM := $(call kokkos_has_string,$(KOKKOS_CUDA_OPTIONS),force_uvm) KOKKOS_INTERNAL_CUDA_USE_RELOC := $(call kokkos_has_string,$(KOKKOS_CUDA_OPTIONS),rdc) KOKKOS_INTERNAL_CUDA_USE_LAMBDA := $(call kokkos_has_string,$(KOKKOS_CUDA_OPTIONS),enable_lambda) +KOKKOS_INTERNAL_ENABLE_ETI := $(call kokkos_has_string,$(KOKKOS_OPTIONS),enable_eti) # Check for Kokkos Host Execution Spaces one of which must be on. @@ -78,7 +84,12 @@ KOKKOS_INTERNAL_USE_OPENMPTARGET := $(call kokkos_has_string,$(KOKKOS_DEVICES),O ifeq ($(KOKKOS_INTERNAL_USE_CUDA), 1) KOKKOS_INTERNAL_NVCC_PATH := $(shell which nvcc) - CUDA_PATH ?= $(KOKKOS_INTERNAL_NVCC_PATH:/bin/nvcc=) + ifeq ($(origin CUDA_PATH), undefined) + CUDA_PATH = $(KOKKOS_INTERNAL_NVCC_PATH:/bin/nvcc=) + endif + ifeq ($(CUDA_PATH),) + CUDA_PATH = $(KOKKOS_INTERNAL_NVCC_PATH:/bin/nvcc=) + endif KOKKOS_INTERNAL_COMPILER_NVCC_VERSION := $(shell nvcc --version 2>&1 | grep release | cut -d' ' -f5 | cut -d',' -f1 | tr -d .) endif @@ -116,7 +127,7 @@ ifeq ($(KOKKOS_INTERNAL_COMPILER_HCC), 1) endif ifeq ($(KOKKOS_INTERNAL_COMPILER_CLANG), 1) - KOKKOS_INTERNAL_COMPILER_CLANG_VERSION := $(shell clang --version | grep version | cut -d ' ' -f3 | tr -d '.') + KOKKOS_INTERNAL_COMPILER_CLANG_VERSION := $(shell $(CXX) --version | grep version | cut -d ' ' -f3 | tr -d '.') ifeq ($(KOKKOS_INTERNAL_USE_CUDA), 1) ifeq ($(shell test $(KOKKOS_INTERNAL_COMPILER_CLANG_VERSION) -lt 400; echo $$?),0) @@ -323,12 +334,13 @@ endif # Generating the list of Flags. -KOKKOS_CPPFLAGS = -I./ -I$(KOKKOS_PATH)/core/src -I$(KOKKOS_PATH)/containers/src -I$(KOKKOS_PATH)/algorithms/src +#CPPFLAGS is now unused +KOKKOS_CPPFLAGS = +KOKKOS_CXXFLAGS = -I./ -I$(KOKKOS_PATH)/core/src -I$(KOKKOS_PATH)/containers/src -I$(KOKKOS_PATH)/algorithms/src -I$(KOKKOS_ETI_PATH) KOKKOS_TPL_INCLUDE_DIRS = KOKKOS_TPL_LIBRARY_DIRS = KOKKOS_TPL_LIBRARY_NAMES = -KOKKOS_CXXFLAGS = ifeq ($(KOKKOS_INTERNAL_ENABLE_COMPILER_WARNINGS), 1) KOKKOS_CXXFLAGS += $(KOKKOS_INTERNAL_COMPILER_WARNINGS) endif @@ -336,6 +348,8 @@ endif KOKKOS_LIBS = -ldl KOKKOS_TPL_LIBRARY_NAMES += dl KOKKOS_LDFLAGS = -L$(shell pwd) +# CXXLDFLAGS is used together with CXXFLAGS in a combined compile/link command +KOKKOS_CXXLDFLAGS = -L$(shell pwd) KOKKOS_LINK_FLAGS = KOKKOS_SRC = KOKKOS_HEADERS = @@ -362,7 +376,7 @@ tmp := $(call kokkos_append_header,'\#endif') tmp := $(call kokkos_append_header,"/* Execution Spaces */") ifeq ($(KOKKOS_INTERNAL_USE_CUDA), 1) - tmp := $(call kokkos_append_header,"\#define KOKKOS_HAVE_CUDA") + tmp := $(call kokkos_append_header,"\#define KOKKOS_ENABLE_CUDA") endif ifeq ($(KOKKOS_INTERNAL_USE_ROCM), 1) @@ -374,19 +388,19 @@ ifeq ($(KOKKOS_INTERNAL_USE_OPENMPTARGET), 1) endif ifeq ($(KOKKOS_INTERNAL_USE_OPENMP), 1) - tmp := $(call kokkos_append_header,'\#define KOKKOS_HAVE_OPENMP') + tmp := $(call kokkos_append_header,'\#define KOKKOS_ENABLE_OPENMP') endif ifeq ($(KOKKOS_INTERNAL_USE_PTHREADS), 1) - tmp := $(call kokkos_append_header,"\#define KOKKOS_HAVE_PTHREAD") + tmp := $(call kokkos_append_header,"\#define KOKKOS_ENABLE_THREADS") endif ifeq ($(KOKKOS_INTERNAL_USE_QTHREADS), 1) - tmp := $(call kokkos_append_header,"\#define KOKKOS_HAVE_QTHREADS") + tmp := $(call kokkos_append_header,"\#define KOKKOS_ENABLE_QTHREADS") endif ifeq ($(KOKKOS_INTERNAL_USE_SERIAL), 1) - tmp := $(call kokkos_append_header,"\#define KOKKOS_HAVE_SERIAL") + tmp := $(call kokkos_append_header,"\#define KOKKOS_ENABLE_SERIAL") endif ifeq ($(KOKKOS_INTERNAL_USE_TM), 1) @@ -422,13 +436,13 @@ endif tmp := $(call kokkos_append_header,"/* General Settings */") ifeq ($(KOKKOS_INTERNAL_ENABLE_CXX11), 1) KOKKOS_CXXFLAGS += $(KOKKOS_INTERNAL_CXX11_FLAG) - tmp := $(call kokkos_append_header,"\#define KOKKOS_HAVE_CXX11") + tmp := $(call kokkos_append_header,"\#define KOKKOS_ENABLE_CXX11") endif ifeq ($(KOKKOS_INTERNAL_ENABLE_CXX1Z), 1) KOKKOS_CXXFLAGS += $(KOKKOS_INTERNAL_CXX1Z_FLAG) - tmp := $(call kokkos_append_header,"\#define KOKKOS_HAVE_CXX11") - tmp := $(call kokkos_append_header,"\#define KOKKOS_HAVE_CXX1Z") + tmp := $(call kokkos_append_header,"\#define KOKKOS_ENABLE_CXX11") + tmp := $(call kokkos_append_header,"\#define KOKKOS_ENABLE_CXX1Z") endif ifeq ($(KOKKOS_INTERNAL_ENABLE_DEBUG), 1) @@ -437,9 +451,9 @@ ifeq ($(KOKKOS_INTERNAL_ENABLE_DEBUG), 1) endif KOKKOS_CXXFLAGS += -g - KOKKOS_LDFLAGS += -g -ldl + KOKKOS_LDFLAGS += -g tmp := $(call kokkos_append_header,"\#define KOKKOS_ENABLE_DEBUG_BOUNDS_CHECK") - tmp := $(call kokkos_append_header,"\#define KOKKOS_HAVE_DEBUG") + tmp := $(call kokkos_append_header,"\#define KOKKOS_ENABLE_DEBUG") ifeq ($(KOKKOS_INTERNAL_DISABLE_DUALVIEW_MODIFY_CHECK), 0) tmp := $(call kokkos_append_header,"\#define KOKKOS_ENABLE_DEBUG_DUALVIEW_MODIFY_CHECK") endif @@ -451,14 +465,15 @@ endif ifeq ($(KOKKOS_INTERNAL_USE_HWLOC), 1) ifneq ($(HWLOC_PATH),) - KOKKOS_CPPFLAGS += -I$(HWLOC_PATH)/include + KOKKOS_CXXFLAGS += -I$(HWLOC_PATH)/include KOKKOS_LDFLAGS += -L$(HWLOC_PATH)/lib + KOKKOS_CXXLDFLAGS += -L$(HWLOC_PATH)/lib KOKKOS_TPL_INCLUDE_DIRS += $(HWLOC_PATH)/include KOKKOS_TPL_LIBRARY_DIRS += $(HWLOC_PATH)/lib endif KOKKOS_LIBS += -lhwloc KOKKOS_TPL_LIBRARY_NAMES += hwloc - tmp := $(call kokkos_append_header,"\#define KOKKOS_HAVE_HWLOC") + tmp := $(call kokkos_append_header,"\#define KOKKOS_ENABLE_HWLOC") endif ifeq ($(KOKKOS_INTERNAL_USE_LIBRT), 1) @@ -469,14 +484,15 @@ endif ifeq ($(KOKKOS_INTERNAL_USE_MEMKIND), 1) ifneq ($(MEMKIND_PATH),) - KOKKOS_CPPFLAGS += -I$(MEMKIND_PATH)/include + KOKKOS_CXXFLAGS += -I$(MEMKIND_PATH)/include KOKKOS_LDFLAGS += -L$(MEMKIND_PATH)/lib + KOKKOS_CXXLDFLAGS += -L$(MEMKIND_PATH)/lib KOKKOS_TPL_INCLUDE_DIRS += $(MEMKIND_PATH)/include KOKKOS_TPL_LIBRARY_DIRS += $(MEMKIND_PATH)/lib endif KOKKOS_LIBS += -lmemkind -lnuma KOKKOS_TPL_LIBRARY_NAMES += memkind numa - tmp := $(call kokkos_append_header,"\#define KOKKOS_HAVE_HBWSPACE") + tmp := $(call kokkos_append_header,"\#define KOKKOS_ENABLE_HBWSPACE") endif ifeq ($(KOKKOS_INTERNAL_DISABLE_PROFILING), 0) @@ -486,6 +502,13 @@ endif ifeq ($(KOKKOS_INTERNAL_DISABLE_DEPRECATED_CODE), 0) tmp := $(call kokkos_append_header,"\#define KOKKOS_ENABLE_DEPRECATED_CODE") endif +ifeq ($(KOKKOS_INTERNAL_ENABLE_ETI), 1) + tmp := $(call kokkos_append_header,"\#define KOKKOS_ENABLE_ETI") +endif + +ifeq ($(KOKKOS_INTERNAL_ENABLE_LARGE_MEM_TESTS), 1) + tmp := $(call kokkos_append_header,"\#define KOKKOS_ENABLE_LARGE_MEM_TESTS") +endif tmp := $(call kokkos_append_header,"/* Optimization Settings */") @@ -497,27 +520,35 @@ tmp := $(call kokkos_append_header,"/* Cuda Settings */") ifeq ($(KOKKOS_INTERNAL_USE_CUDA), 1) ifeq ($(KOKKOS_INTERNAL_CUDA_USE_LDG), 1) - tmp := $(call kokkos_append_header,"\#define KOKKOS_CUDA_USE_LDG_INTRINSIC") + tmp := $(call kokkos_append_header,"\#define KOKKOS_ENABLE_CUDA_LDG_INTRINSIC") else ifeq ($(KOKKOS_INTERNAL_COMPILER_CLANG), 1) - tmp := $(call kokkos_append_header,"\#define KOKKOS_CUDA_USE_LDG_INTRINSIC") + tmp := $(call kokkos_append_header,"\#define KOKKOS_ENABLE_CUDA_LDG_INTRINSIC") endif endif ifeq ($(KOKKOS_INTERNAL_CUDA_USE_UVM), 1) - tmp := $(call kokkos_append_header,"\#define KOKKOS_CUDA_USE_UVM") + tmp := $(call kokkos_append_header,"\#define KOKKOS_ENABLE_CUDA_UVM") endif ifeq ($(KOKKOS_INTERNAL_CUDA_USE_RELOC), 1) - tmp := $(call kokkos_append_header,"\#define KOKKOS_CUDA_USE_RELOCATABLE_DEVICE_CODE") + tmp := $(call kokkos_append_header,"\#define KOKKOS_ENABLE_CUDA_RELOCATABLE_DEVICE_CODE") KOKKOS_CXXFLAGS += --relocatable-device-code=true KOKKOS_LDFLAGS += --relocatable-device-code=true endif + ifeq ($(KOKKOS_INTERNAL_COMPILER_NVCC), 1) + ifeq ($(shell test $(KOKKOS_INTERNAL_COMPILER_NVCC_VERSION) -ge 90; echo $$?),0) + # This diagnostic is just plain wrong in CUDA 9 + # See https://github.com/kokkos/kokkos/issues/1470 + KOKKOS_CXXFLAGS += -Xcudafe --diag_suppress=esa_on_defaulted_function_ignored + endif + endif + ifeq ($(KOKKOS_INTERNAL_CUDA_USE_LAMBDA), 1) ifeq ($(KOKKOS_INTERNAL_COMPILER_NVCC), 1) ifeq ($(shell test $(KOKKOS_INTERNAL_COMPILER_NVCC_VERSION) -gt 70; echo $$?),0) - tmp := $(call kokkos_append_header,"\#define KOKKOS_CUDA_USE_LAMBDA") + tmp := $(call kokkos_append_header,"\#define KOKKOS_ENABLE_CUDA_LAMBDA") KOKKOS_CXXFLAGS += -expt-extended-lambda else $(warning Warning: Cuda Lambda support was requested but NVCC version is too low. This requires NVCC for Cuda version 7.5 or higher. Disabling Lambda support now.) @@ -525,12 +556,12 @@ ifeq ($(KOKKOS_INTERNAL_USE_CUDA), 1) endif ifeq ($(KOKKOS_INTERNAL_COMPILER_CLANG), 1) - tmp := $(call kokkos_append_header,"\#define KOKKOS_CUDA_USE_LAMBDA") + tmp := $(call kokkos_append_header,"\#define KOKKOS_ENABLE_CUDA_LAMBDA") endif endif ifeq ($(KOKKOS_INTERNAL_COMPILER_CLANG), 1) - tmp := $(call kokkos_append_header,"\#define KOKKOS_CUDA_CLANG_WORKAROUND") + tmp := $(call kokkos_append_header,"\#define KOKKOS_IMPL_CUDA_CLANG_WORKAROUND") endif endif @@ -907,10 +938,14 @@ ifeq ($(KOKKOS_INTERNAL_USE_ROCM), 1) KOKKOS_CXXFLAGS += $(shell $(ROCM_HCC_PATH)/bin/hcc-config --cxxflags) KOKKOS_LDFLAGS += $(shell $(ROCM_HCC_PATH)/bin/hcc-config --ldflags) -lhc_am -lm + KOKKOS_CXXLDFLAGS += $(shell $(ROCM_HCC_PATH)/bin/hcc-config --ldflags) -lhc_am -lm KOKKOS_TPL_LIBRARY_NAMES += hc_am m KOKKOS_LDFLAGS += $(KOKKOS_INTERNAL_ROCM_ARCH_FLAG) KOKKOS_SRC += $(wildcard $(KOKKOS_PATH)/core/src/ROCm/*.cpp) +ifeq ($(KOKKOS_INTERNAL_ENABLE_ETI), 1) + KOKKOS_SRC += $(wildcard $(KOKKOS_ETI_PATH)/ROCm/*.cpp) +endif KOKKOS_HEADERS += $(wildcard $(KOKKOS_PATH)/core/src/ROCm/*.hpp) endif @@ -937,10 +972,14 @@ KOKKOS_SRC += $(wildcard $(KOKKOS_PATH)/containers/src/impl/*.cpp) ifeq ($(KOKKOS_INTERNAL_USE_CUDA), 1) KOKKOS_SRC += $(wildcard $(KOKKOS_PATH)/core/src/Cuda/*.cpp) +ifeq ($(KOKKOS_INTERNAL_ENABLE_ETI), 1) + KOKKOS_SRC += $(wildcard $(KOKKOS_ETI_PATH)/Cuda/*.cpp) +endif KOKKOS_HEADERS += $(wildcard $(KOKKOS_PATH)/core/src/Cuda/*.hpp) ifneq ($(CUDA_PATH),) - KOKKOS_CPPFLAGS += -I$(CUDA_PATH)/include + KOKKOS_CXXFLAGS += -I$(CUDA_PATH)/include KOKKOS_LDFLAGS += -L$(CUDA_PATH)/lib64 + KOKKOS_CXXLDFLAGS += -L$(CUDA_PATH)/lib64 KOKKOS_TPL_INCLUDE_DIRS += $(CUDA_PATH)/include KOKKOS_TPL_LIBRARY_DIRS += $(CUDA_PATH)/lib64 ifeq ($(KOKKOS_INTERNAL_COMPILER_CLANG), 1) @@ -964,6 +1003,9 @@ endif ifeq ($(KOKKOS_INTERNAL_USE_OPENMP), 1) KOKKOS_SRC += $(wildcard $(KOKKOS_PATH)/core/src/OpenMP/*.cpp) +ifeq ($(KOKKOS_INTERNAL_ENABLE_ETI), 1) + KOKKOS_SRC += $(wildcard $(KOKKOS_ETI_PATH)/OpenMP/*.cpp) +endif KOKKOS_HEADERS += $(wildcard $(KOKKOS_PATH)/core/src/OpenMP/*.hpp) ifeq ($(KOKKOS_INTERNAL_COMPILER_NVCC), 1) @@ -978,6 +1020,9 @@ endif ifeq ($(KOKKOS_INTERNAL_USE_PTHREADS), 1) KOKKOS_SRC += $(wildcard $(KOKKOS_PATH)/core/src/Threads/*.cpp) +ifeq ($(KOKKOS_INTERNAL_ENABLE_ETI), 1) + KOKKOS_SRC += $(wildcard $(KOKKOS_ETI_PATH)/Threads/*.cpp) +endif KOKKOS_HEADERS += $(wildcard $(KOKKOS_PATH)/core/src/Threads/*.hpp) KOKKOS_LIBS += -lpthread KOKKOS_TPL_LIBRARY_NAMES += pthread @@ -987,8 +1032,9 @@ ifeq ($(KOKKOS_INTERNAL_USE_QTHREADS), 1) KOKKOS_SRC += $(wildcard $(KOKKOS_PATH)/core/src/Qthreads/*.cpp) KOKKOS_HEADERS += $(wildcard $(KOKKOS_PATH)/core/src/Qthreads/*.hpp) ifneq ($(QTHREADS_PATH),) - KOKKOS_CPPFLAGS += -I$(QTHREADS_PATH)/include + KOKKOS_CXXFLAGS += -I$(QTHREADS_PATH)/include KOKKOS_LDFLAGS += -L$(QTHREADS_PATH)/lib + KOKKOS_CXXLDFLAGS += -L$(QTHREADS_PATH)/lib KOKKOS_TPL_INCLUDE_DIRS += $(QTHREADS_PATH)/include KOKKOS_TPL_LIBRARY_DIRS += $(QTHREADS_PATH)/lib64 endif @@ -1011,6 +1057,11 @@ endif # Don't include Kokkos_Serial.cpp or Kokkos_Serial_Task.cpp if not using Serial # device to avoid a link warning. +ifeq ($(KOKKOS_INTERNAL_USE_SERIAL), 1) +ifeq ($(KOKKOS_INTERNAL_ENABLE_ETI), 1) + KOKKOS_SRC += $(wildcard $(KOKKOS_ETI_PATH)/Serial/*.cpp) +endif +endif ifneq ($(KOKKOS_INTERNAL_USE_SERIAL), 1) KOKKOS_SRC := $(filter-out $(KOKKOS_PATH)/core/src/impl/Kokkos_Serial.cpp,$(KOKKOS_SRC)) KOKKOS_SRC := $(filter-out $(KOKKOS_PATH)/core/src/impl/Kokkos_Serial_Task.cpp,$(KOKKOS_SRC)) diff --git a/lib/kokkos/Makefile.targets b/lib/kokkos/Makefile.targets index a63598577c720a7b522a5f10192aa413046564b9..44da1e082ae12f6293a3fb925fb8b30e76e28248 100644 --- a/lib/kokkos/Makefile.targets +++ b/lib/kokkos/Makefile.targets @@ -31,6 +31,12 @@ Kokkos_SharedAlloc.o: $(KOKKOS_CPP_DEPENDS) $(KOKKOS_PATH)/core/src/impl/Kokkos_ Kokkos_MemoryPool.o: $(KOKKOS_CPP_DEPENDS) $(KOKKOS_PATH)/core/src/impl/Kokkos_MemoryPool.cpp $(CXX) $(KOKKOS_CPPFLAGS) $(KOKKOS_CXXFLAGS) $(CXXFLAGS) -c $(KOKKOS_PATH)/core/src/impl/Kokkos_MemoryPool.cpp +ifeq ($(KOKKOS_INTERNAL_USE_SERIAL), 1) +ifeq ($(KOKKOS_INTERNAL_ENABLE_ETI), 1) + include $(KOKKOS_ETI_PATH)/Serial/Makefile.eti_Serial +endif +endif + ifeq ($(KOKKOS_INTERNAL_USE_CUDA), 1) Kokkos_Cuda_Impl.o: $(KOKKOS_CPP_DEPENDS) $(KOKKOS_PATH)/core/src/Cuda/Kokkos_Cuda_Impl.cpp $(CXX) $(KOKKOS_CPPFLAGS) $(KOKKOS_CXXFLAGS) $(CXXFLAGS) -c $(KOKKOS_PATH)/core/src/Cuda/Kokkos_Cuda_Impl.cpp @@ -40,6 +46,9 @@ Kokkos_Cuda_Task.o: $(KOKKOS_CPP_DEPENDS) $(KOKKOS_PATH)/core/src/Cuda/Kokkos_Cu $(CXX) $(KOKKOS_CPPFLAGS) $(KOKKOS_CXXFLAGS) $(CXXFLAGS) -c $(KOKKOS_PATH)/core/src/Cuda/Kokkos_Cuda_Task.cpp Kokkos_Cuda_Locks.o: $(KOKKOS_CPP_DEPENDS) $(KOKKOS_PATH)/core/src/Cuda/Kokkos_Cuda_Locks.cpp $(CXX) $(KOKKOS_CPPFLAGS) $(KOKKOS_CXXFLAGS) $(CXXFLAGS) -c $(KOKKOS_PATH)/core/src/Cuda/Kokkos_Cuda_Locks.cpp +ifeq ($(KOKKOS_INTERNAL_ENABLE_ETI), 1) + include $(KOKKOS_ETI_PATH)/Cuda/Makefile.eti_Cuda +endif endif ifeq ($(KOKKOS_INTERNAL_USE_ROCM), 1) @@ -51,6 +60,9 @@ Kokkos_ROCm_Task.o: $(KOKKOS_CPP_DEPENDS) $(KOKKOS_PATH)/core/src/ROCm/Kokkos_RO $(CXX) $(KOKKOS_CPPFLAGS) $(KOKKOS_CXXFLAGS) $(CXXFLAGS) -c $(KOKKOS_PATH)/core/src/ROCm/Kokkos_ROCm_Task.cpp Kokkos_ROCm_Impl.o: $(KOKKOS_CPP_DEPENDS) $(KOKKOS_PATH)/core/src/ROCm/Kokkos_ROCm_Impl.cpp $(CXX) $(KOKKOS_CPPFLAGS) $(KOKKOS_CXXFLAGS) $(CXXFLAGS) -c $(KOKKOS_PATH)/core/src/ROCm/Kokkos_ROCm_Impl.cpp +ifeq ($(KOKKOS_INTERNAL_ENABLE_ETI), 1) + include $(KOKKOS_ETI_PATH)/ROCm/Makefile.eti_ROCm +endif endif ifeq ($(KOKKOS_INTERNAL_USE_PTHREADS), 1) @@ -58,6 +70,9 @@ Kokkos_ThreadsExec_base.o: $(KOKKOS_CPP_DEPENDS) $(KOKKOS_PATH)/core/src/Threads $(CXX) $(KOKKOS_CPPFLAGS) $(KOKKOS_CXXFLAGS) $(CXXFLAGS) -c $(KOKKOS_PATH)/core/src/Threads/Kokkos_ThreadsExec_base.cpp Kokkos_ThreadsExec.o: $(KOKKOS_CPP_DEPENDS) $(KOKKOS_PATH)/core/src/Threads/Kokkos_ThreadsExec.cpp $(CXX) $(KOKKOS_CPPFLAGS) $(KOKKOS_CXXFLAGS) $(CXXFLAGS) -c $(KOKKOS_PATH)/core/src/Threads/Kokkos_ThreadsExec.cpp +ifeq ($(KOKKOS_INTERNAL_ENABLE_ETI), 1) + include $(KOKKOS_ETI_PATH)/Threads/Makefile.eti_Threads +endif endif ifeq ($(KOKKOS_INTERNAL_USE_QTHREADS), 1) @@ -72,6 +87,9 @@ Kokkos_OpenMP_Exec.o: $(KOKKOS_CPP_DEPENDS) $(KOKKOS_PATH)/core/src/OpenMP/Kokko $(CXX) $(KOKKOS_CPPFLAGS) $(KOKKOS_CXXFLAGS) $(CXXFLAGS) -c $(KOKKOS_PATH)/core/src/OpenMP/Kokkos_OpenMP_Exec.cpp Kokkos_OpenMP_Task.o: $(KOKKOS_CPP_DEPENDS) $(KOKKOS_PATH)/core/src/OpenMP/Kokkos_OpenMP_Task.cpp $(CXX) $(KOKKOS_CPPFLAGS) $(KOKKOS_CXXFLAGS) $(CXXFLAGS) -c $(KOKKOS_PATH)/core/src/OpenMP/Kokkos_OpenMP_Task.cpp +ifeq ($(KOKKOS_INTERNAL_ENABLE_ETI), 1) + include $(KOKKOS_ETI_PATH)/OpenMP/Makefile.eti_OpenMP +endif endif ifeq ($(KOKKOS_INTERNAL_USE_OPENMPTARGET), 1) diff --git a/lib/kokkos/algorithms/src/Kokkos_Random.hpp b/lib/kokkos/algorithms/src/Kokkos_Random.hpp index 1c659e44a45e5cd1d972f441f956ce4e47de2d2e..5f1d88bfffbfa90dd2a68361d8ef17e853a5bc8a 100644 --- a/lib/kokkos/algorithms/src/Kokkos_Random.hpp +++ b/lib/kokkos/algorithms/src/Kokkos_Random.hpp @@ -702,7 +702,11 @@ namespace Kokkos { } Random_XorShift64_Pool(uint64_t seed) { num_states_ = 0; +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE init(seed,DeviceType::max_hardware_threads()); +#else + init(seed,DeviceType::impl_max_hardware_threads()); +#endif } Random_XorShift64_Pool(const Random_XorShift64_Pool& src): @@ -751,7 +755,11 @@ namespace Kokkos { KOKKOS_INLINE_FUNCTION Random_XorShift64 get_state() const { +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE const int i = DeviceType::hardware_thread_id();; +#else + const int i = DeviceType::impl_hardware_thread_id();; +#endif return Random_XorShift64(state_(i),i); } @@ -957,7 +965,11 @@ namespace Kokkos { inline Random_XorShift1024_Pool(uint64_t seed){ num_states_ = 0; +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE init(seed,DeviceType::max_hardware_threads()); +#else + init(seed,DeviceType::impl_max_hardware_threads()); +#endif } Random_XorShift1024_Pool(const Random_XorShift1024_Pool& src): @@ -1012,7 +1024,11 @@ namespace Kokkos { KOKKOS_INLINE_FUNCTION Random_XorShift1024 get_state() const { +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE const int i = DeviceType::hardware_thread_id(); +#else + const int i = DeviceType::impl_hardware_thread_id(); +#endif return Random_XorShift1024(state_,p_(i),i); }; diff --git a/lib/kokkos/algorithms/src/Kokkos_Sort.hpp b/lib/kokkos/algorithms/src/Kokkos_Sort.hpp index 888476045b8ee4a424e668a5da71567c408af934..c952b1e541af767d346ae2fdd2cea6c16c84a3eb 100644 --- a/lib/kokkos/algorithms/src/Kokkos_Sort.hpp +++ b/lib/kokkos/algorithms/src/Kokkos_Sort.hpp @@ -288,6 +288,7 @@ public: Kokkos::abort("BinSort::sort: values range length != permutation vector length"); } +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE scratch_view_type sorted_values("Scratch", len, @@ -298,6 +299,18 @@ public: values.extent(5), values.extent(6), values.extent(7)); +#else + scratch_view_type + sorted_values("Scratch", + values.rank_dynamic > 0 ? len : KOKKOS_IMPL_CTOR_DEFAULT_ARG, + values.rank_dynamic > 1 ? values.extent(1) : KOKKOS_IMPL_CTOR_DEFAULT_ARG , + values.rank_dynamic > 2 ? values.extent(2) : KOKKOS_IMPL_CTOR_DEFAULT_ARG, + values.rank_dynamic > 3 ? values.extent(3) : KOKKOS_IMPL_CTOR_DEFAULT_ARG, + values.rank_dynamic > 4 ? values.extent(4) : KOKKOS_IMPL_CTOR_DEFAULT_ARG, + values.rank_dynamic > 5 ? values.extent(5) : KOKKOS_IMPL_CTOR_DEFAULT_ARG, + values.rank_dynamic > 6 ? values.extent(6) : KOKKOS_IMPL_CTOR_DEFAULT_ARG, + values.rank_dynamic > 7 ? values.extent(7) : KOKKOS_IMPL_CTOR_DEFAULT_ARG); +#endif { copy_permute_functor< scratch_view_type /* DstViewType */ @@ -362,8 +375,10 @@ public: KOKKOS_INLINE_FUNCTION void operator() (const bin_sort_bins_tag& tag, const int&i ) const { + auto bin_size = bin_count_const(i); + if (bin_size <= 1) return; + int upper_bound = bin_offsets(i)+bin_size; bool sorted = false; - int upper_bound = bin_offsets(i)+bin_count_const(i); while(!sorted) { sorted = true; int old_idx = sort_order(bin_offsets(i)); @@ -501,7 +516,7 @@ bool try_std_sort(ViewType view) { template struct min_max_functor { - typedef Kokkos::Experimental::MinMaxScalar minmax_scalar; + typedef Kokkos::MinMaxScalar minmax_scalar; ViewType view; min_max_functor(const ViewType& view_):view(view_) {} @@ -523,8 +538,8 @@ void sort( ViewType const & view , bool const always_use_kokkos_sort = false) } typedef BinOp1D CompType; - Kokkos::Experimental::MinMaxScalar result; - Kokkos::Experimental::MinMax reducer(result); + Kokkos::MinMaxScalar result; + Kokkos::MinMax reducer(result); parallel_reduce("Kokkos::Sort::FindExtent",Kokkos::RangePolicy(0,view.extent(0)), Impl::min_max_functor(view),reducer); if(result.min_val == result.max_val) return; @@ -542,8 +557,8 @@ void sort( ViewType view typedef Kokkos::RangePolicy range_policy ; typedef BinOp1D CompType; - Kokkos::Experimental::MinMaxScalar result; - Kokkos::Experimental::MinMax reducer(result); + Kokkos::MinMaxScalar result; + Kokkos::MinMax reducer(result); parallel_reduce("Kokkos::Sort::FindExtent", range_policy( begin , end ) , Impl::min_max_functor(view),reducer ); diff --git a/lib/kokkos/cmake/kokkos_build.cmake b/lib/kokkos/cmake/kokkos_build.cmake index 76d0655adb76753d09df3bbfd88ff6d98d7da013..94dd733ca33ecf4d0ea35957ffcf99b9e1dc00b6 100644 --- a/lib/kokkos/cmake/kokkos_build.cmake +++ b/lib/kokkos/cmake/kokkos_build.cmake @@ -76,7 +76,11 @@ IF(KOKKOS_SEPARATE_LIBS) ) foreach(lib IN LISTS KOKKOS_TPL_LIBRARY_NAMES) - find_library(LIB_${lib} ${lib} PATHS ${KOKKOS_TPL_LIBRARY_DIRS}) + if ("${lib}" STREQUAL "cuda") + set(LIB_cuda "-lcuda") + else() + find_library(LIB_${lib} ${lib} PATHS ${KOKKOS_TPL_LIBRARY_DIRS}) + endif() target_link_libraries(kokkoscore PUBLIC ${LIB_${lib}}) endforeach() @@ -154,7 +158,11 @@ ELSE() ) foreach(lib IN LISTS KOKKOS_TPL_LIBRARY_NAMES) - find_library(LIB_${lib} ${lib} PATHS ${KOKKOS_TPL_LIBRARY_DIRS}) + if ("${lib}" STREQUAL "cuda") + set(LIB_cuda "-lcuda") + else() + find_library(LIB_${lib} ${lib} PATHS ${KOKKOS_TPL_LIBRARY_DIRS}) + endif() target_link_libraries(kokkos PUBLIC ${LIB_${lib}}) endforeach() diff --git a/lib/kokkos/cmake/kokkos_options.cmake b/lib/kokkos/cmake/kokkos_options.cmake index 25eb8e86ced2e549eefa1d1c0788ffc41654cbd8..80a091bb98b1d57a24b50f0f22e902133062ecee 100644 --- a/lib/kokkos/cmake/kokkos_options.cmake +++ b/lib/kokkos/cmake/kokkos_options.cmake @@ -31,6 +31,7 @@ list(APPEND KOKKOS_INTERNAL_ENABLE_OPTIONS_LIST Profiling_Load_Print Aggressive_Vectorization Deprecated_Code + Explicit_Instantiation ) #------------------------------------------------------------------------------- @@ -40,6 +41,7 @@ list(APPEND KOKKOS_INTERNAL_ENABLE_OPTIONS_LIST foreach(opt ${KOKKOS_INTERNAL_ENABLE_OPTIONS_LIST}) string(TOUPPER ${opt} OPT ) IF(DEFINED Kokkos_ENABLE_${opt}) + MESSAGE("Kokkos_ENABLE_${opt} is defined!") IF(DEFINED KOKKOS_ENABLE_${OPT}) IF(NOT ("${KOKKOS_ENABLE_${OPT}}" STREQUAL "${Kokkos_ENABLE_${opt}}")) IF(DEFINED KOKKOS_ENABLE_${OPT}_INTERNAL) @@ -57,18 +59,16 @@ foreach(opt ${KOKKOS_INTERNAL_ENABLE_OPTIONS_LIST}) ENDIF() ELSE() SET(KOKKOS_INTERNAL_ENABLE_${OPT}_DEFAULT ${Kokkos_ENABLE_${opt}}) + MESSAGE("set KOKKOS_INTERNAL_ENABLE_${OPT}_DEFAULT!") ENDIF() ENDIF() endforeach() +IF(DEFINED Kokkos_ARCH) + MESSAGE(FATAL_ERROR "Defined Kokkos_ARCH, use KOKKOS_ARCH instead!") +ENDIF() IF(DEFINED Kokkos_Arch) - IF(DEFINED KOKKOS_ARCH) - IF(NOT (${KOKKOS_ARCH} STREQUAL "${Kokkos_Arch}")) - MESSAGE(FATAL_ERROR "Defined both Kokkos_Arch and KOKKOS_ARCH and they differ!") - ENDIF() - ELSE() - SET(KOKKOS_ARCH ${Kokkos_Arch}) - ENDIF() + MESSAGE(FATAL_ERROR "Defined Kokkos_Arch, use KOKKOS_ARCH instead!") ENDIF() #------------------------------------------------------------------------------- @@ -103,6 +103,8 @@ list(APPEND KOKKOS_ARCH_LIST Maxwell53 # (GPU) NVIDIA Maxwell generation CC 5.3 Pascal60 # (GPU) NVIDIA Pascal generation CC 6.0 Pascal61 # (GPU) NVIDIA Pascal generation CC 6.1 + Volta70 # (GPU) NVIDIA Volta generation CC 7.0 + Volta72 # (GPU) NVIDIA Volta generation CC 7.2 ) # List of possible device architectures. @@ -267,6 +269,8 @@ set(KOKKOS_ENABLE_PROFILING_LOAD_PRINT ${KOKKOS_INTERNAL_ENABLE_PROFILING_LOAD_P set_kokkos_default_default(DEPRECATED_CODE ON) set(KOKKOS_ENABLE_DEPRECATED_CODE ${KOKKOS_INTERNAL_ENABLE_DEPRECATED_CODE_DEFAULT} CACHE BOOL "Enable deprecated code.") +set_kokkos_default_default(EXPLICIT_INSTANTIATION ON) +set(KOKKOS_ENABLE_EXPLICIT_INSTANTIATION ${KOKKOS_INTERNAL_ENABLE_EXPLICIT_INSTANTIATION_DEFAULT} CACHE BOOL "Enable explicit template instantiation.") #------------------------------------------------------------------------------- #------------------------------- KOKKOS_USE_TPLS ------------------------------- diff --git a/lib/kokkos/cmake/kokkos_settings.cmake b/lib/kokkos/cmake/kokkos_settings.cmake index 579fab0c954c230d7690cda72a854d332d598be1..21c9d75a96aa52fd349c751aa0c2fa0f12316c91 100644 --- a/lib/kokkos/cmake/kokkos_settings.cmake +++ b/lib/kokkos/cmake/kokkos_settings.cmake @@ -74,6 +74,9 @@ endif() if(${KOKKOS_ENABLE_PROFILING_LOAD_PRINT}) list(APPEND KOKKOS_OPTIONSl enable_profile_load_print) endif() +if(${KOKKOS_ENABLE_EXPLICIT_INSTANTIATION}) + list(APPEND KOKKOS_OPTIONSl enable_eti) +endif() # List needs to be comma-delimitted string(REPLACE ";" "," KOKKOS_GMAKE_OPTIONS "${KOKKOS_OPTIONSl}") @@ -158,6 +161,19 @@ if (NOT "${KOKKOS_INTERNAL_ADDTOPATH}" STREQUAL "") set(KOKKOS_SETTINGS ${KOKKOS_SETTINGS} "PATH=\"${KOKKOS_INTERNAL_ADDTOPATH}:$ENV{PATH}\"") endif() +if (CMAKE_CXX_STANDARD) + if (CMAKE_CXX_STANDARD STREQUAL "98") + message(FATAL_ERROR "Kokkos requires C++11 or newer!") + endif() + set(KOKKOS_CXX_STANDARD "c++${CMAKE_CXX_STANDARD}") + if (CMAKE_CXX_EXTENSIONS) + if (CMAKE_CXX_COMPILER_ID STREQUAL "GNU") + set(KOKKOS_CXX_STANDARD "gnu++${CMAKE_CXX_STANDARD}") + endif() + endif() + set(KOKKOS_SETTINGS ${KOKKOS_SETTINGS} "KOKKOS_CXX_STANDARD=\"${KOKKOS_CXX_STANDARD}\"") +endif() + # Final form that gets passed to make set(KOKKOS_SETTINGS env ${KOKKOS_SETTINGS}) diff --git a/lib/kokkos/cmake/tribits.cmake b/lib/kokkos/cmake/tribits.cmake index 1b5a7b2adb4572cf3b454af49ce7db13332fc7ce..f8eebc29f8b01dab2fe1baa831d30306b0eef376 100644 --- a/lib/kokkos/cmake/tribits.cmake +++ b/lib/kokkos/cmake/tribits.cmake @@ -300,7 +300,9 @@ FUNCTION(TRIBITS_ADD_EXECUTABLE EXE_NAME) ENDIF() ENDFUNCTION() -ADD_CUSTOM_TARGET(check COMMAND ${CMAKE_CTEST_COMMAND} -VV -C ${CMAKE_CFG_INTDIR}) +IF(NOT TARGET check) + ADD_CUSTOM_TARGET(check COMMAND ${CMAKE_CTEST_COMMAND} -VV -C ${CMAKE_CFG_INTDIR}) +ENDIF() FUNCTION(TRIBITS_ADD_TEST) ENDFUNCTION() diff --git a/lib/kokkos/config/test_all_sandia b/lib/kokkos/config/test_all_sandia index 28b4a64b10819538c7f5fc672ada210671343cf4..15e6049afbaf238d132a28335f811576d2d0c72d 100755 --- a/lib/kokkos/config/test_all_sandia +++ b/lib/kokkos/config/test_all_sandia @@ -22,30 +22,38 @@ if [[ "$HOSTNAME" =~ .*bowman.* ]]; then module load git fi -if [[ "$HOSTNAME" =~ n.* ]]; then # Warning: very generic name +if [[ "$HOSTNAME" == n* ]]; then # Warning: very generic name if [[ "$PROCESSOR" = "aarch64" ]]; then MACHINE=sullivan module load git fi fi -if [[ "$HOSTNAME" =~ node.* ]]; then # Warning: very generic name +if [[ "$HOSTNAME" == node* ]]; then # Warning: very generic name if [[ "$MACHINE" = "" ]]; then MACHINE=shepard module load git fi fi -if [[ "$HOSTNAME" =~ apollo ]]; then +if [[ "$HOSTNAME" == apollo\.* ]]; then MACHINE=apollo module load git fi -if [[ "$HOSTNAME" =~ sullivan ]]; then +if [[ "$HOSTNAME" == sullivan ]]; then MACHINE=sullivan module load git fi +if [[ "$HOSTNAME" == mayer\.* ]]; then + MACHINE=mayer +# module load git +fi +if [[ "$HOSTNAME" == cn* ]]; then # Warning: very generic name + MACHINE=mayer +fi + if [ ! -z "$SEMS_MODULEFILES_ROOT" ]; then if [[ "$MACHINE" = "" ]]; then MACHINE=sems @@ -83,7 +91,7 @@ CUSTOM_BUILD_LIST="" QTHREADS_PATH="" DRYRUN=False BUILD_ONLY=False -declare -i NUM_JOBS_TO_RUN_IN_PARALLEL=3 +declare -i NUM_JOBS_TO_RUN_IN_PARALLEL=1 TEST_SCRIPT=False SKIP_HWLOC=False SPOT_CHECK=False @@ -142,6 +150,9 @@ do --with-cuda-options*) KOKKOS_CUDA_OPTIONS="--with-cuda-options=${key#*=}" ;; + --with-options*) + KOKKOS_OPTIONS="--with-options=enable_large_mem_tests,${key#*=}" + ;; --cxxflags-extra*) CXX_FLAGS_EXTRA="${key#*=}" ;; @@ -247,7 +258,7 @@ elif [ "$MACHINE" = "white" ]; then ARCH_FLAG="--arch=Power8,Kepler37" fi - NUM_JOBS_TO_RUN_IN_PARALLEL=2 + NUM_JOBS_TO_RUN_IN_PARALLEL=1 elif [ "$MACHINE" = "bowman" ]; then source /etc/profile.d/modules.sh @@ -268,7 +279,7 @@ elif [ "$MACHINE" = "bowman" ]; then ARCH_FLAG="--arch=KNL" fi - NUM_JOBS_TO_RUN_IN_PARALLEL=2 + NUM_JOBS_TO_RUN_IN_PARALLEL=1 elif [ "$MACHINE" = "sullivan" ]; then source /etc/profile.d/modules.sh @@ -284,7 +295,24 @@ elif [ "$MACHINE" = "sullivan" ]; then ARCH_FLAG="--arch=ARMv8-ThunderX" fi - NUM_JOBS_TO_RUN_IN_PARALLEL=2 + NUM_JOBS_TO_RUN_IN_PARALLEL=1 + +elif [ "$MACHINE" = "mayer" ]; then + SKIP_HWLOC=True + export SLURM_TASKS_PER_NODE=96 + + BASE_MODULE_LIST="/" + ARM_MODULE_LIST="/" + + # Format: (compiler module-list build-list exe-name warning-flag) + COMPILERS=("gcc/7.2.0 $BASE_MODULE_LIST $ARM_GCC_BUILD_LIST g++ $GCC_WARNING_FLAGS" + "arm/1.4.0 $ARM_MODULE_LIST $ARM_GCC_BUILD_LIST armclang++ $CLANG_WARNING_FLAGS") + + if [ -z "$ARCH_FLAG" ]; then + ARCH_FLAG="--arch=ARMv8-TX2" + fi + + NUM_JOBS_TO_RUN_IN_PARALLEL=1 elif [ "$MACHINE" = "shepard" ]; then source /etc/profile.d/modules.sh @@ -303,7 +331,7 @@ elif [ "$MACHINE" = "shepard" ]; then if [ -z "$ARCH_FLAG" ]; then ARCH_FLAG="--arch=HSW" fi - NUM_JOBS_TO_RUN_IN_PARALLEL=2 + NUM_JOBS_TO_RUN_IN_PARALLEL=1 elif [ "$MACHINE" = "apollo" ]; then source /projects/sems/modulefiles/utils/sems-modules-init.sh @@ -331,7 +359,7 @@ elif [ "$MACHINE" = "apollo" ]; then if [ "$SPOT_CHECK" = "True" ]; then # Format: (compiler module-list build-list exe-name warning-flag) COMPILERS=("gcc/4.8.4 $BASE_MODULE_LIST "OpenMP,Pthread" g++ $GCC_WARNING_FLAGS" - "gcc/5.1.0 $BASE_MODULE_LIST "Serial" g++ $GCC_WARNING_FLAGS" + "gcc/5.3.0 $BASE_MODULE_LIST "Serial" g++ $GCC_WARNING_FLAGS" "intel/16.0.1 $BASE_MODULE_LIST "OpenMP" icpc $INTEL_WARNING_FLAGS" "clang/3.9.0 $BASE_MODULE_LIST "Pthread_Serial" clang++ $CLANG_WARNING_FLAGS" "clang/6.0 $CLANG_MODULE_LIST "Cuda_Pthread" clang++ $CUDA_WARNING_FLAGS" @@ -358,7 +386,7 @@ elif [ "$MACHINE" = "apollo" ]; then ARCH_FLAG="--arch=SNB,Volta70" fi - NUM_JOBS_TO_RUN_IN_PARALLEL=2 + NUM_JOBS_TO_RUN_IN_PARALLEL=1 else echo "Unhandled machine $MACHINE" >&2 @@ -627,6 +655,11 @@ single_build_and_test() { if [[ "$KOKKOS_CUDA_OPTIONS" != "" ]]; then local extra_args="$extra_args $KOKKOS_CUDA_OPTIONS" fi + if [[ "$KOKKOS_OPTIONS" != "" ]]; then + local extra_args="$extra_args $KOKKOS_OPTIONS" + else + local extra_args="$extra_args --with-options=enable_large_mem_tests" + fi echo " Starting job $desc" @@ -642,7 +675,7 @@ single_build_and_test() { else run_cmd ${KOKKOS_PATH}/generate_makefile.bash --with-devices=$build $ARCH_FLAG --compiler=$(which $compiler_exe) --cxxflags=\"$cxxflags\" --ldflags=\"$ldflags\" $extra_args &>> ${desc}.configure.log || { report_and_log_test_result 1 ${desc} configure && return 0; } local -i build_start_time=$(date +%s) - run_cmd make -j 32 build-test >& ${desc}.build.log || { report_and_log_test_result 1 ${desc} build && return 0; } + run_cmd make -j 48 build-test >& ${desc}.build.log || { report_and_log_test_result 1 ${desc} build && return 0; } local -i build_end_time=$(date +%s) comment="build_time=$(($build_end_time-$build_start_time))" @@ -682,6 +715,9 @@ run_in_background() { if [[ "$compiler" == cuda* ]]; then num_jobs=1 fi + if [[ "$compiler" == clang ]]; then + num_jobs=1 + fi # fi wait_for_jobs $num_jobs diff --git a/lib/kokkos/containers/performance_tests/TestCuda.cpp b/lib/kokkos/containers/performance_tests/TestCuda.cpp index 682f3f52f70a6d7c37ef1006267d3bddffc9cd70..351fb86df3d48836ff45286401e52a6f1f684fd7 100644 --- a/lib/kokkos/containers/performance_tests/TestCuda.cpp +++ b/lib/kokkos/containers/performance_tests/TestCuda.cpp @@ -70,13 +70,12 @@ protected: static void SetUpTestCase() { std::cout << std::setprecision(5) << std::scientific; - Kokkos::HostSpace::execution_space::initialize(); - Kokkos::Cuda::initialize( Kokkos::Cuda::SelectDevice(0) ); + Kokkos::InitArguments args(-1, -1, 0); + Kokkos::initialize(args); } static void TearDownTestCase() { - Kokkos::Cuda::finalize(); - Kokkos::HostSpace::execution_space::finalize(); + Kokkos::finalize(); } }; diff --git a/lib/kokkos/containers/performance_tests/TestOpenMP.cpp b/lib/kokkos/containers/performance_tests/TestOpenMP.cpp index 66d497552e3de2dc09cd409bed204c35bc0ddb68..e6218074eafd965554c3ebe60136386c488fd882 100644 --- a/lib/kokkos/containers/performance_tests/TestOpenMP.cpp +++ b/lib/kokkos/containers/performance_tests/TestOpenMP.cpp @@ -70,13 +70,13 @@ protected: { std::cout << std::setprecision(5) << std::scientific; - Kokkos::OpenMP::initialize(); + Kokkos::initialize(); Kokkos::OpenMP::print_configuration( std::cout ); } static void TearDownTestCase() { - Kokkos::OpenMP::finalize(); + Kokkos::finalize(); } }; diff --git a/lib/kokkos/containers/performance_tests/TestThreads.cpp b/lib/kokkos/containers/performance_tests/TestThreads.cpp index a951a5ca56e4f30fcbe39d0ae797f6bd0d126882..6a02e67b252781f7320223ad8fb87cadfef3355e 100644 --- a/lib/kokkos/containers/performance_tests/TestThreads.cpp +++ b/lib/kokkos/containers/performance_tests/TestThreads.cpp @@ -81,12 +81,12 @@ protected: std::cout << "Threads: " << num_threads << std::endl; - Kokkos::Threads::initialize( num_threads ); + Kokkos::initialize( Kokkos::InitArguments(num_threads) ); } static void TearDownTestCase() { - Kokkos::Threads::finalize(); + Kokkos::finalize(); } }; diff --git a/lib/kokkos/containers/src/Kokkos_Bitset.hpp b/lib/kokkos/containers/src/Kokkos_Bitset.hpp index c48058d75d774feaf424cf196522dfceadf9e150..bfe8080f3b92ce9d11a0c040c123745acaf7e66b 100644 --- a/lib/kokkos/containers/src/Kokkos_Bitset.hpp +++ b/lib/kokkos/containers/src/Kokkos_Bitset.hpp @@ -271,7 +271,7 @@ private: block = Impl::rotate_right(block, offset); return ((( !(scan_direction & BIT_SCAN_REVERSE) ? Impl::bit_scan_forward(block) : - Impl::bit_scan_reverse(block) + ::Kokkos::log2(block) ) + offset ) & block_mask ) + block_start; diff --git a/lib/kokkos/containers/src/Kokkos_DualView.hpp b/lib/kokkos/containers/src/Kokkos_DualView.hpp index 74fe4418f870854498cd65cfe8e346d7d183e716..548e96d251f13224f4b82e13dda1c233fe36a577 100644 --- a/lib/kokkos/containers/src/Kokkos_DualView.hpp +++ b/lib/kokkos/containers/src/Kokkos_DualView.hpp @@ -209,14 +209,14 @@ public: /// the first three integer arguments will be nonzero, and you may /// omit the integer arguments that follow. DualView (const std::string& label, - const size_t n0 = 0, - const size_t n1 = 0, - const size_t n2 = 0, - const size_t n3 = 0, - const size_t n4 = 0, - const size_t n5 = 0, - const size_t n6 = 0, - const size_t n7 = 0) + const size_t n0 = KOKKOS_IMPL_CTOR_DEFAULT_ARG, + const size_t n1 = KOKKOS_IMPL_CTOR_DEFAULT_ARG, + const size_t n2 = KOKKOS_IMPL_CTOR_DEFAULT_ARG, + const size_t n3 = KOKKOS_IMPL_CTOR_DEFAULT_ARG, + const size_t n4 = KOKKOS_IMPL_CTOR_DEFAULT_ARG, + const size_t n5 = KOKKOS_IMPL_CTOR_DEFAULT_ARG, + const size_t n6 = KOKKOS_IMPL_CTOR_DEFAULT_ARG, + const size_t n7 = KOKKOS_IMPL_CTOR_DEFAULT_ARG) : d_view (label, n0, n1, n2, n3, n4, n5, n6, n7) , h_view (create_mirror_view (d_view)) // without UVM, host View mirrors , modified_device (View ("DualView::modified_device")) @@ -464,14 +464,14 @@ public: /// This discards any existing contents of the objects, and resets /// their modified flags. It does not copy the old contents /// of either View into the new View objects. - void realloc( const size_t n0 = 0 , - const size_t n1 = 0 , - const size_t n2 = 0 , - const size_t n3 = 0 , - const size_t n4 = 0 , - const size_t n5 = 0 , - const size_t n6 = 0 , - const size_t n7 = 0 ) { + void realloc( const size_t n0 = KOKKOS_IMPL_CTOR_DEFAULT_ARG , + const size_t n1 = KOKKOS_IMPL_CTOR_DEFAULT_ARG , + const size_t n2 = KOKKOS_IMPL_CTOR_DEFAULT_ARG , + const size_t n3 = KOKKOS_IMPL_CTOR_DEFAULT_ARG , + const size_t n4 = KOKKOS_IMPL_CTOR_DEFAULT_ARG , + const size_t n5 = KOKKOS_IMPL_CTOR_DEFAULT_ARG , + const size_t n6 = KOKKOS_IMPL_CTOR_DEFAULT_ARG , + const size_t n7 = KOKKOS_IMPL_CTOR_DEFAULT_ARG ) { ::Kokkos::realloc(d_view,n0,n1,n2,n3,n4,n5,n6,n7); h_view = create_mirror_view( d_view ); @@ -483,14 +483,14 @@ public: /// /// This method only copies the old contents into the new View /// objects for the device which was last marked as modified. - void resize( const size_t n0 = 0 , - const size_t n1 = 0 , - const size_t n2 = 0 , - const size_t n3 = 0 , - const size_t n4 = 0 , - const size_t n5 = 0 , - const size_t n6 = 0 , - const size_t n7 = 0 ) { + void resize( const size_t n0 = KOKKOS_IMPL_CTOR_DEFAULT_ARG , + const size_t n1 = KOKKOS_IMPL_CTOR_DEFAULT_ARG , + const size_t n2 = KOKKOS_IMPL_CTOR_DEFAULT_ARG , + const size_t n3 = KOKKOS_IMPL_CTOR_DEFAULT_ARG , + const size_t n4 = KOKKOS_IMPL_CTOR_DEFAULT_ARG , + const size_t n5 = KOKKOS_IMPL_CTOR_DEFAULT_ARG , + const size_t n6 = KOKKOS_IMPL_CTOR_DEFAULT_ARG , + const size_t n7 = KOKKOS_IMPL_CTOR_DEFAULT_ARG ) { if(modified_device() >= modified_host()) { /* Resize on Device */ ::Kokkos::resize(d_view,n0,n1,n2,n3,n4,n5,n6,n7); @@ -533,10 +533,21 @@ public: //! \name Methods for getting capacity, stride, or dimension(s). //@{ +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE //! The allocation size (same as Kokkos::View::capacity). size_t capacity() const { return d_view.span(); } +#endif + + //! The allocation size (same as Kokkos::View::span). + KOKKOS_INLINE_FUNCTION constexpr size_t span() const { + return d_view.span(); + } + + KOKKOS_INLINE_FUNCTION bool span_is_contiguous() const { + return d_view.span_is_contiguous(); + } //! Get stride(s) for each dimension. template< typename iType> @@ -556,6 +567,11 @@ public: extent_int( const iType & r ) const { return static_cast(d_view.extent(r)); } +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE + /* Deprecate all 'dimension' functions in favor of + * ISO/C++ vocabulary 'extent'. + */ + /* \brief return size of dimension 0 */ size_t dimension_0() const {return d_view.extent(0);} /* \brief return size of dimension 1 */ @@ -572,6 +588,7 @@ public: size_t dimension_6() const {return d_view.extent(6);} /* \brief return size of dimension 7 */ size_t dimension_7() const {return d_view.extent(7);} +#endif //@} }; diff --git a/lib/kokkos/containers/src/Kokkos_DynRankView.hpp b/lib/kokkos/containers/src/Kokkos_DynRankView.hpp index ccf53b3d5096574f4446d4cf9d7eee71b6e45c08..b30009a999e108235a1318431df08fd3115ec965 100644 --- a/lib/kokkos/containers/src/Kokkos_DynRankView.hpp +++ b/lib/kokkos/containers/src/Kokkos_DynRankView.hpp @@ -64,7 +64,7 @@ namespace Impl { template struct DynRankDimTraits { - enum : size_t{unspecified = ~size_t(0)}; + enum : size_t{unspecified =KOKKOS_INVALID_INDEX}; // Compute the rank of the view from the nonzero dimension arguments. KOKKOS_INLINE_FUNCTION @@ -192,14 +192,14 @@ struct DynRankDimTraits { static typename std::enable_if< (std::is_same::value || std::is_same::value) && std::is_integral::value , Layout >::type reconstructLayout( const Layout& layout , iType dynrank ) { - return Layout( dynrank > 0 ? layout.dimension[0] : ~size_t(0) - , dynrank > 1 ? layout.dimension[1] : ~size_t(0) - , dynrank > 2 ? layout.dimension[2] : ~size_t(0) - , dynrank > 3 ? layout.dimension[3] : ~size_t(0) - , dynrank > 4 ? layout.dimension[4] : ~size_t(0) - , dynrank > 5 ? layout.dimension[5] : ~size_t(0) - , dynrank > 6 ? layout.dimension[6] : ~size_t(0) - , dynrank > 7 ? layout.dimension[7] : ~size_t(0) + return Layout( dynrank > 0 ? layout.dimension[0] :KOKKOS_INVALID_INDEX + , dynrank > 1 ? layout.dimension[1] :KOKKOS_INVALID_INDEX + , dynrank > 2 ? layout.dimension[2] :KOKKOS_INVALID_INDEX + , dynrank > 3 ? layout.dimension[3] :KOKKOS_INVALID_INDEX + , dynrank > 4 ? layout.dimension[4] :KOKKOS_INVALID_INDEX + , dynrank > 5 ? layout.dimension[5] :KOKKOS_INVALID_INDEX + , dynrank > 6 ? layout.dimension[6] :KOKKOS_INVALID_INDEX + , dynrank > 7 ? layout.dimension[7] :KOKKOS_INVALID_INDEX ); } @@ -209,21 +209,21 @@ struct DynRankDimTraits { static typename std::enable_if< (std::is_same::value) && std::is_integral::value , Layout >::type reconstructLayout( const Layout& layout , iType dynrank ) { - return Layout( dynrank > 0 ? layout.dimension[0] : ~size_t(0) + return Layout( dynrank > 0 ? layout.dimension[0] :KOKKOS_INVALID_INDEX , dynrank > 0 ? layout.stride[0] : (0) - , dynrank > 1 ? layout.dimension[1] : ~size_t(0) + , dynrank > 1 ? layout.dimension[1] :KOKKOS_INVALID_INDEX , dynrank > 1 ? layout.stride[1] : (0) - , dynrank > 2 ? layout.dimension[2] : ~size_t(0) + , dynrank > 2 ? layout.dimension[2] :KOKKOS_INVALID_INDEX , dynrank > 2 ? layout.stride[2] : (0) - , dynrank > 3 ? layout.dimension[3] : ~size_t(0) + , dynrank > 3 ? layout.dimension[3] :KOKKOS_INVALID_INDEX , dynrank > 3 ? layout.stride[3] : (0) - , dynrank > 4 ? layout.dimension[4] : ~size_t(0) + , dynrank > 4 ? layout.dimension[4] :KOKKOS_INVALID_INDEX , dynrank > 4 ? layout.stride[4] : (0) - , dynrank > 5 ? layout.dimension[5] : ~size_t(0) + , dynrank > 5 ? layout.dimension[5] :KOKKOS_INVALID_INDEX , dynrank > 5 ? layout.stride[5] : (0) - , dynrank > 6 ? layout.dimension[6] : ~size_t(0) + , dynrank > 6 ? layout.dimension[6] :KOKKOS_INVALID_INDEX , dynrank > 6 ? layout.stride[6] : (0) - , dynrank > 7 ? layout.dimension[7] : ~size_t(0) + , dynrank > 7 ? layout.dimension[7] :KOKKOS_INVALID_INDEX , dynrank > 7 ? layout.stride[7] : (0) ); } @@ -501,6 +501,7 @@ public: * ISO/C++ vocabulary 'extent'. */ +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE template< typename iType > KOKKOS_INLINE_FUNCTION constexpr typename std::enable_if< std::is_integral::value , size_t >::type @@ -514,17 +515,18 @@ public: KOKKOS_INLINE_FUNCTION constexpr size_t dimension_5() const { return m_map.dimension_5(); } KOKKOS_INLINE_FUNCTION constexpr size_t dimension_6() const { return m_map.dimension_6(); } KOKKOS_INLINE_FUNCTION constexpr size_t dimension_7() const { return m_map.dimension_7(); } +#endif //---------------------------------------- - KOKKOS_INLINE_FUNCTION constexpr size_t size() const { return m_map.dimension_0() * - m_map.dimension_1() * - m_map.dimension_2() * - m_map.dimension_3() * - m_map.dimension_4() * - m_map.dimension_5() * - m_map.dimension_6() * - m_map.dimension_7(); } + KOKKOS_INLINE_FUNCTION constexpr size_t size() const { return m_map.extent(0) * + m_map.extent(1) * + m_map.extent(2) * + m_map.extent(3) * + m_map.extent(4) * + m_map.extent(5) * + m_map.extent(6) * + m_map.extent(7); } KOKKOS_INLINE_FUNCTION constexpr size_t stride_0() const { return m_map.stride_0(); } KOKKOS_INLINE_FUNCTION constexpr size_t stride_1() const { return m_map.stride_1(); } @@ -547,15 +549,19 @@ public: enum { reference_type_is_lvalue_reference = std::is_lvalue_reference< reference_type >::value }; KOKKOS_INLINE_FUNCTION constexpr size_t span() const { return m_map.span(); } +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE // Deprecated, use 'span()' instead KOKKOS_INLINE_FUNCTION constexpr size_t capacity() const { return m_map.span(); } +#endif KOKKOS_INLINE_FUNCTION constexpr bool span_is_contiguous() const { return m_map.span_is_contiguous(); } KOKKOS_INLINE_FUNCTION constexpr pointer_type data() const { return m_map.data(); } +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE // Deprecated, use 'span_is_contigous()' instead KOKKOS_INLINE_FUNCTION constexpr bool is_contiguous() const { return m_map.span_is_contiguous(); } // Deprecated, use 'data()' instead KOKKOS_INLINE_FUNCTION constexpr pointer_type ptr_on_device() const { return m_map.data(); } +#endif //---------------------------------------- // Allow specializations to query their specialized map @@ -998,7 +1004,7 @@ public: //---------------------------------------- // Allocation according to allocation properties and array layout - // unused arg_layout dimensions must be set to ~size_t(0) so that rank deduction can properly take place + // unused arg_layout dimensions must be set toKOKKOS_INVALID_INDEX so that rank deduction can properly take place template< class ... P > explicit inline DynRankView( const Kokkos::Impl::ViewCtorProp< P ... > & arg_prop @@ -1038,7 +1044,12 @@ public: , "View allocation constructor requires managed memory" ); if ( alloc_prop::initialize && - ! alloc_prop::execution_space::is_initialized() ) { +#ifdef KOKKOS_ENABLE_DEPRECATED_CODE + ! alloc_prop::execution_space::is_initialized() +#else + ! alloc_prop::execution_space::impl_is_initialized() +#endif + ) { // If initializing view data then // the execution space must be initialized. Kokkos::Impl::throw_runtime_exception("Constructing DynRankView and initializing data with uninitialized execution space"); @@ -1104,14 +1115,14 @@ public: DynRankView( const Kokkos::Impl::ViewCtorProp< P ... > & arg_prop , typename std::enable_if< ! Kokkos::Impl::ViewCtorProp< P... >::has_pointer , size_t - >::type const arg_N0 = ~size_t(0) - , const size_t arg_N1 = ~size_t(0) - , const size_t arg_N2 = ~size_t(0) - , const size_t arg_N3 = ~size_t(0) - , const size_t arg_N4 = ~size_t(0) - , const size_t arg_N5 = ~size_t(0) - , const size_t arg_N6 = ~size_t(0) - , const size_t arg_N7 = ~size_t(0) + >::type const arg_N0 =KOKKOS_INVALID_INDEX + , const size_t arg_N1 =KOKKOS_INVALID_INDEX + , const size_t arg_N2 =KOKKOS_INVALID_INDEX + , const size_t arg_N3 =KOKKOS_INVALID_INDEX + , const size_t arg_N4 =KOKKOS_INVALID_INDEX + , const size_t arg_N5 =KOKKOS_INVALID_INDEX + , const size_t arg_N6 =KOKKOS_INVALID_INDEX + , const size_t arg_N7 =KOKKOS_INVALID_INDEX ) : DynRankView( arg_prop , typename traits::array_layout @@ -1124,14 +1135,14 @@ public: DynRankView( const Kokkos::Impl::ViewCtorProp< P ... > & arg_prop , typename std::enable_if< Kokkos::Impl::ViewCtorProp< P... >::has_pointer , size_t - >::type const arg_N0 = ~size_t(0) - , const size_t arg_N1 = ~size_t(0) - , const size_t arg_N2 = ~size_t(0) - , const size_t arg_N3 = ~size_t(0) - , const size_t arg_N4 = ~size_t(0) - , const size_t arg_N5 = ~size_t(0) - , const size_t arg_N6 = ~size_t(0) - , const size_t arg_N7 = ~size_t(0) + >::type const arg_N0 =KOKKOS_INVALID_INDEX + , const size_t arg_N1 =KOKKOS_INVALID_INDEX + , const size_t arg_N2 =KOKKOS_INVALID_INDEX + , const size_t arg_N3 =KOKKOS_INVALID_INDEX + , const size_t arg_N4 =KOKKOS_INVALID_INDEX + , const size_t arg_N5 =KOKKOS_INVALID_INDEX + , const size_t arg_N6 =KOKKOS_INVALID_INDEX + , const size_t arg_N7 =KOKKOS_INVALID_INDEX ) : DynRankView( arg_prop , typename traits::array_layout @@ -1156,14 +1167,14 @@ public: DynRankView( const Label & arg_label , typename std::enable_if< Kokkos::Impl::is_view_label