example_pybind11

cuvec/CMakeLists.txt

@@ -157,7 +157,7 @@ endif()
 # example projects
 
 add_subdirectory(src/example_mod)
-#add_subdirectory(src/example_pybind11)
+add_subdirectory(src/example_pybind11)
 add_subdirectory(src/example_swig)
 
 # install project

cuvec/src/example_pybind11/CMakeLists.txt

@@ -0,0 +1,25 @@
+project(example_pybind11)
+file(GLOB SRC LIST_DIRECTORIES false "*.cu")
+
+if(CUDAToolkit_FOUND)
+  include_directories(${CUDAToolkit_INCLUDE_DIRS})
+endif()
+python_add_library(${PROJECT_NAME} MODULE WITH_SOABI ${SRC})
+target_include_directories(${PROJECT_NAME} PUBLIC
+  "$<BUILD_INTERFACE:${${CMAKE_PROJECT_NAME}_INCLUDE_DIRS}>"
+  "$<INSTALL_INTERFACE:${CMAKE_PROJECT_NAME}/include>")
+if(CUDAToolkit_FOUND)
+  target_link_libraries(${PROJECT_NAME} PRIVATE pybind11::headers CUDA::cudart_static)
+else()
+  set_source_files_properties(${SRC} PROPERTIES LANGUAGE CXX)
+  target_link_libraries(${PROJECT_NAME} PRIVATE pybind11::headers)
+endif()
+
+set_target_properties(${PROJECT_NAME} PROPERTIES
+  CXX_STANDARD 11
+  VERSION ${CMAKE_PROJECT_VERSION} SOVERSION ${CMAKE_PROJECT_VERSION_MAJOR}
+  INTERFACE_${PROJECT_NAME}_MAJOR_VERSION ${CMAKE_PROJECT_VERSION_MAJOR})
+set_property(TARGET ${PROJECT_NAME} APPEND PROPERTY COMPATIBLE_INTERFACE_STRING ${PROJECT_NAME}_MAJOR_VERSION)
+install(TARGETS ${PROJECT_NAME}
+  INCLUDES DESTINATION ${CMAKE_PROJECT_NAME}/include
+  LIBRARY DESTINATION ${CMAKE_PROJECT_NAME})

cuvec/src/example_pybind11/example_pybind11.cu

@@ -0,0 +1,72 @@
+/**
+ * Example external pybind11 extension module using CuVec.
+ *
+ * Copyright (2021) Casper da Costa-Luis
+ */
+#include "cuvec_pybind11.cuh"  // NDCuVec
+#include <pybind11/pybind11.h> // pybind11, PYBIND11_MODULE
+#include <stdexcept>           // std::length_error
+#ifdef CUVEC_DISABLE_CUDA
+#include <chrono> // std::chrono
+#else
+/// dst = src + 1
+__global__ void _d_incr(float *dst, float *src, int X, int Y) {
+  int x = threadIdx.x + blockDim.x * blockIdx.x;
+  if (x >= X) return;
+  int y = threadIdx.y + blockDim.y * blockIdx.y;
+  if (y >= Y) return;
+  dst[y * X + x] = src[y * X + x] + 1;
+}
+#endif // CUVEC_DISABLE_CUDA
+NDCuVec<float> *increment2d_f(NDCuVec<float> &src, NDCuVec<float> *output, bool timing) {
+  auto &N = src.shape;
+  if (N.size() != 2) throw std::length_error("`src` must be 2D");
+
+#ifndef CUVEC_DISABLE_CUDA
+  cudaEvent_t eStart, eAlloc, eKern;
+  cudaEventCreate(&eStart);
+  cudaEventCreate(&eAlloc);
+  cudaEventCreate(&eKern);
+  cudaEventRecord(eStart);
+#else
+  auto eStart = std::chrono::steady_clock::now();
+#endif
+
+  if (!output)
+    output = new NDCuVec<float>(N);
+  else if (N != output->shape)
+    throw std::length_error("`output` must be same shape as `src`");
+
+#ifndef CUVEC_DISABLE_CUDA
+  cudaEventRecord(eAlloc);
+  dim3 thrds((N[1] + 31) / 32, (N[0] + 31) / 32);
+  dim3 blcks(32, 32);
+  _d_incr<<<thrds, blcks>>>(output->vec.data(), src.vec.data(), N[1], N[0]);
+  cuvec::HandleError(cudaGetLastError(), __FILE__, __LINE__);
+  // cudaDeviceSynchronize();
+  cudaEventRecord(eKern);
+  cudaEventSynchronize(eKern);
+  float alloc_ms, kernel_ms;
+  cudaEventElapsedTime(&alloc_ms, eStart, eAlloc);
+  cudaEventElapsedTime(&kernel_ms, eAlloc, eKern);
+// fprintf(stderr, "%.3f ms, %.3f ms\n", alloc_ms, kernel_ms);
+#else
+  auto eAlloc = std::chrono::steady_clock::now();
+  for (size_t i = 0; i < src.vec.size(); i++) output->vec[i] = src.vec[i] + 1;
+  auto eKern = std::chrono::steady_clock::now();
+  double alloc_ms = std::chrono::duration<double, std::milli>(eAlloc - eStart).count();
+  double kernel_ms = std::chrono::duration<double, std::milli>(eKern - eAlloc).count();
+// fprintf(stderr, "%.3lf ms, %.3lf ms\n", alloc_ms, kernel_ms);
+#endif
+  if (timing) {
+    // hack: store times in first two elements of output
+    output->vec[0] = alloc_ms;
+    output->vec[1] = kernel_ms;
+  }
+  return output;
+}
+
+using namespace pybind11::literals; // _a
+PYBIND11_MODULE(example_pybind11, m) {
+  m.def("increment2d_f", &increment2d_f, "src"_a, "output"_a = nullptr, "timing"_a = false);
+}

cuvec/src/example_swig/CMakeLists.txt

@@ -2,8 +2,6 @@ project(example_swig)
 file(GLOB SRC LIST_DIRECTORIES false "*.cu")
 file(GLOB ISRC LIST_DIRECTORIES false "*.i")
 
-#include_directories(${Python_INCLUDE_DIRS})
-
 if(SWIG_FOUND)
   if(CUDAToolkit_FOUND)
     include_directories(${CUDAToolkit_INCLUDE_DIRS})

tests/test_perf.py

@@ -17,6 +17,14 @@
 except ImportError:
     sw, example_swig = None, None  # type: ignore # yapf: disable
 
+try:
+    # `cuvec.pybind11` alternative to `cuvec.cpython`
+    # `example_pybind11` is defined in ../cuvec/src/example_pybind11/
+    from cuvec import example_pybind11  # type: ignore # yapf: disable
+    from cuvec import pybind11 as py
+except ImportError:
+    py, example_pybind11 = None, None  # type: ignore # yapf: disable
+
 
 def _time_overhead():
     tic = time()
@@ -53,7 +61,7 @@ def test_inner(*args, **kwargs):
     return wrapper
 
 
-@mark.parametrize("cu,ex", [(cu, example_mod), (sw, example_swig)])
+@mark.parametrize("cu,ex", [(cu, example_mod), (sw, example_swig), (py, example_pybind11)])
 @retry_on_except()
 def test_perf(cu, ex, shape=(1337, 42), quiet=False, return_time=False):
     if cu is None:

tests/test_pybind11.py

@@ -145,3 +145,25 @@ def test_cuda_array_interface():
     assert ndarr.dtype == v.dtype
     with raises(AttributeError):
         ndarr.__cuda_array_interface__
+
+
+def test_increment():
+    # `example_pybind11` is defined in ../cuvec/src/example_pybind11/
+    from cuvec.example_pybind11 import increment2d_f
+    a = cu.zeros((1337, 42), 'f')
+    assert (a == 0).all()
+    increment2d_f(a.cuvec, a.cuvec)
+    assert (a == 1).all()
+
+    a[:] = 0
+    assert (a == 0).all()
+
+    b = cu.retarray(increment2d_f(a.cuvec))
+    assert (b == 1).all()
+
+    c = cu.retarray(increment2d_f(b.cuvec, a.cuvec), a)
+    assert (a == 2).all()
+    assert c.cuvec == a.cuvec
+    assert (c == a).all()
+    assert str(c.pyvec) == str(a.pyvec)
+    assert np.asarray(c.pyvec).data == np.asarray(a.pyvec).data