Support MoltenVK kernels and cudified CPU backends in OpenFPM data structures

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()
-

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)
-

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<unsigned int dim, typename grid_type>
 __global__ void copy_ndim_grid_block_device(grid_type src, grid_type dst)

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)

src/Grid/cuda/cuda_grid_unit_tests_func.cu

@@ -9,29 +9,29 @@ __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<typename grid_type>
 __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<typename grid_type>
 __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<typename grid_type>
@@ -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);
 }
-

src/Grid/grid_base_implementation.hpp

@@ -36,7 +36,7 @@ struct skip_init<true,T>
 	}
 };
 
-#ifdef __NVCC__
+#if defined(__NVCC__) || defined(CUDIFY_USE_METAL)
 
 template<bool active>
 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<dim,T,S,layout_base,ord_type> & 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.
 

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<unsigned int dim,
 	typename T,
@@ -460,7 +465,8 @@ class CellList_gpu<dim,T,Memory,transform_type,false> : 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<dim,T,Memory,transform_type,false> : 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<dim,T,Memory,transform_type,true> : 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<dim,T,Memory,transform_type,true> : 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<layout_base,T,4>
 		T::is_sparse_::value> type;
 };
 
+#undef OPENFPM_CELL_LIST_CONSTEXPR_IF
+
 #endif
 
 #endif /* OPENFPM_DATA_SRC_NN_CELLLIST_CELLLIST_GPU_HPP_ */

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);
+			}
 		}
 	}
 

src/NN/CellList/tests/CellDecomposer_gpu_ker_unit_test.cu

@@ -19,30 +19,30 @@ BOOST_AUTO_TEST_SUITE( CellDecomposer_gpu_test_suite )
 BOOST_AUTO_TEST_CASE( CellDecomposer_gpu_test_use )
 {
 	//! Spacing
-	openfpm::array<float,3> spacing_c = {0.1,0.1,0.1};
+	openfpm::array<float,3> spacing_c = {0.1f,0.1f,0.1f};
 
 	//! \brief number of sub-divisions in each direction
 	openfpm::array<unsigned int,3> div_c = {10,10,10};
 
 	//! \brief cell offset
 	openfpm::array<unsigned int,3> 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);
 
 	CellDecomposer_gpu_ker<3,float,unsigned int,shift_only<3,float>> clk(spacing_c,div_c,off,t);
 
 	openfpm::vector_gpu<aggregate<grid_key_dx<3,unsigned int>>> 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>();