Skip line-length check in cmake-lint

* Update GPU memory management
* Add allocate_memory and free_memory

.github/workflows/test.yml

@@ -13,6 +13,11 @@ jobs:
             gres: gpu:A4000
             cmake_flags: "-DTRIGDX_USE_MKL=1 -DTRIGDX_USE_GPU=1 -DTRIGDX_USE_MKL=1 -DTRIGDX_USE_XSIMD=1 -DCMAKE_CUDA_ARCHITECTURES=86"
             environment_modules: "spack/20250403 intel-oneapi-mkl cuda python"
+          - name: Intel, NVIDIA A4000
+            partition: defq
+            gres: gpu:A4000
+            cmake_flags: "-DTRIGDX_USE_MKL=1 -DTRIGDX_USE_GPU=1 -DTRIGDX_USE_MKL=1 -DTRIGDX_USE_XSIMD=1 -DCMAKE_CUDA_ARCHITECTURES=86"
+            environment_modules: "spack/20250403 intel-oneapi-compilers intel-oneapi-mkl cuda python"
           - name: NVIDIA GH200
             partition: ghq
             gres: gpu:GH200
@@ -38,7 +43,7 @@ jobs:
             cmake -S . -B build ${CMAKE_FLAGS}
             make -C build -j
       - name: Upload build
-        uses: actions/upload-artifact@v4
+        uses: pyTooling/upload-artifact@v4
         with:
           name: build-${{ matrix.name }}
           path: build
@@ -53,7 +58,7 @@ jobs:
       PARTITION_NAME: ${{ matrix.partition }}
     steps:
       - *cleanup
-      - uses: actions/download-artifact@v4
+      - uses: pyTooling/download-artifact@v4
         with:
           name: build-${{ matrix.name }}
       - uses: astron-rd/slurm-action@v1.2
@@ -61,7 +66,7 @@ jobs:
           partition: ${{ matrix.partition }}
           gres: ${{ matrix.gres }}
           commands: |
-            find tests -type f -executable -exec {} \;
+            find build/tests -type f -executable -exec {} \;
 
   benchmark:
     runs-on: [slurm]
@@ -73,7 +78,7 @@ jobs:
       PARTITION_NAME: ${{ matrix.partition }}
     steps:
       - *cleanup
-      - uses: actions/download-artifact@v4
+      - uses: pyTooling/download-artifact@v4
         with:
           name: build-${{ matrix.name }}
       - uses: astron-rd/slurm-action@v1.2
@@ -81,4 +86,4 @@ jobs:
           partition: ${{ matrix.partition }}
           gres: ${{ matrix.gres }}
           commands: |
-            find benchmarks -type f -executable -exec {} \;
+            find build/benchmarks -type f -executable -exec {} \;

.pre-commit-config.yaml

@@ -8,3 +8,4 @@ repos:
     hooks:
       - id: cmake-format
       - id: cmake-lint
+        args: [--disabled-codes=C0301]

benchmarks/benchmark_utils.hpp

@@ -2,13 +2,14 @@
 
 #include <chrono>
 #include <cmath>
+#include <stdexcept>
 #include <string>
 #include <vector>
 
 #include <benchmark/benchmark.h>
 
-void init_x(std::vector<float> &x) {
-  for (size_t i = 0; i < x.size(); ++i) {
+void init_x(float *x, size_t n) {
+  for (size_t i = 0; i < n; ++i) {
     x[i] = (i % 360) * 0.0174533f; // degrees to radians
   }
 }
@@ -16,24 +17,31 @@ void init_x(std::vector<float> &x) {
 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.data(), s.data());
+    backend.compute_sinf(N, x, s);
     benchmark::DoNotOptimize(s);
   }
 
+  backend.free_memory(x);
+  backend.free_memory(s);
+
   state.SetItemsProcessed(static_cast<int64_t>(state.iterations()) *
                           static_cast<int64_t>(N));
 }
@@ -41,24 +49,35 @@ 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.data(), c.data());
+    backend.compute_cosf(N, x, c);
     benchmark::DoNotOptimize(c);
   }
 
+  backend.free_memory(x);
+  backend.free_memory(c);
+
   state.SetItemsProcessed(static_cast<int64_t>(state.iterations()) *
                           static_cast<int64_t>(N));
 }
@@ -66,25 +85,38 @@ 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.data(), s.data(), c.data());
+    backend.compute_sincosf(N, x, s, c);
     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,7 +11,8 @@ public:
   GPUBackend();
   ~GPUBackend() override;
 
-  void init(size_t n = 0) override;
+  void *allocate_memory(size_t bytes) const 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,

include/trigdx/interface.hpp

@@ -1,6 +1,8 @@
 #pragma once
 
 #include <cstddef>
+#include <cstdint>
+#include <cstdlib>
 
 // Base interface for all math backends
 class Backend {
@@ -10,6 +12,12 @@ 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;
 

python/CMakeLists.txt

@@ -1,4 +1,6 @@
-if(NOT TARGET pybind11)
+find_package(pybind11 CONFIG QUIET)
+
+if(NOT pybind11_FOUND)
   FetchContent_Declare(
     pybind11
     GIT_REPOSITORY https://github.com/pybind/pybind11.git
@@ -6,5 +8,16 @@ if(NOT TARGET pybind11)
   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

@@ -0,0 +1,16 @@
+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,7 +72,9 @@ void bind_backend(py::module &m, const char *name) {
       .def("compute_sincosf", &compute_sincos<float>);
 }
 
-PYBIND11_MODULE(pytrigdx, m) {
+PYBIND11_MODULE(trigdx, m) {
+  m.doc() = "TrigDx python bindings";
+
   py::class_<Backend, std::shared_ptr<Backend>>(m, "Backend")
       .def("init", &Backend::init);
 

src/gpu.cpp

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