diff --git a/CMakeLists.txt b/CMakeLists.txt index 27355f27..21398b23 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -7,6 +7,10 @@ endif () set(CMAKE_CXX_STANDARD 14) set(CMAKE_CUDA_STANDARD 14) +if (CUDA_ON_BACKEND STREQUAL "METAL") + set(CMAKE_CXX_STANDARD 17) + set(CMAKE_CXX_STANDARD_REQUIRED ON) +endif() ###### CONFIG.h FILE ###### @@ -29,4 +33,3 @@ if(hasParent) set(DEFINE_HAVE_VCDEVEL ${DEFINE_HAVE_VCDEVEL} CACHE INTERNAL "") set(DEFINE_HAVE_LIBHILBERT ${DEFINE_HAVE_LIBHILBERT} CACHE INTERNAL "") endif() - diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 691fb6f6..153850b6 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -20,7 +20,7 @@ if (TEST_PERFORMANCE) endif () -if (NOT CUDA_ON_BACKEND STREQUAL "None") +if (NOT CUDA_ON_BACKEND STREQUAL "NONE") set(CUDA_SOURCES ${CUDA_SOURCES} Vector/map_vector_sparse_unit_tests.cu Vector/vector_gpu_unit_tests.cu @@ -218,6 +218,54 @@ if (NOT APPLE) target_link_libraries(mem_map rt) endif () +# The Metal backend deliberately consumes the same HIP/CUDA translation units +# as CUDA and HIP. Compile both halves of every non-performance .cu source: +# HIP clang emits the ordinary host-side Boost tests and kernel handles, while +# openfpm_add_moltenvk_kernels lowers the instantiated device kernels to +# embedded Vulkan SPIR-V for MoltenVK. Keeping CUDA_SOURCES authoritative here +# prevents a second Metal-only algorithm/test catalogue from drifting away from +# the core implementation. +if (CUDA_ON_BACKEND STREQUAL "METAL") + include("${CMAKE_CURRENT_SOURCE_DIR}/../../openfpm_devices/cmake/OpenFPMMoltenVK.cmake") + + set(_openfpm_data_metal_include_directories + "${CMAKE_CURRENT_SOURCE_DIR}" + "${CMAKE_CURRENT_SOURCE_DIR}/../../openfpm_devices/src" + "${CMAKE_CURRENT_SOURCE_DIR}/../../openfpm_io/src" + "${CMAKE_BINARY_DIR}/config" + ${VULKAN_INCLUDE_DIR} + ${LIBHILBERT_INCLUDE_DIRS} + ${Boost_INCLUDE_DIRS} + ${Vc_INCLUDE_DIR} + ${MPI_C_INCLUDE_DIRS} + ${MPI_CXX_INCLUDE_DIRS} + ${HDF5_INCLUDE_DIRS}) + if (ALPAKA_ROOT) + list(APPEND _openfpm_data_metal_include_directories + "${ALPAKA_ROOT}/include") + endif() + list(REMOVE_DUPLICATES _openfpm_data_metal_include_directories) + + foreach(_openfpm_data_cuda_source IN LISTS CUDA_SOURCES) + if (_openfpm_data_cuda_source MATCHES "\\.cu$") + openfpm_add_moltenvk_kernels( + TARGET mem_map + SOURCE "${_openfpm_data_cuda_source}" + HIP_CLANG "${OPENFPM_MOLTENVK_HIP_CLANG}" + LLVM_OPT "${OPENFPM_MOLTENVK_LLVM_OPT}" + LLVM_LINK "${OPENFPM_MOLTENVK_LLVM_LINK}" + CLSPV "${OPENFPM_MOLTENVK_CLSPV}" + SPIRV_VAL "${OPENFPM_MOLTENVK_SPIRV_VAL}" + ABI_TOOL "${OPENFPM_MOLTENVK_ABI_TOOL}" + CHIPSTAR_ROOT "${OPENFPM_CHIPSTAR_ROOT}" + CHIPSTAR_GENERATED_INCLUDE "${OPENFPM_CHIPSTAR_GENERATED_INCLUDE}" + COMPILE_DEFINITIONS CUDIFY_USE_METAL CUDA_GPU + BOOST_ALLOW_DEPRECATED_HEADERS + INCLUDE_DIRECTORIES ${_openfpm_data_metal_include_directories}) + endif() + endforeach() +endif() + install(FILES Grid/comb.hpp Grid/copy_grid_fast.hpp Grid/grid_base_implementation.hpp @@ -486,4 +534,3 @@ install (FILES SparseGridGpu/TemplateUtils/mathUtils.hpp SparseGridGpu/TemplateUtils/encap_shmem.hpp DESTINATION openfpm_data/include/SparseGridGpu/TemplateUtils/ COMPONENT OpenFPM) - diff --git a/src/Grid/cuda/cuda_grid_gpu_funcs.cuh b/src/Grid/cuda/cuda_grid_gpu_funcs.cuh index 26a48869..31a05061 100644 --- a/src/Grid/cuda/cuda_grid_gpu_funcs.cuh +++ b/src/Grid/cuda/cuda_grid_gpu_funcs.cuh @@ -12,7 +12,7 @@ #include "util/cuda_util.hpp" #include "map_grid_cuda_ker.cuh" -#if defined(CUDA_GPU) && defined(__NVCC__) +#if defined(CUDA_GPU) && (defined(__NVCC__) || defined(CUDIFY_USE_METAL)) template __global__ void copy_ndim_grid_block_device(grid_type src, grid_type dst) diff --git a/src/Grid/cuda/cuda_grid_gpu_tests.cu b/src/Grid/cuda/cuda_grid_gpu_tests.cu index 91248b3d..46043b29 100644 --- a/src/Grid/cuda/cuda_grid_gpu_tests.cu +++ b/src/Grid/cuda/cuda_grid_gpu_tests.cu @@ -184,15 +184,15 @@ BOOST_AUTO_TEST_CASE (gpu_computation_lambda) blockIdx.y * blockDim.y + threadIdx.y, blockIdx.z * blockDim.z + threadIdx.z}); - c3_k.template get<0>(p) = 5.0; + c3_k.template get<0>(p) = 5.0f; - c3_k.template get<1>(p)[0] = 5.0; - c3_k.template get<1>(p)[1] = 5.0; + c3_k.template get<1>(p)[0] = 5.0f; + c3_k.template get<1>(p)[1] = 5.0f; - c3_k.template get<2>(p)[0][0] = 5.0; - c3_k.template get<2>(p)[0][1] = 5.0; - c3_k.template get<2>(p)[1][0] = 5.0; - c3_k.template get<2>(p)[1][1] = 5.0; + c3_k.template get<2>(p)[0][0] = 5.0f; + c3_k.template get<2>(p)[0][1] = 5.0f; + c3_k.template get<2>(p)[1][0] = 5.0f; + c3_k.template get<2>(p)[1][1] = 5.0f; }; auto ite = c3.getGPUIterator({0,0,0},{63,63,63}); @@ -700,7 +700,7 @@ __global__ void test_se1_crash_gt2(grid_type gt1, grid_type gt2) { grid_key_dx<2> k({10000,12345}); - gt1.template get<1>(k)[2] = 6.0; + gt1.template get<1>(k)[2] = 6.0f; } } @@ -709,7 +709,7 @@ __global__ void test_se1_crash_gt3(grid_type gt1, grid_type gt2) { grid_key_dx<2> k({10000,12345}); - gt1.template get<2>(k)[2][2] = 6.0; + gt1.template get<2>(k)[2][2] = 6.0f; } BOOST_AUTO_TEST_CASE (gpu_grid_test_se_class1) diff --git a/src/Grid/cuda/cuda_grid_unit_tests_func.cu b/src/Grid/cuda/cuda_grid_unit_tests_func.cu index c43741cd..9758a36a 100644 --- a/src/Grid/cuda/cuda_grid_unit_tests_func.cu +++ b/src/Grid/cuda/cuda_grid_unit_tests_func.cu @@ -9,9 +9,9 @@ __global__ void grid_gradient_vector(grid_type1 g1, grid_type2 g2, ite_gpu<3> it { GRID_ID_3(ite_gpu); - g2.template get<4>(key)[0] = (g1.template get<0>(key.move(0,1)) - g1.template get<0>(key.move(0,-1))) / 2.0; - g2.template get<4>(key)[1] = (g1.template get<0>(key.move(1,1)) - g1.template get<0>(key.move(1,-1))) / 2.0; - g2.template get<4>(key)[2] = (g1.template get<0>(key.move(2,1)) - g1.template get<0>(key.move(2,-1))) / 2.0; + g2.template get<4>(key)[0] = (g1.template get<0>(key.move(0,1)) - g1.template get<0>(key.move(0,-1))) / 2.0f; + g2.template get<4>(key)[1] = (g1.template get<0>(key.move(1,1)) - g1.template get<0>(key.move(1,-1))) / 2.0f; + g2.template get<4>(key)[2] = (g1.template get<0>(key.move(2,1)) - g1.template get<0>(key.move(2,-1))) / 2.0f; } template @@ -19,9 +19,9 @@ __global__ void grid_fill_vector(grid_type g1, ite_gpu<3> ite_gpu) { GRID_ID_3(ite_gpu); - g1.template get<4>(key)[0] = 1.0; - g1.template get<4>(key)[1] = 2.0; - g1.template get<4>(key)[2] = 3.0; + g1.template get<4>(key)[0] = 1.0f; + g1.template get<4>(key)[1] = 2.0f; + g1.template get<4>(key)[2] = 3.0f; } template @@ -29,9 +29,9 @@ __global__ void grid_fill_vector2(grid_type g1, ite_gpu<3> ite_gpu) { GRID_ID_3(ite_gpu); - g1.template get<4>(key)[0] = 1001.0; - g1.template get<4>(key)[1] = 1002.0; - g1.template get<4>(key)[2] = 1003.0; + g1.template get<4>(key)[0] = 1001.0f; + g1.template get<4>(key)[1] = 1002.0f; + g1.template get<4>(key)[2] = 1003.0f; } template @@ -42,7 +42,7 @@ __global__ void compute_stencil_grid(grid_type g1, grid_type g2, ite_gpu<3> ite_ g2.template get<0>(key) = g1.template get<0>(key.move(0,1)) + g1.template get<0>(key.move(0,-1)) + g1.template get<0>(key.move(1,1)) + g1.template get<0>(key.move(1,-1)) + g1.template get<0>(key.move(2,1)) + g1.template get<0>(key.move(2,-1)) - - 6.0*g1.template get<0>(key); + 6.0f*g1.template get<0>(key); } __global__ void compute_stencil(float * prp_0, float * prp_1, int sz, grid_key_dx<3> start, grid_key_dx<3> stop) @@ -52,7 +52,7 @@ __global__ void compute_stencil(float * prp_0, float * prp_1, int sz, grid_key_d prp_1[tz*sz*sz + ty*sz + tx] = prp_0[tz*sz*sz + ty*sz + tx + 1] + prp_0[tz*sz*sz + ty*sz + tx - 1] + prp_0[tz*sz*sz + (ty + 1)*sz + tx] + prp_0[tz*sz*sz + (ty - 1)*sz + tx] + prp_0[(tz + 1)*sz*sz + ty*sz + tx + 1] + prp_0[(tz - 1)*sz*sz + ty*sz + tx - 1] - - 6.0*prp_0[tz*sz*sz + ty*sz + tx]; + 6.0f*prp_0[tz*sz*sz + ty*sz + tx]; } __global__ void fill_one(float * prp_0,int sz) @@ -145,4 +145,3 @@ void gpu_grid_gradient_vector(grid_gpu<3,Point_aggr_test> & g1, grid_gpu<3,Point CUDA_LAUNCH_DIM3(grid_gradient_vector, gpu_it.wthr, gpu_it.thr ,g1.toKernel(),g2.toKernel(),gpu_it); } - diff --git a/src/Grid/grid_base_implementation.hpp b/src/Grid/grid_base_implementation.hpp index c111c600..33c9df23 100644 --- a/src/Grid/grid_base_implementation.hpp +++ b/src/Grid/grid_base_implementation.hpp @@ -36,7 +36,7 @@ struct skip_init } }; -#ifdef __NVCC__ +#if defined(__NVCC__) || defined(CUDIFY_USE_METAL) template struct copy_ndim_grid_device_active_impl @@ -554,7 +554,7 @@ class grid_base_impl void resize_impl_device(const size_t (& sz)[dim],grid_base_impl & grid_new, unsigned int blockSize = 1) { -#if defined(CUDA_GPU) && defined(__NVCC__) +#if defined(CUDA_GPU) && (defined(__NVCC__) || defined(CUDIFY_USE_METAL)) // Compile time-cheking that make sense to call a GPU kernel to copy. diff --git a/src/NN/CellList/cuda/CellList_gpu.hpp b/src/NN/CellList/cuda/CellList_gpu.hpp index 33eafcb4..4c4f4fab 100644 --- a/src/NN/CellList/cuda/CellList_gpu.hpp +++ b/src/NN/CellList/cuda/CellList_gpu.hpp @@ -16,6 +16,11 @@ #include "NN/CellList/CellList.hpp" #include "util/cuda/scan_ofp.cuh" +#if __cplusplus >= 201703L +#define OPENFPM_CELL_LIST_CONSTEXPR_IF if constexpr +#else +#define OPENFPM_CELL_LIST_CONSTEXPR_IF if +#endif template : public CellDecomposer_sm ); } - if (opt & CL_GPU_REORDER_PROPERTY && sizeof...(prp)) { + OPENFPM_CELL_LIST_CONSTEXPR_IF (sizeof...(prp) != 0) + if (opt & CL_GPU_REORDER_PROPERTY) { CUDA_LAUNCH( (reorderParticlesPrp< decltype(vPrp.toKernel()), @@ -745,7 +751,8 @@ class CellList_gpu : public CellDecomposer_sm ); } - if (opt & CL_GPU_RESTORE_PROPERTY && sizeof...(prp)) { + OPENFPM_CELL_LIST_CONSTEXPR_IF (sizeof...(prp) != 0) + if (opt & CL_GPU_RESTORE_PROPERTY) { CUDA_LAUNCH( (reorderParticlesPrp< decltype(vPrpReordered.toKernel()), @@ -1215,7 +1222,8 @@ class CellList_gpu : public CellDecomposer_sm< ); } - if (opt & CL_GPU_REORDER_PROPERTY && sizeof...(prp)) { + OPENFPM_CELL_LIST_CONSTEXPR_IF (sizeof...(prp) != 0) + if (opt & CL_GPU_REORDER_PROPERTY) { CUDA_LAUNCH( (reorderParticlesPrp< decltype(vPrp.toKernel()), @@ -1480,7 +1488,8 @@ class CellList_gpu : public CellDecomposer_sm< ); } - if (opt & CL_GPU_RESTORE_PROPERTY && sizeof...(prp)) { + OPENFPM_CELL_LIST_CONSTEXPR_IF (sizeof...(prp) != 0) + if (opt & CL_GPU_RESTORE_PROPERTY) { CUDA_LAUNCH( (reorderParticlesPrp< decltype(vPrpReordered.toKernel()), @@ -1529,6 +1538,8 @@ struct toKernel_transform T::is_sparse_::value> type; }; +#undef OPENFPM_CELL_LIST_CONSTEXPR_IF + #endif #endif /* OPENFPM_DATA_SRC_NN_CELLLIST_CELLLIST_GPU_HPP_ */ diff --git a/src/NN/CellList/cuda/CellList_gpu_ker.cuh b/src/NN/CellList/cuda/CellList_gpu_ker.cuh index 26ef674b..a3910292 100644 --- a/src/NN/CellList/cuda/CellList_gpu_ker.cuh +++ b/src/NN/CellList/cuda/CellList_gpu_ker.cuh @@ -20,17 +20,18 @@ class NN_gpu_it inline __device__ void SelectValid() { - while (neighborPartIndexStart == neighborPartIndexStop && isNext()) + while (neighborPartIndexStart == neighborPartIndexStop) { - this->nextCell(); + ++boxNeighborCellOffset_i; + if (boxNeighborCellOffset_i >= neighborCellOffset.size()) return; - if (isNext() == false) break; - - if (cellPositionIndex+this->neighborCellIndexAct+1 >= numPartInCellPrefixSum.size() || cellPositionIndex+this->neighborCellIndexAct < 0) - continue; - - neighborPartIndexStart = numPartInCellPrefixSum.template get<0>(cellPositionIndex+this->neighborCellIndexAct); - neighborPartIndexStop = numPartInCellPrefixSum.template get<0>(cellPositionIndex+this->neighborCellIndexAct+1); + neighborCellIndexAct = neighborCellOffset.template get<0>(boxNeighborCellOffset_i); + int neighborCell = cellPositionIndex + neighborCellIndexAct; + if (neighborCell >= 0 && neighborCell + 1 < numPartInCellPrefixSum.size()) + { + neighborPartIndexStart = numPartInCellPrefixSum.template get<0>(neighborCell); + neighborPartIndexStop = numPartInCellPrefixSum.template get<0>(neighborCell+1); + } } } diff --git a/src/NN/CellList/tests/CellDecomposer_gpu_ker_unit_test.cu b/src/NN/CellList/tests/CellDecomposer_gpu_ker_unit_test.cu index f49dc295..c8039722 100644 --- a/src/NN/CellList/tests/CellDecomposer_gpu_ker_unit_test.cu +++ b/src/NN/CellList/tests/CellDecomposer_gpu_ker_unit_test.cu @@ -19,7 +19,7 @@ BOOST_AUTO_TEST_SUITE( CellDecomposer_gpu_test_suite ) BOOST_AUTO_TEST_CASE( CellDecomposer_gpu_test_use ) { //! Spacing - openfpm::array spacing_c = {0.1,0.1,0.1}; + openfpm::array spacing_c = {0.1f,0.1f,0.1f}; //! \brief number of sub-divisions in each direction openfpm::array div_c = {10,10,10}; @@ -27,7 +27,7 @@ BOOST_AUTO_TEST_CASE( CellDecomposer_gpu_test_use ) //! \brief cell offset openfpm::array off = {2,2,2}; - Point<3,float> trans({0.0,0.0,0.0}); + Point<3,float> trans({0.0f,0.0f,0.0f}); shift_only<3,float> t(Matrix<3,float>::identity(),trans); @@ -35,14 +35,14 @@ BOOST_AUTO_TEST_CASE( CellDecomposer_gpu_test_use ) openfpm::vector_gpu>> output(8); - CUDA_LAUNCH_DIM3(check,1,1,output.toKernel(),clk,(unsigned int)0,Point<3,float>({0.2,0.2,0.2})); - CUDA_LAUNCH_DIM3(check,1,1,output.toKernel(),clk,(unsigned int)1,Point<3,float>({0.1,0.2,0.3})); - CUDA_LAUNCH_DIM3(check,1,1,output.toKernel(),clk,(unsigned int)2,Point<3,float>({0.25,0.55,0.45})); - CUDA_LAUNCH_DIM3(check,1,1,output.toKernel(),clk,(unsigned int)3,Point<3,float>({0.15,0.15,0.95})); - CUDA_LAUNCH_DIM3(check,1,1,output.toKernel(),clk,(unsigned int)4,Point<3,float>({1.05,1.05,1.05})); - CUDA_LAUNCH_DIM3(check,1,1,output.toKernel(),clk,(unsigned int)5,Point<3,float>({1.15,1.15,1.15})); - CUDA_LAUNCH_DIM3(check,1,1,output.toKernel(),clk,(unsigned int)6,Point<3,float>({-0.05,-0.05,-0.05})); - CUDA_LAUNCH_DIM3(check,1,1,output.toKernel(),clk,(unsigned int)7,Point<3,float>({-0.15,-0.15,-0.15})); + CUDA_LAUNCH_DIM3(check,1,1,output.toKernel(),clk,(unsigned int)0,Point<3,float>({0.2f,0.2f,0.2f})); + CUDA_LAUNCH_DIM3(check,1,1,output.toKernel(),clk,(unsigned int)1,Point<3,float>({0.1f,0.2f,0.3f})); + CUDA_LAUNCH_DIM3(check,1,1,output.toKernel(),clk,(unsigned int)2,Point<3,float>({0.25f,0.55f,0.45f})); + CUDA_LAUNCH_DIM3(check,1,1,output.toKernel(),clk,(unsigned int)3,Point<3,float>({0.15f,0.15f,0.95f})); + CUDA_LAUNCH_DIM3(check,1,1,output.toKernel(),clk,(unsigned int)4,Point<3,float>({1.05f,1.05f,1.05f})); + CUDA_LAUNCH_DIM3(check,1,1,output.toKernel(),clk,(unsigned int)5,Point<3,float>({1.15f,1.15f,1.15f})); + CUDA_LAUNCH_DIM3(check,1,1,output.toKernel(),clk,(unsigned int)6,Point<3,float>({-0.05f,-0.05f,-0.05f})); + CUDA_LAUNCH_DIM3(check,1,1,output.toKernel(),clk,(unsigned int)7,Point<3,float>({-0.15f,-0.15f,-0.15f})); output.template deviceToHost<0>(); diff --git a/src/NN/CellList/tests/CellList_gpu_test.cu b/src/NN/CellList/tests/CellList_gpu_test.cu index 81351ee3..b9511bdf 100644 --- a/src/NN/CellList/tests/CellList_gpu_test.cu +++ b/src/NN/CellList/tests/CellList_gpu_test.cu @@ -612,20 +612,20 @@ template void Test_cell_gpu(Box,CudaMemory,memory_traits_inte> vPos; + openfpm::vector,CudaMemory,memory_traits_inte> vPos; openfpm::vector,CudaMemory,memory_traits_inte> vPrp; // create 3 particles - Point p1({0.2,0.2,0.2}); - Point p2({0.9,0.2,0.2}); - Point p3({0.2,0.9,0.2}); - Point p4({0.2,0.2,0.9}); - Point p5({0.9,0.9,0.2}); - Point p6({0.9,0.2,0.9}); - Point p7({0.2,0.9,0.9}); - Point p8({0.9,0.9,0.9}); - Point p9({0.0,0.0,0.0}); + Point p1({0.2,0.2,0.2}); + Point p2({0.9,0.2,0.2}); + Point p3({0.2,0.9,0.2}); + Point p4({0.2,0.2,0.9}); + Point p5({0.9,0.9,0.2}); + Point p6({0.9,0.2,0.9}); + Point p7({0.2,0.9,0.9}); + Point p8({0.9,0.9,0.9}); + Point p9({0.0,0.0,0.0}); vPos.add(p1); vPos.add(p2); @@ -665,7 +665,7 @@ template void Test_cell_gpu(Box(); vPrp.template hostToDevice<0,1,2>(); - openfpm::vector,CudaMemory,memory_traits_inte> vPosReorder(vPos.size()); + openfpm::vector,CudaMemory,memory_traits_inte> vPosReorder(vPos.size()); openfpm::vector,CudaMemory,memory_traits_inte> vPrpReorder(vPrp.size()); // create an gpu context @@ -686,7 +686,7 @@ template void Test_cell_gpu(Box,CudaMemory,memory_traits_inte> pl_correct; + openfpm::vector,CudaMemory,memory_traits_inte> pl_correct; pl_correct.add(p9); pl_correct.add(p1); @@ -756,10 +756,15 @@ BOOST_AUTO_TEST_CASE( CellList_gpu_use) { std::cout << "Test cell list GPU" << "\n"; - Box<3,double> box({0.0f,0.0f,0.0f},{1.0f,1.0f,1.0f}); - Box<3,double> box2({-1.0f,-1.0f,-1.0f},{1.0f,1.0f,1.0f}); +#ifdef CUDIFY_USE_METAL + using cell_scalar = float; +#else + using cell_scalar = double; +#endif + Box<3,cell_scalar> box({0.0f,0.0f,0.0f},{1.0f,1.0f,1.0f}); + Box<3,cell_scalar> box2({-1.0f,-1.0f,-1.0f},{1.0f,1.0f,1.0f}); - Test_cell_gpu<3,double,CellList_gpu<3,double,CudaMemory>>(box); + Test_cell_gpu<3,cell_scalar,CellList_gpu<3,cell_scalar,CudaMemory>>(box); std::cout << "End cell list GPU" << "\n"; @@ -770,10 +775,16 @@ BOOST_AUTO_TEST_CASE( CellList_gpu_use_sparse ) { std::cout << "Test cell list GPU sparse" << "\n"; - Box<3,double> box({0.0f,0.0f,0.0f},{1.0f,1.0f,1.0f}); - Box<3,double> box2({-1.0f,-1.0f,-1.0f},{1.0f,1.0f,1.0f}); +#ifdef CUDIFY_USE_METAL + using cell_scalar = float; +#else + using cell_scalar = double; +#endif + Box<3,cell_scalar> box({0.0f,0.0f,0.0f},{1.0f,1.0f,1.0f}); + Box<3,cell_scalar> box2({-1.0f,-1.0f,-1.0f},{1.0f,1.0f,1.0f}); - Test_cell_gpu<3,double,CellList_gpu<3,double,CudaMemory,no_transform_only<3,double>,true>> (box); + Test_cell_gpu<3,cell_scalar,CellList_gpu<3,cell_scalar,CudaMemory, + no_transform_only<3,cell_scalar>,true>> (box); std::cout << "End cell list GPU sparse" << "\n"; @@ -1996,4 +2007,3 @@ BOOST_AUTO_TEST_CASE( CellList_swap_test ) } BOOST_AUTO_TEST_SUITE_END() - diff --git a/src/Space/Shape/Point.hpp b/src/Space/Shape/Point.hpp index 84eee06a..044e28fb 100644 --- a/src/Space/Shape/Point.hpp +++ b/src/Space/Shape/Point.hpp @@ -159,6 +159,19 @@ template class Point {get(i) = x;i++;} } + /*! \brief Constructor from a homogeneous list with a convertible scalar type + * + * \param p1 initializer list + * + */ + template::value>::type> + __device__ __host__ inline Point(std::initializer_list p1) + { + size_t i = 0; + for(S x : p1) + {get(i) = static_cast(x);i++;} + } + //! Default contructor __device__ __host__ inline Point() {} diff --git a/src/Space/Shape/Sphere.hpp b/src/Space/Shape/Sphere.hpp index f2e1ef85..cb03481d 100644 --- a/src/Space/Shape/Sphere.hpp +++ b/src/Space/Shape/Sphere.hpp @@ -96,7 +96,8 @@ template class Sphere * \param radius * */ - Sphere(Point & c, T radius) + template + __device__ __host__ Sphere(const Point & c, T radius) { // Copy the center for (int i = 0 ; i < dim ; i++) diff --git a/src/Space/Shape/Sphere_unit_test.cpp b/src/Space/Shape/Sphere_unit_test.cpp index 61a0458a..5b44c784 100644 --- a/src/Space/Shape/Sphere_unit_test.cpp +++ b/src/Space/Shape/Sphere_unit_test.cpp @@ -27,7 +27,7 @@ BOOST_AUTO_TEST_CASE( Sphere_test_use) BOOST_REQUIRE_EQUAL(s.isInside(p3),false); double dist = s.distance(p3); - BOOST_REQUIRE_EQUAL(dist,0.15980762113533162); + BOOST_REQUIRE_CLOSE_FRACTION(dist,0.15980762113533162,1.0e-15); } BOOST_AUTO_TEST_SUITE_END() diff --git a/src/SparseGrid/SparseGrid_conv_opt.hpp b/src/SparseGrid/SparseGrid_conv_opt.hpp index c8e70295..1376aea6 100644 --- a/src/SparseGrid/SparseGrid_conv_opt.hpp +++ b/src/SparseGrid/SparseGrid_conv_opt.hpp @@ -453,7 +453,7 @@ struct conv_impl<3> long int sumzp = (v == sz::value-1)?offset_jump[5] - (sz::value - 1)*sx::value*sy::value:sx::value*sy::value; sumzp += s2; - if (Vc::Vector::Size == 2 || Vc::Vector::Size == 4 || Vc::Vector::Size == 8) + if (Vc::Vector::Size == 1 || Vc::Vector::Size == 2 || Vc::Vector::Size == 4 || Vc::Vector::Size == 8) { mxm.i = *(typename data_il::Size>::type *)&mask.mask[s2]; mxm.i = mxm.i << 8; @@ -898,7 +898,7 @@ struct conv_impl<3> for (int k = 0 ; k < sx::value ; k += Vc::Vector::Size) { // we do only id exist the point - if (*(int *)&mask.mask[s2] == 0) {s2 += Vc::Vector::Size; continue;} + if (*(typename data_il::Size>::type *)&mask.mask[s2] == 0) {s2 += Vc::Vector::Size; continue;} data_il<4> mxm; data_il<4> mxp; @@ -928,7 +928,17 @@ struct conv_impl<3> ids.s2 = s2; - if (Vc::Vector::Size == 2) + if (Vc::Vector::Size == 1) + { + mxm.i = mxp.i = mym.i = myp.i = mzm.i = mzp.i = 0; + mxm.uc[0] = mask.mask[ids.sumdm[0]]; + mxp.uc[0] = mask.mask[ids.sumdp[0]]; + mym.uc[0] = mask.mask[ids.sumdm[1]]; + myp.uc[0] = mask.mask[ids.sumdp[1]]; + mzm.uc[0] = mask.mask[ids.sumdm[2]]; + mzp.uc[0] = mask.mask[ids.sumdp[2]]; + } + else if (Vc::Vector::Size == 2) { mxm.i = *(short int *)&mask.mask[s2]; mxm.i = mxm.i << 8; diff --git a/src/SparseGrid/SparseGrid_unit_tests.cpp b/src/SparseGrid/SparseGrid_unit_tests.cpp index 78a0ee16..e585e3b8 100644 --- a/src/SparseGrid/SparseGrid_unit_tests.cpp +++ b/src/SparseGrid/SparseGrid_unit_tests.cpp @@ -920,24 +920,8 @@ BOOST_AUTO_TEST_CASE( sparse_grid_fast_stencil_vectorized) for (int i = it.start_b(0) ; i < it.stop_b(0) ; i += Vc::double_v::Size) { Vc::Mask cmp; - - if (Vc::double_v::Size == 2) - { - cmp[0] = mask[c] == true; - cmp[1] = mask[c+1] == true; - } - else if (Vc::double_v::Size == 4) - { - cmp[0] = mask[c] == true; - cmp[1] = mask[c+1] == true; - cmp[2] = mask[c+2] == true; - cmp[3] = mask[c+3] == true; - } - else - { - std::cout << "UNSUPPORTED" << std::endl; - exit(1); - } + for (int lane = 0; lane < Vc::double_v::Size; ++lane) + cmp[lane] = mask[c+lane] == true; // we do only id exist the point @@ -957,18 +941,8 @@ BOOST_AUTO_TEST_CASE( sparse_grid_fast_stencil_vectorized) ypd + ymd + zpd + zmd - 6.0*cmd; - if (Vc::double_v::Size == 2) - { - surround[0] = (mask_sum[c] == 7); - surround[1] = (mask_sum[c+1] == 7); - } - else if (Vc::double_v::Size == 4) - { - surround[0] = (mask_sum[c] == 7); - surround[1] = (mask_sum[c+1] == 7); - surround[2] = (mask_sum[c+2] == 7); - surround[3] = (mask_sum[c+3] == 7); - } + for (int lane = 0; lane < Vc::double_v::Size; ++lane) + surround[lane] = mask_sum[c+lane] == 7; Lap = Vc::iif(surround,Lap,six); @@ -1126,7 +1100,12 @@ BOOST_AUTO_TEST_CASE( sparse_grid_fast_stencil_vectorized_block_skip) long int sumzp = (v == sz::value-1)?offset_jump[5] - (sz::value - 1)*sx::value*sy::value:sx::value*sy::value; sumzp += s2; - if (Vc::double_v::Size == 2) + if (Vc::double_v::Size == 1) + { + mxm.uc[0] = mask.mask[sumxm]; + mxp.uc[0] = mask.mask[sumxp]; + } + else if (Vc::double_v::Size == 2) { mxm.uc[0] = mask.mask[sumxm]; mxm.uc[1] = mask.mask[s2]; @@ -1146,11 +1125,10 @@ BOOST_AUTO_TEST_CASE( sparse_grid_fast_stencil_vectorized_block_skip) mxp.uc[2] = mask.mask[s2+3]; mxp.uc[3] = mask.mask[sumxp]; } - else - { - std::cout << "UNSUPPORTED" << std::endl; - exit(1); - } + else + { + BOOST_FAIL("Unsupported Vc vector width"); + } mym.i = *(int *)&mask.mask[sumym]; myp.i = *(int *)&mask.mask[sumyp]; @@ -1158,7 +1136,12 @@ BOOST_AUTO_TEST_CASE( sparse_grid_fast_stencil_vectorized_block_skip) mzm.i = *(int *)&mask.mask[sumzm]; mzp.i = *(int *)&mask.mask[sumzp]; - if (Vc::double_v::Size == 2) + if (Vc::double_v::Size == 1) + { + xm[0] = chunk.template get<0>()[sumxm]; + xp[0] = chunk.template get<0>()[sumxp]; + } + else if (Vc::double_v::Size == 2) { xm[0] = chunk.template get<0>()[sumxm]; xm[1] = cmd[0]; @@ -1191,18 +1174,8 @@ BOOST_AUTO_TEST_CASE( sparse_grid_fast_stencil_vectorized_block_skip) data_il tot_m; tot_m.i = mxm.i + mxp.i + mym.i + myp.i + mzm.i + mzp.i; - if (Vc::double_v::Size == 2) - { - surround[0] = (tot_m.uc[0] == 6); - surround[1] = (tot_m.uc[1] == 6); - } - else if (Vc::double_v::Size == 4) - { - surround[0] = (tot_m.uc[0] == 6); - surround[1] = (tot_m.uc[1] == 6); - surround[2] = (tot_m.uc[2] == 6); - surround[3] = (tot_m.uc[3] == 6); - } + for (int lane = 0; lane < Vc::double_v::Size; ++lane) + surround[lane] = tot_m.uc[lane] == 6; Vc::double_v Lap = xp + xm + yp + ym + @@ -2435,4 +2408,3 @@ BOOST_AUTO_TEST_CASE( sparse_testing_clear ) } BOOST_AUTO_TEST_SUITE_END() - diff --git a/src/SparseGridGpu/SparseGridGpu.hpp b/src/SparseGridGpu/SparseGridGpu.hpp index 1f0f9500..95cd1cc0 100644 --- a/src/SparseGridGpu/SparseGridGpu.hpp +++ b/src/SparseGridGpu/SparseGridGpu.hpp @@ -319,7 +319,7 @@ struct NNFull template __device__ static inline bool isPadding(sparseGrid_type & sparseGrid, coord_type & coord, Mask_type (& enlargedBlock)[eb_size]) { - return NNfull_is_padding_impl<3>::template is_padding(sparseGrid,coord,enlargedBlock); + return NNfull_is_padding_impl<3>::is_padding(sparseGrid,coord,enlargedBlock); } /*! \brief given a coordinate writtel in local coordinate for a given it return the neighborhood chunk position and the offset diff --git a/src/SparseGridGpu/tests/SparseGridGpu_tests.cu b/src/SparseGridGpu/tests/SparseGridGpu_tests.cu index b7569737..5190c326 100644 --- a/src/SparseGridGpu/tests/SparseGridGpu_tests.cu +++ b/src/SparseGridGpu/tests/SparseGridGpu_tests.cu @@ -135,7 +135,8 @@ __global__ void insertBoundaryValuesHeat(SparseGridType sparseGrid) if (y == 0 || y == bDimY * gridDim.y - 1) { - value = 10.0 * x / (bDimX * gridDim.x - 1); + value = 10.0f * static_cast(x) / + static_cast(bDimX * gridDim.x - 1); } sparseGrid.template insert

(coord) = value; @@ -180,7 +181,8 @@ struct LaplacianStencil const auto coord = sparseGrid.getCoordInEnlargedBlock(offset); const auto linId = sparseGrid.getLinIdInEnlargedBlock(offset); ScalarT cur = enlargedBlock[linId]; - ScalarT res = -2.0*dim*cur; // The central part of the stencil + ScalarT res = static_cast(-2) * + static_cast(dim) * cur; // The central part of the stencil for (int d=0; d>(sparseGrid.getBox(),STENCIL_MODE_INPLACE,0.0 ,gridSize.x * blockEdgeSize, 0.0, 10.0); + sparseGrid.template applyStencils>(sparseGrid.getBox(),STENCIL_MODE_INPLACE,0.0f ,gridSize.x * blockEdgeSize, 0.0f, 10.0f); // Now apply the laplacian operator const unsigned int maxIter = 1000; // const unsigned int maxIter = 100; for (unsigned int iter=0; iter>(sparseGrid.getBox(),STENCIL_MODE_INPLACE, 0.1); - sparseGrid.applyStencils>(sparseGrid.getBox(),STENCIL_MODE_INPLACE, 0.01); + sparseGrid.applyStencils>(sparseGrid.getBox(),STENCIL_MODE_INPLACE, 0.1f); + sparseGrid.applyStencils>(sparseGrid.getBox(),STENCIL_MODE_INPLACE, 0.01f); } sparseGrid.template deviceToHost<0,1>(); @@ -540,7 +542,7 @@ BOOST_AUTO_TEST_CASE(testStencil_lap_simplified) sparseGrid.findNeighbours(); // Pre-compute the neighbours pos for each block! sparseGrid.tagBoundaries(gpuContext); - sparseGrid.template applyStencils>(sparseGrid.getBox(),STENCIL_MODE_INPLACE,0.0 ,gridSize.x * blockEdgeSize, 0.0, 10.0); + sparseGrid.template applyStencils>(sparseGrid.getBox(),STENCIL_MODE_INPLACE,0.0f ,gridSize.x * blockEdgeSize, 0.0f, 10.0f); // Now apply the laplacian operator const unsigned int maxIter = 1000; @@ -549,11 +551,11 @@ BOOST_AUTO_TEST_CASE(testStencil_lap_simplified) { sparseGrid.conv_cross<0, 1, 1>({0,0},{16,16},[] __device__ (float & u, cross_stencil<2,float> & cs){ return u + (cs.xm[0] + cs.xp[0] + - cs.xm[1] + cs.xp[1] - 4.0*u)*0.1; + cs.xm[1] + cs.xp[1] - 4.0f*u)*0.1f; }); sparseGrid.conv_cross<1, 0, 1>({0,0},{16,16},[] __device__ (float & u, cross_stencil<2,float> & cs){ return u + (cs.xm[0] + cs.xp[0] + - cs.xm[1] + cs.xp[1] - 4.0*u)*0.1; + cs.xm[1] + cs.xp[1] - 4.0f*u)*0.1f; }); } @@ -567,7 +569,7 @@ BOOST_AUTO_TEST_CASE(testStencil_lap_simplified) for (size_t i = 0; i < 64*4; i++) { auto coord = sparseGrid.getCoord(i); - float expectedValue = 10.0 * coord.get(0) / (gridSize.x * blockEdgeSize - 1); + float expectedValue = 10.0f * coord.get(0) / (gridSize.x * blockEdgeSize - 1); match &= fabs(sparseGrid.template get<0>(coord) - expectedValue) < 1e-2; } @@ -598,8 +600,8 @@ BOOST_AUTO_TEST_CASE(testStencil_lap_no_cross_simplified) sparseGrid.findNeighbours(); // Pre-compute the neighbours pos for each block! sparseGrid.tagBoundaries(gpuContext); - sparseGrid.template applyStencils>(sparseGrid.getBox(),STENCIL_MODE_INPLACE,0.0 ,gridSize.x * blockEdgeSize, 0.0, 10.0); - sparseGrid.template applyStencils>(sparseGrid.getBox(),STENCIL_MODE_INPLACE,0.0 ,gridSize.x * blockEdgeSize, 0.0, 10.0); + sparseGrid.template applyStencils>(sparseGrid.getBox(),STENCIL_MODE_INPLACE,0.0f ,gridSize.x * blockEdgeSize, 0.0f, 10.0f); + sparseGrid.template applyStencils>(sparseGrid.getBox(),STENCIL_MODE_INPLACE,0.0f ,gridSize.x * blockEdgeSize, 0.0f, 10.0f); typedef typename GetCpBlockType::type CpBlockType; @@ -625,13 +627,13 @@ BOOST_AUTO_TEST_CASE(testStencil_lap_no_cross_simplified) sparseGrid.conv<0, 1, 1>({0,0},{16,16},[] __device__ (CpBlockType & u,int i, int j){ float c = u(i,j); return c + (u(i-1,j) + u(i+1,j) + - u(i,j-1) + u(i,j+1) - 4.0*c)*0.1; + u(i,j-1) + u(i,j+1) - 4.0f*c)*0.1f; }); sparseGrid.conv<1, 0, 1>({0,0},{16,16},[] __device__ (CpBlockType & u,int i, int j){ float c = u(i,j); return c + (u(i-1,j) + u(i+1,j) + - u(i,j-1) + u(i,j+1) - 4.0*c)*0.1; + u(i,j-1) + u(i,j+1) - 4.0f*c)*0.1f; }); } @@ -676,10 +678,10 @@ BOOST_AUTO_TEST_CASE(testStencil_lap_no_cross_simplified2) sparseGrid.findNeighbours(); // Pre-compute the neighbours pos for each block! sparseGrid.tagBoundaries(gpuContext); - sparseGrid.template applyStencils>(sparseGrid.getBox(),STENCIL_MODE_INPLACE,0.0 ,gridSize.x * blockEdgeSize, 0.0, 10.0); - sparseGrid.template applyStencils>(sparseGrid.getBox(),STENCIL_MODE_INPLACE,0.0 ,gridSize.x * blockEdgeSize, 0.0, 5.0); - sparseGrid.template applyStencils>(sparseGrid.getBox(),STENCIL_MODE_INPLACE,0.0 ,gridSize.x * blockEdgeSize, 0.0, 10.0); - sparseGrid.template applyStencils>(sparseGrid.getBox(),STENCIL_MODE_INPLACE,0.0 ,gridSize.x * blockEdgeSize, 0.0, 5.0); + sparseGrid.template applyStencils>(sparseGrid.getBox(),STENCIL_MODE_INPLACE,0.0f ,gridSize.x * blockEdgeSize, 0.0f, 10.0f); + sparseGrid.template applyStencils>(sparseGrid.getBox(),STENCIL_MODE_INPLACE,0.0f ,gridSize.x * blockEdgeSize, 0.0f, 5.0f); + sparseGrid.template applyStencils>(sparseGrid.getBox(),STENCIL_MODE_INPLACE,0.0f ,gridSize.x * blockEdgeSize, 0.0f, 10.0f); + sparseGrid.template applyStencils>(sparseGrid.getBox(),STENCIL_MODE_INPLACE,0.0f ,gridSize.x * blockEdgeSize, 0.0f, 5.0f); typedef typename GetCpBlockType::type CpBlockType; @@ -706,20 +708,20 @@ BOOST_AUTO_TEST_CASE(testStencil_lap_no_cross_simplified2) float cu = u(i,j); float cv = v(i,j); u_out = cu + (u(i-1,j) + u(i+1,j) + - u(i,j-1) + u(i,j+1) - 4.0*cu)*0.1; + u(i,j-1) + u(i,j+1) - 4.0f*cu)*0.1f; v_out = cv + (v(i-1,j) + v(i+1,j) + - v(i,j-1) + v(i,j+1) - 4.0*cv)*0.1; + v(i,j-1) + v(i,j+1) - 4.0f*cv)*0.1f; }); sparseGrid.conv2<2,3,0,1,1>({0,0},{16,16},[] __device__ (float & u_out, float & v_out ,CpBlockType & u, CpBlockType & v, int i, int j){ float cu = u(i,j); float cv = v(i,j); u_out = cu + (u(i-1,j) + u(i+1,j) + - u(i,j-1) + u(i,j+1) - 4.0*cu)*0.1; + u(i,j-1) + u(i,j+1) - 4.0f*cu)*0.1f; v_out = cv + (v(i-1,j) + v(i+1,j) + - v(i,j-1) + v(i,j+1) - 4.0*cv)*0.1; + v(i,j-1) + v(i,j+1) - 4.0f*cv)*0.1f; }); } @@ -766,7 +768,7 @@ BOOST_AUTO_TEST_CASE(testStencil_lap_no_cross_simplified_subset) typedef typename GetCpBlockType::type CpBlockType; sparseGrid.conv<0, 1, 1>({3,3},{11,11},[] __device__ (CpBlockType & u,int i, int j){ - return 5.0; + return 5.0f; }); @@ -921,7 +923,7 @@ struct Conv3x3x3 { const auto coord = sparseGrid.getCoordInEnlargedBlock(offset); const auto linId = sparseGrid.getLinIdInEnlargedBlock(offset); - ScalarT tot = 0.0; + ScalarT tot = static_cast(0); for (int i = 0; i < dim; ++i) { for (int j = 0; j < dim; ++j) @@ -977,7 +979,7 @@ struct Conv3x3x3_noshared if ((curMask & mask_sparse::EXIST) && !(curMask & mask_sparse::PADDING)) { - ScalarT tot = 0.0; + ScalarT tot = static_cast(0); for (int i = 0; i < dim; ++i) { for (int j = 0; j < dim; ++j) @@ -1989,7 +1991,7 @@ BOOST_AUTO_TEST_CASE(testSparseGridGpuOutput3DHeatStencil) // Now apply some boundary conditions sparseGrid.template applyStencils>(sparseGrid.getBox(),STENCIL_MODE_INPLACE, 192, 384, - 0.0, 10.0); + 0.0f, 10.0f); // Now apply the laplacian operator @@ -1999,9 +2001,9 @@ BOOST_AUTO_TEST_CASE(testSparseGridGpuOutput3DHeatStencil) { for (int innerIter=0; innerIter<10; ++innerIter) { - sparseGrid.applyStencils>(sparseGrid.getBox(),STENCIL_MODE_INPLACE, 0.1); + sparseGrid.applyStencils>(sparseGrid.getBox(),STENCIL_MODE_INPLACE, 0.1f); - sparseGrid.applyStencils>(sparseGrid.getBox(),STENCIL_MODE_INPLACE, 0.1); + sparseGrid.applyStencils>(sparseGrid.getBox(),STENCIL_MODE_INPLACE, 0.1f); } } diff --git a/src/SparseGridGpu/tests/utils/SparseGridGpu_util_test.cuh b/src/SparseGridGpu/tests/utils/SparseGridGpu_util_test.cuh index 610414a6..cf73a45e 100644 --- a/src/SparseGridGpu/tests/utils/SparseGridGpu_util_test.cuh +++ b/src/SparseGridGpu/tests/utils/SparseGridGpu_util_test.cuh @@ -254,7 +254,8 @@ struct HeatStencil // const auto linId = sparseGrid.getLinIdInEnlargedBlock(offset); linId = sparseGrid.getLinIdInEnlargedBlock(offset); ScalarT cur = enlargedBlock[linId]; - ScalarT laplacian = -2.0 * dim * cur; // The central part of the stencil + ScalarT laplacian = static_cast(-2) * + static_cast(dim) * cur; // The central part of the stencil for (int d = 0; d < dim; ++d) { @@ -388,7 +389,7 @@ struct Conv3x3x3 __shared__ ScalarT enlargedBlock[enlargedBlockSize]; - sparseGrid.loadGhostBlock(dataBlockLoad,dataBlockIdPos,enlargedBlock); + sparseGrid.template loadGhostBlock(dataBlockLoad,dataBlockIdPos,enlargedBlock); __syncthreads(); @@ -860,7 +861,7 @@ struct SkeletonStencil SparseGridT::getBlockEdgeSize() + 2 * supportRadius, dim>::value; __shared__ ScalarT enlargedBlock[enlargedBlockSize]; - sparseGrid.loadGhostBlock(dataBlockLoad, dataBlockIdPos, enlargedBlock); + sparseGrid.template loadGhostBlock(dataBlockLoad, dataBlockIdPos, enlargedBlock); __syncthreads(); @@ -880,7 +881,7 @@ struct SkeletonStencil enlargedBlock[linId] = res; } __syncthreads(); - sparseGrid.storeBlock(dataBlockStore, enlargedBlock); + sparseGrid.template storeBlock(dataBlockStore, enlargedBlock); } /*! \brief Stencil Host function diff --git a/src/Vector/cuda/map_vector_cuda_ker.cuh b/src/Vector/cuda/map_vector_cuda_ker.cuh index fef74515..9672624b 100644 --- a/src/Vector/cuda/map_vector_cuda_ker.cuh +++ b/src/Vector/cuda/map_vector_cuda_ker.cuh @@ -389,7 +389,7 @@ namespace openfpm vector_gpu_ker(const vector_gpu_ker_ref & vref) { - this->operator=(vref.vref); + this->operator=(*vref.vref); } /*! \brief implementation of the constructor @@ -640,7 +640,7 @@ namespace openfpm typedef typename apply_transform::type T_; //! vector reference - vector_gpu_ker & vref; + vector_gpu_ker * vref = nullptr; public: @@ -656,7 +656,7 @@ namespace openfpm __device__ __host__ unsigned int size() const { - return vref.size(); + return vref->size(); } __host__ __device__ size_t size_local() const @@ -666,89 +666,97 @@ namespace openfpm __device__ __host__ unsigned int capacity() const { - return vref.capacity; + return vref->capacity(); } template - __device__ __host__ inline auto get(unsigned int id) const -> decltype(vref.template get

(id)) + __device__ __host__ inline auto get(unsigned int id) const -> decltype(vref->template get

(id)) { - return vref.template get

(id); + return vref->template get

(id); } - __device__ __host__ inline auto get(unsigned int id) -> decltype(vref.get(id)) + __device__ __host__ inline auto get(unsigned int id) -> decltype(vref->get(id)) { - return vref.get(id); + return vref->get(id); } - inline __device__ __host__ auto get(unsigned int id) const -> decltype(vref.get(id)) + inline __device__ __host__ auto get(unsigned int id) const -> decltype(vref->get(id)) { - return vref.get(id); + return vref->get(id); } - inline __device__ __host__ auto get_o(unsigned int id) const -> decltype(vref.get_o(id)) + inline __device__ __host__ auto get_o(unsigned int id) const -> decltype(vref->get_o(id)) { - return vref.get_o(id); + return vref->get_o(id); } - inline __device__ __host__ auto get_o(unsigned int id) -> decltype(vref.get_o(id)) + inline __device__ __host__ auto get_o(unsigned int id) -> decltype(vref->get_o(id)) { - return vref.get_o(id); + return vref->get_o(id); } - inline auto last() const -> decltype(vref.last()) + inline auto last() const -> decltype(vref->last()) { - return vref.last(); + return vref->last(); } template - __device__ __host__ inline auto get(unsigned int id) -> decltype(vref.template get

(id)) + __device__ __host__ inline auto get(unsigned int id) -> decltype(vref->template get

(id)) { - return vref.template get

(id); + return vref->template get

(id); } - inline auto last() -> decltype(vref.last()) + inline auto last() -> decltype(vref->last()) { - return vref.last(); + return vref->last(); } + vector_gpu_ker_ref() = default; + vector_gpu_ker_ref(vector_gpu_ker & vref) - :vref(vref) + :vref(&vref) {} + vector_gpu_ker_ref & operator=(const vector_gpu_ker_ref & other) + { + vref = other.vref; + return *this; + } + __device__ void set(int id, const container & obj) { - vref.set(id,obj); + vref->set(id,obj); } template __device__ __host__ void * getPointer() { - return vref.template getPointer

(); + return vref->template getPointer

(); } template __device__ __host__ const void * getPointer() const { - return vref.template getPointer

(); + return vref->template getPointer

(); } template void set_o(unsigned int i, const encap_S & obj) { - vref.set(i,obj); + vref->set(i,obj); } __device__ void set(unsigned int id, const vector_gpu_ker & v, unsigned int src) { - vref.set(id,v,src); + vref->set(id,v,src); } template __device__ void set(unsigned int id, const vector_gpu_ker & v, unsigned int src) { - vref.template set(id,v,src); + vref->template set(id,v,src); } __host__ ite_gpu<1> getGPUIterator(size_t n_thr = default_kernel_wg_threads_) const { - return vref.getGPUIterator(n_thr); + return vref->getGPUIterator(n_thr); } /*! \brief Get an iterator for the GPU @@ -757,40 +765,63 @@ namespace openfpm */ ite_gpu<1> getGPUIteratorTo(size_t stop, size_t n_thr = default_kernel_wg_threads_) const { - return vref.getGPUItertatorTo(stop,n_thr); + return vref->getGPUItertatorTo(stop,n_thr); + } + + ite_gpu<1> getDomainIteratorGPU(size_t n_thr = default_kernel_wg_threads_) const + { + return vref->getDomainIteratorGPU(n_thr); + } + + template + __device__ __host__ auto getProp(unsigned int id) const + -> decltype(vref->template getProp

(id)) + { + return vref->template getProp

(id); + } + + void init() const + { + vref->init(); + } + + __host__ __device__ auto value(unsigned int p) + -> decltype(vref->value(p)) + { + return vref->value(p); } vector_gpu_ker & getVector() { - return *this; + return *vref; } const vector_gpu_ker & getVector() const { - return *this; + return *vref; } __host__ vector_gpu_ker_ref & operator=(const vector_gpu_ker & v) { - vref.operator=(v); - return this; + if (vref != nullptr) {vref->operator=(v);} + return *this; } __device__ grid_gpu_ker<1,T_,layout_base, grid_sm<1,void>> & getBase() { - return vref.getBase(); + return vref->getBase(); } pointer_check check_device_pointer(void * ptr) { - return vref.check_device_pointer(ptr); + return vref->check_device_pointer(ptr); } - void * internal_get_size_pointer() {return &vref.internal_get_size_pointer();} + void * internal_get_size_pointer() {return vref->internal_get_size_pointer();} void print_size() { - return vref.print_size(); + return vref->print_size(); } }; diff --git a/src/Vector/cuda/map_vector_sparse_cuda_ker_unit_tests.cu b/src/Vector/cuda/map_vector_sparse_cuda_ker_unit_tests.cu index 6fd752c7..be5c8d6c 100644 --- a/src/Vector/cuda/map_vector_sparse_cuda_ker_unit_tests.cu +++ b/src/Vector/cuda/map_vector_sparse_cuda_ker_unit_tests.cu @@ -189,9 +189,16 @@ __global__ void test_sparse_get_test(vd_sparse_type vd_test, vector_out_type out BOOST_AUTO_TEST_SUITE( vector_cuda_sparse ) +#ifdef CUDIFY_USE_METAL +using sparse_test_real = float; +#else +using sparse_test_real = double; +#endif +using sparse_test_properties = aggregate; + BOOST_AUTO_TEST_CASE( vector_sparse_cuda_gpu ) { - openfpm::vector_sparse_gpu> vs; + openfpm::vector_sparse_gpu vs; vs.template setBackground<0>(17); @@ -217,7 +224,7 @@ BOOST_AUTO_TEST_CASE( vector_sparse_cuda_gpu ) vs.flush,smin_<1>,smax_<2> >(gpuContext,flush_type::FLUSH_ON_DEVICE); - openfpm::vector_gpu> output; + openfpm::vector_gpu output; output.resize(1500); @@ -260,7 +267,7 @@ BOOST_AUTO_TEST_CASE( vector_sparse_cuda_gpu ) BOOST_AUTO_TEST_CASE( vector_sparse_cuda_gpu_incremental_add ) { - openfpm::vector_sparse_gpu> vs; + openfpm::vector_sparse_gpu vs; vs.template setBackground<0>(17); @@ -338,7 +345,7 @@ BOOST_AUTO_TEST_CASE( vector_sparse_cuda_gpu_incremental_add ) BOOST_AUTO_TEST_CASE( vector_sparse_cuda_gpu_get ) { - openfpm::vector_sparse_gpu> vs; + openfpm::vector_sparse_gpu vs; vs.template setBackground<0>(0); @@ -459,7 +466,7 @@ BOOST_AUTO_TEST_CASE( vector_sparse_cuda_gpu_get ) BOOST_AUTO_TEST_CASE( vector_sparse_cuda_gpu_special_function ) { - openfpm::vector_sparse_gpu> vs; + openfpm::vector_sparse_gpu vs; vs.template setBackground<0>(17); @@ -509,7 +516,7 @@ BOOST_AUTO_TEST_CASE( vector_sparse_cuda_gpu_special_function ) //////////////////////////////// REMOVE test section -void check_lines(openfpm::vector_sparse_gpu> & vs, +void check_lines(openfpm::vector_sparse_gpu & vs, bool s1, bool s2, bool s3, @@ -591,7 +598,7 @@ void check_lines(openfpm::vector_sparse_gpu> & vs BOOST_AUTO_TEST_CASE( vector_sparse_cuda_gpu_remove ) { - openfpm::vector_sparse_gpu> vs; + openfpm::vector_sparse_gpu vs; vs.template setBackground<0>(17); @@ -670,7 +677,7 @@ BOOST_AUTO_TEST_CASE( vector_sparse_cuda_gpu_remove ) BOOST_AUTO_TEST_CASE( vector_sparse_cuda_gpu_remove_incremental ) { - openfpm::vector_sparse_gpu> vs; + openfpm::vector_sparse_gpu vs; vs.template setBackground<0>(17); diff --git a/src/Vector/cuda/map_vector_sparse_cuda_kernels.cuh b/src/Vector/cuda/map_vector_sparse_cuda_kernels.cuh index d78e2bf3..6bec9fde 100644 --- a/src/Vector/cuda/map_vector_sparse_cuda_kernels.cuh +++ b/src/Vector/cuda/map_vector_sparse_cuda_kernels.cuh @@ -8,12 +8,12 @@ #ifndef MAP_VECTOR_SPARSE_CUDA_KERNELS_CUH_ #define MAP_VECTOR_SPARSE_CUDA_KERNELS_CUH_ +#include + #ifdef __NVCC__ #include "config.h" -#include - #if CUDART_VERSION < 11000 #include "util/cuda/cub_old/util_type.cuh" #include "util/cuda/cub_old/block/block_scan.cuh" @@ -26,17 +26,43 @@ #endif +template +struct sparse_reduction_make_void +{ + typedef void type; +}; + +template +using sparse_reduction_void_t = typename sparse_reduction_make_void::type; + +template +struct sparse_reduction_scalar +{ + typedef T type; +}; + +template +struct sparse_reduction_scalar> +{ + typedef typename T::scalarType type; +}; + template struct zero_t { __device__ __host__ type_t operator()() const { - return 0; + type_t value{}; + value = 0; + return value; } }; template struct limit_max_t { __device__ __host__ type_t operator()() const { - return std::numeric_limits::max(); + type_t value{}; + typedef typename sparse_reduction_scalar::type scalar_type; + value = std::numeric_limits::max(); + return value; } }; @@ -142,6 +168,8 @@ __global__ void set_one_insert_buffer(vect_type vadd) vadd.template get<0>(p) = 1; } +#endif + template struct plus_block_t { __device__ __host__ type_t operator()(type_t a, type_t b) const { @@ -154,8 +182,6 @@ struct plus_block_t { } }; -#endif - template struct sadd_block_ { @@ -214,22 +240,18 @@ struct smax_ {} }; -#ifdef __NVCC__ - template struct maximum_block_t { __forceinline__ __device__ __host__ type_t operator()(type_t a, type_t b) const { type_t res; for (int i=0; i struct smax_block_ { @@ -290,22 +312,18 @@ struct smin_ {} }; -#ifdef __NVCC__ - template struct minimum_block_t { __forceinline__ __device__ __host__ type_t operator()(type_t a, type_t b) const { type_t res; for (int i=0; i struct smin_block_ { @@ -338,9 +356,6 @@ struct smin_block_ {} }; - -#ifdef __NVCC__ - template struct bitwiseOr_t { __forceinline__ __device__ __host__ type_t operator()(type_t a, type_t b) const { @@ -348,6 +363,8 @@ struct bitwiseOr_t { } }; +#ifdef __NVCC__ + template struct sBitwiseOr_ { diff --git a/src/memory_ly/memory_array.hpp b/src/memory_ly/memory_array.hpp index e7469f15..d07f3e68 100644 --- a/src/memory_ly/memory_array.hpp +++ b/src/memory_ly/memory_array.hpp @@ -70,13 +70,15 @@ class memory_array */ void initialize(void * ptr, size_t sz, bool init) { - this->ptr = static_cast(ptr); - - - // Initialize the constructors - + // A placement array-new is allowed to reserve an implementation-defined + // cookie before the first element. Apple Clang does so for arrays of + // non-trivial OpenFPM objects. Keep the pointer returned by new[] rather + // than assuming that it is identical to the raw allocation; the memory + // backends deliberately provide a small allocation guard for this cookie. if (init == false) - new (ptr)T[sz]; + this->ptr = new (ptr)T[sz]; + else + this->ptr = static_cast(ptr); this->sz = sz; } diff --git a/src/timer.hpp b/src/timer.hpp index f87a0e1c..45199cbf 100644 --- a/src/timer.hpp +++ b/src/timer.hpp @@ -157,7 +157,24 @@ class timer cstart = cstop; } - #if defined(__NVCC__) && !defined(CUDA_ON_CPU) + #if defined(__NVCC__) && defined(CUDA_ON_CPU) + + void startGPU() + { + start(); + } + + void stopGPU() + { + stop(); + } + + double getwctGPU() + { + return getwct() * 1000.0; + } + + #elif defined(__NVCC__) void startGPU() { @@ -201,4 +218,3 @@ class timer #endif - diff --git a/src/util/cuda/sort_ofp.cuh b/src/util/cuda/sort_ofp.cuh index 4cec54a6..e45fbf66 100644 --- a/src/util/cuda/sort_ofp.cuh +++ b/src/util/cuda/sort_ofp.cuh @@ -13,6 +13,8 @@ #include "util/cuda_util.hpp" #include "util/ofp_context.hpp" +#include +#include #if CUDART_VERSION >= 11000 // Here we have for sure CUDA >= 11 @@ -265,9 +267,25 @@ namespace openfpm { #ifdef CUDA_ON_CPU - key_val_it kv(keys_input,vals_input); + // libc++ rejects this file's proxy iterator as a random-access iterator. + // Materialize the pairs for the CPU backend; accelerator paths stay intact. + std::vector> pairs; + pairs.reserve(static_cast(count)); + for (int i = 0; i < count; ++i) + pairs.emplace_back(keys_input[i],vals_input[i]); - std::sort(kv,kv+count,comp); + std::sort(pairs.begin(),pairs.end(), + [&comp](const key_val & a, + const key_val & b) + { + return comp(a.key,b.key); + }); + + for (int i = 0; i < count; ++i) + { + keys_input[i] = pairs[static_cast(i)].key; + vals_input[i] = pairs[static_cast(i)].val; + } #else #ifdef __HIP__ diff --git a/src/util/for_each_ref.hpp b/src/util/for_each_ref.hpp index 7e264e25..7ed4f5c0 100644 --- a/src/util/for_each_ref.hpp +++ b/src/util/for_each_ref.hpp @@ -75,10 +75,16 @@ struct for_each_ref_impl typedef typename deref::type item; typedef typename apply1::type arg; - // dwa 2002/9/10 -- make sure not to invoke undefined behavior - // when we pass arg. + // SPIR-V HIP compilation cannot call Boost's host-only value_init/unwrap + // helpers. MPL iteration values are empty integral/type wrappers, so + // direct value initialization has the same semantics. +#if defined(__HIP_DEVICE_COMPILE__) && defined(__HIP_PLATFORM_SPIRV__) + arg x{}; + f(x); +#else value_initialized x; aux::unwrap(f, 0)(boost::get(x)); +#endif typedef typename mpl::next::type iter; for_each_ref_impl::value> diff --git a/src/util/mathutil.hpp b/src/util/mathutil.hpp index c2235f76..7019e375 100644 --- a/src/util/mathutil.hpp +++ b/src/util/mathutil.hpp @@ -197,11 +197,11 @@ namespace openfpm * \return the bound number * */ - template + template __device__ __host__ static inline - T periodic_l(const T & pos, const T & p2, const T & p1) + T periodic_l(const Position & pos, const T & p2, const T & p1) { - T pos_tmp = pos; + T pos_tmp = static_cast(pos); if (pos >= p2) {