TODO: first changes to add expf

.pre-commit-config.yaml

@@ -7,5 +7,4 @@ repos:
     rev: v0.6.13
     hooks:
       - id: cmake-format
-      - id: cmake-lint
+      - id: cmake-lint
-        args: [--disabled-codes=C0301]

benchmarks/benchmark_utils.hpp

@@ -2,14 +2,13 @@
 
 #include <chrono>
 #include <cmath>
-#include <stdexcept>
 #include <string>
 #include <vector>
 
 #include <benchmark/benchmark.h>
 
-void init_x(float *x, size_t n) {
+void init_x(std::vector<float> &x) {
-  for (size_t i = 0; i < n; ++i) {
+  for (size_t i = 0; i < x.size(); ++i) {
     x[i] = (i % 360) * 0.0174533f; // degrees to radians
   }
 }
@@ -17,31 +16,24 @@ void init_x(float *x, size_t n) {
 template <typename Backend>
 static void benchmark_sinf(benchmark::State &state) {
   const size_t N = static_cast<size_t>(state.range(0));
+  std::vector<float> x(N), s(N);
+  init_x(x);
 
   Backend backend;
 
   auto start = std::chrono::high_resolution_clock::now();
   backend.init(N);
-  float *x =
-      reinterpret_cast<float *>(backend.allocate_memory(N * sizeof(float)));
-  float *s =
-      reinterpret_cast<float *>(backend.allocate_memory(N * sizeof(float)));
   auto end = std::chrono::high_resolution_clock::now();
   state.counters["init_ms"] =
       std::chrono::duration_cast<std::chrono::microseconds>(end - start)
           .count() /
       1.e3;
 
-  init_x(x, N);
-
   for (auto _ : state) {
-    backend.compute_sinf(N, x, s);
+    backend.compute_sinf(N, x.data(), s.data());
     benchmark::DoNotOptimize(s);
   }
 
-  backend.free_memory(x);
-  backend.free_memory(s);
-
   state.SetItemsProcessed(static_cast<int64_t>(state.iterations()) *
                           static_cast<int64_t>(N));
 }
@@ -49,35 +41,24 @@ static void benchmark_sinf(benchmark::State &state) {
 template <typename Backend>
 static void benchmark_cosf(benchmark::State &state) {
   const size_t N = static_cast<size_t>(state.range(0));
+  std::vector<float> x(N), c(N);
+  init_x(x);
 
   Backend backend;
 
   auto start = std::chrono::high_resolution_clock::now();
   backend.init(N);
-  float *x =
-      reinterpret_cast<float *>(backend.allocate_memory(N * sizeof(float)));
-  float *c =
-      reinterpret_cast<float *>(backend.allocate_memory(N * sizeof(float)));
-
-  if (!x || !c) {
-    throw std::runtime_error("Buffer allocation failed");
-  }
   auto end = std::chrono::high_resolution_clock::now();
   state.counters["init_ms"] =
       std::chrono::duration_cast<std::chrono::microseconds>(end - start)
           .count() /
       1.e3;
 
-  init_x(x, N);
-
   for (auto _ : state) {
-    backend.compute_cosf(N, x, c);
+    backend.compute_cosf(N, x.data(), c.data());
     benchmark::DoNotOptimize(c);
   }
 
-  backend.free_memory(x);
-  backend.free_memory(c);
-
   state.SetItemsProcessed(static_cast<int64_t>(state.iterations()) *
                           static_cast<int64_t>(N));
 }
@@ -85,38 +66,25 @@ static void benchmark_cosf(benchmark::State &state) {
 template <typename Backend>
 static void benchmark_sincosf(benchmark::State &state) {
   const size_t N = static_cast<size_t>(state.range(0));
+  std::vector<float> x(N), s(N), c(N);
+  init_x(x);
 
   Backend backend;
 
   auto start = std::chrono::high_resolution_clock::now();
   backend.init(N);
-  float *x =
-      reinterpret_cast<float *>(backend.allocate_memory(N * sizeof(float)));
-  float *s =
-      reinterpret_cast<float *>(backend.allocate_memory(N * sizeof(float)));
-  float *c =
-      reinterpret_cast<float *>(backend.allocate_memory(N * sizeof(float)));
-  if (!x || !s || !c) {
-    throw std::runtime_error("Buffer allocation failed");
-  }
   auto end = std::chrono::high_resolution_clock::now();
   state.counters["init_ms"] =
       std::chrono::duration_cast<std::chrono::microseconds>(end - start)
           .count() /
       1.e3;
 
-  init_x(x, N);
-
   for (auto _ : state) {
-    backend.compute_sincosf(N, x, s, c);
+    backend.compute_sincosf(N, x.data(), s.data(), c.data());
     benchmark::DoNotOptimize(s);
     benchmark::DoNotOptimize(c);
   }
 
-  backend.free_memory(x);
-  backend.free_memory(s);
-  backend.free_memory(c);
-
   state.SetItemsProcessed(static_cast<int64_t>(state.iterations()) *
                           static_cast<int64_t>(N));
 }

include/trigdx/gpu.hpp

@@ -11,12 +11,12 @@ public:
   GPUBackend();
   ~GPUBackend() override;
 
-  void *allocate_memory(size_t bytes) const override;
+  void init(size_t n = 0) override;
-  void free_memory(void *ptr) const override;
   void compute_sinf(size_t n, const float *x, float *s) const override;
   void compute_cosf(size_t n, const float *x, float *c) const override;
   void compute_sincosf(size_t n, const float *x, float *s,
                        float *c) const override;
+  void compute_expf(size_t n, const float *x, float *e) const override;
 
 private:
   struct Impl;

include/trigdx/interface.hpp

@@ -1,8 +1,6 @@
 #pragma once
 
 #include <cstddef>
-#include <cstdint>
-#include <cstdlib>
 
 // Base interface for all math backends
 class Backend {
@@ -12,12 +10,6 @@ public:
   // Optional initialization
   virtual void init(size_t n = 0) {}
 
-  virtual void *allocate_memory(size_t bytes) const {
-    return static_cast<void *>(new uint8_t[bytes]);
-  };
-
-  virtual void free_memory(void *ptr) const { std::free(ptr); };
-
   // Compute sine for n elements
   virtual void compute_sinf(size_t n, const float *x, float *s) const = 0;
 
@@ -27,4 +19,7 @@ public:
   // Compute sine and cosine for n elements
   virtual void compute_sincosf(size_t n, const float *x, float *s,
                                float *c) const = 0;
+
+  // Compute exponent for n elements
+  virtual void compute_expf(size_t n, const float *x, float *e) const = 0;
 };

include/trigdx/mkl.hpp

@@ -10,4 +10,6 @@ public:
 
   void compute_sincosf(size_t n, const float *x, float *s,
                        float *c) const override;
+
+  void compute_expf(size_t n, const float *x, float *e) const override;
 };

include/trigdx/reference.hpp

@@ -10,4 +10,6 @@ public:
 
   void compute_sincosf(size_t n, const float *x, float *s,
                        float *c) const override;
+
+  void compute_expf(size_t n, const float *x, float *e) const override;
 };

python/CMakeLists.txt

@@ -8,16 +8,5 @@ if(NOT pybind11_FOUND)
   FetchContent_MakeAvailable(pybind11)
 endif()
 
-# Needed to set ${Python_VERSION_MAJOR} and ${Python_VERSION_MINOR}
-find_package(Python REQUIRED)
-
 pybind11_add_module(pytrigdx bindings.cpp)
 target_link_libraries(pytrigdx PRIVATE trigdx)
-set_target_properties(pytrigdx PROPERTIES OUTPUT_NAME "trigdx")
-
-set(PYTHON_SITE_PACKAGES
-    "${CMAKE_INSTALL_LIBDIR}/python${Python_VERSION_MAJOR}.${Python_VERSION_MINOR}/site-packages/trigdx"
-)
-
-install(TARGETS pytrigdx DESTINATION ${PYTHON_SITE_PACKAGES})
-install(FILES __init__.py DESTINATION ${PYTHON_SITE_PACKAGES})

python/__init__.py

@@ -1,16 +0,0 @@
-from .trigdx import Reference, Lookup16K, Lookup32K, LookupAVX16K, LookupAVX32K
-
-try:
-    from .trigdx import MKL
-except ImportError:
-    pass
-
-try:
-    from .trigdx import GPU
-except ImportError:
-    pass
-
-try:
-    from .trigdx import LookupXSIMD16K, LookupXSIMD32K
-except ImportError:
-    pass

python/bindings.cpp

@@ -72,9 +72,7 @@ void bind_backend(py::module &m, const char *name) {
       .def("compute_sincosf", &compute_sincos<float>);
 }
 
-PYBIND11_MODULE(trigdx, m) {
+PYBIND11_MODULE(pytrigdx, m) {
-  m.doc() = "TrigDx python bindings";
-
   py::class_<Backend, std::shared_ptr<Backend>>(m, "Backend")
       .def("init", &Backend::init);
 
@@ -93,4 +91,4 @@ PYBIND11_MODULE(trigdx, m) {
   bind_backend<LookupXSIMDBackend<16384>>(m, "LookupXSIMD16K");
   bind_backend<LookupXSIMDBackend<32768>>(m, "LookupXSIMD32K");
 #endif
-}
+}

src/gpu.cpp

@@ -10,63 +10,98 @@
 
 struct GPUBackend::Impl {
 
-  void *allocate_memory(size_t bytes) const {
+  ~Impl() {
-    void *ptr;
+    if (h_x) {
-    cudaMallocHost(&ptr, bytes);
+      cudaFreeHost(h_x);
-    return ptr;
+    }
+    if (h_s) {
+      cudaFreeHost(h_s);
+    }
+    if (h_c) {
+      cudaFreeHost(h_c);
+    }
+    if (h_e) {
+      cudaFreeHost(h_e);
+    }
+    if (d_x) {
+      cudaFree(d_x);
+    }
+    if (d_s) {
+      cudaFree(d_s);
+    }
+    if (d_c) {
+      cudaFree(d_c);
+    }
+    if (d_e) {
+      cudaFree(d_e);
+    }
   }
 
-  void free_memory(void *ptr) const { cudaFreeHost(ptr); }
+  void init(size_t n) {
+    const size_t bytes = n * sizeof(float);
+    cudaMallocHost(&h_x, bytes);
+    cudaMallocHost(&h_s, bytes);
+    cudaMallocHost(&h_c, bytes);
+    cudaMallocHost(&h_e, bytes);
+    cudaMalloc(&d_x, bytes);
+    cudaMalloc(&d_s, bytes);
+    cudaMalloc(&d_c, bytes);
+    cudaMalloc(&d_e, bytes);
+  }
 
   void compute_sinf(size_t n, const float *x, float *s) const {
     const size_t bytes = n * sizeof(float);
-    float *d_x, *d_s;
+    std::memcpy(h_x, x, bytes);
-    cudaMalloc(&d_x, bytes);
+    cudaMemcpy(d_x, h_x, bytes, cudaMemcpyHostToDevice);
-    cudaMalloc(&d_s, bytes);
-    cudaMemcpy(d_x, x, bytes, cudaMemcpyHostToDevice);
     launch_sinf_kernel(d_x, d_s, n);
-    cudaMemcpy(s, d_s, bytes, cudaMemcpyDeviceToHost);
+    cudaMemcpy(h_s, d_s, bytes, cudaMemcpyDeviceToHost);
-    cudaFree(d_x);
+    std::memcpy(s, h_s, bytes);
-    cudaFree(d_s);
   }
 
   void compute_cosf(size_t n, const float *x, float *c) const {
     const size_t bytes = n * sizeof(float);
-    float *d_x, *d_c;
+    std::memcpy(h_x, x, bytes);
-    cudaMalloc(&d_x, bytes);
+    cudaMemcpy(d_x, h_x, bytes, cudaMemcpyHostToDevice);
-    cudaMalloc(&d_c, bytes);
-    cudaMemcpy(d_x, x, bytes, cudaMemcpyHostToDevice);
     launch_cosf_kernel(d_x, d_c, n);
-    cudaMemcpy(c, d_c, bytes, cudaMemcpyDeviceToHost);
+    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);
-    cudaFree(d_x);
+    std::memcpy(c, h_c, bytes);
-    cudaFree(d_c);
   }
 
   void compute_sincosf(size_t n, const float *x, float *s, float *c) const {
     const size_t bytes = n * sizeof(float);
-    float *d_x, *d_s, *d_c;
+    std::memcpy(h_x, x, bytes);
-    cudaMalloc(&d_x, bytes);
+    cudaMemcpy(d_x, h_x, bytes, cudaMemcpyHostToDevice);
-    cudaMalloc(&d_s, bytes);
-    cudaMalloc(&d_c, bytes);
-    cudaMemcpy(d_x, x, bytes, cudaMemcpyHostToDevice);
     launch_sincosf_kernel(d_x, d_s, d_c, n);
-    cudaMemcpy(s, d_s, bytes, cudaMemcpyDeviceToHost);
+    cudaMemcpy(h_s, d_s, bytes, cudaMemcpyDeviceToHost);
-    cudaMemcpy(c, d_c, bytes, cudaMemcpyDeviceToHost);
+    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);
-    cudaFree(d_x);
+    std::memcpy(s, h_s, bytes);
-    cudaFree(d_s);
+    std::memcpy(c, h_c, bytes);
-    cudaFree(d_c);
   }
+
+  void compute_expf(size_t n, const float *x, float *e) const {
+    const size_t bytes = n * sizeof(float);
+    std::memcpy(h_x, x, bytes);
+    cudaMemcpy(d_x, h_x, bytes, cudaMemcpyHostToDevice);
+    launch_expf_kernel(d_x, d_e, n);
+    cudaMemcpy(h_e, d_e, bytes, cudaMemcpyDeviceToHost);
+    std::memcpy(e, h_e, bytes);
+  }
+
+  float *h_x = nullptr;
+  float *h_s = nullptr;
+  float *h_c = nullptr;
+  float *h_e = nullptr;
+  float *d_x = nullptr;
+  float *d_s = nullptr;
+  float *d_c = nullptr;
+  float *d_e = nullptr;
 };
 
 GPUBackend::GPUBackend() : impl(std::make_unique<Impl>()) {}
 
 GPUBackend::~GPUBackend() = default;
 
-void *GPUBackend::allocate_memory(size_t bytes) const {
+void GPUBackend::init(size_t n) { impl->init(n); }
-  return impl->allocate_memory(bytes);
-}
-
-void GPUBackend::free_memory(void *ptr) const { impl->free_memory(ptr); }
 
 void GPUBackend::compute_sinf(size_t n, const float *x, float *s) const {
   impl->compute_sinf(n, x, s);
@@ -79,4 +114,8 @@ void GPUBackend::compute_cosf(size_t n, const float *x, float *c) const {
 void GPUBackend::compute_sincosf(size_t n, const float *x, float *s,
                                  float *c) const {
   impl->compute_sincosf(n, x, s, c);
-}
+}
+
+void GPUBackend::compute_expf(size_t n, const float *x, float *e) const {
+  impl->compute_expf(n, x, e);
+}

src/gpu/gpu.cu

@@ -31,6 +31,15 @@ __global__ void kernel_sincosf(const float *__restrict__ x,
   }
 }
 
+__global__ void kernel_expf(const float *__restrict__ x, float *__restrict__ e,
+                            size_t n) {
+  size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
+  if (idx < n) {
+    // e[idx] = __expf(x[idx]);
+    e[idx] = expf(x[idx]);
+  }
+}
+
 namespace {
 inline dim3 make_grid(size_t n, size_t threadsPerBlock = 256) {
   return dim3((n + threadsPerBlock - 1) / threadsPerBlock);
@@ -54,3 +63,9 @@ void launch_sincosf_kernel(const float *d_x, float *d_s, float *d_c, size_t n) {
   dim3 grid = make_grid(n, blocks.x);
   kernel_sincosf<<<grid, blocks>>>(d_x, d_s, d_c, n);
 }
+
+void launch_expf_kernel(const float *d_x, float *d_e, size_t n) {
+  dim3 blocks(256);
+  dim3 grid = make_grid(n, blocks.x);
+  kernel_expf<<<grid, blocks>>>(d_x, d_e, n);
+}

src/gpu/gpu.cuh

@@ -6,3 +6,4 @@ void launch_sinf_kernel(const float *d_x, float *d_s, size_t n);
 void launch_cosf_kernel(const float *d_x, float *d_c, size_t n);
 void launch_sincosf_kernel(const float *d_x, float *d_s, float *d_c,
                            std::size_t n);
+void launch_expf_kernel(const float *d_x, float *d_e, size_t n);

src/mkl.cpp

@@ -14,3 +14,7 @@ void MKLBackend::compute_sincosf(size_t n, const float *x, float *s,
                                  float *c) const {
   vmsSinCos(static_cast<MKL_INT>(n), x, s, c, VML_HA);
 }
+
+void MKLBackend::compute_expf(size_t n, const float *x, float *e) const {
+  vmsExp(static_cast<MKL_INT>(n), x, e, VML_HA);
+}

src/reference.cpp

@@ -21,3 +21,9 @@ void ReferenceBackend::compute_sincosf(size_t n, const float *x, float *s,
     c[i] = cosf(x[i]);
   }
 }
+
+void ReferenceBackend::compute_expf(size_t n, const float *x, float *e) const {
+  for (size_t i = 0; i < n; ++i) {
+    e[i] = expf(x[i]);
+  }
+}

tests/test_mkl.cpp

@@ -8,3 +8,5 @@ TEST_CASE("sinf") { test_sinf<MKLBackend>(1e-6f); }
 TEST_CASE("cosf") { test_cosf<MKLBackend>(1e-6f); }
 
 TEST_CASE("sincosf") { test_sincosf<MKLBackend>(1e-6f); }
+
+TEST_CASE("expf") { test_expf<MKLBackend>(1e-6f); }

tests/test_utils.hpp

@@ -63,3 +63,19 @@ template <typename Backend> inline void test_sincosf(float tol) {
     REQUIRE_THAT(c[i], Catch::Matchers::WithinAbs(c_ref[i], tol));
   }
 }
+
+template <typename Backend> inline void test_expf(float tol) {
+  std::vector<float> x(N), e_ref(N), e(N);
+  init_x(x);
+
+  ReferenceBackend ref;
+  Backend backend;
+  backend.init(N);
+
+  ref.compute_expf(N, x.data(), e_ref.data());
+  backend.compute_expf(N, x.data(), e.data());
+
+  for (size_t i = 0; i < N; ++i) {
+    REQUIRE_THAT(e[i], Catch::Matchers::WithinAbs(e_ref[i], tol));
+  }
+}