7 Commits

Author SHA1 Message Date
Wiebe van Breukelen
38664f6acb Fix compiler warnings 2025-10-22 15:20:50 +02:00
mmancini-skao
76998a137a Fix error in Taylor expansion (#36)
Replaced term3 with term4 in calculations for t4.
2025-10-20 17:09:35 +02:00
Bram Veenboer
500d35070e Fix formatting (#35)
* Run pre-commit

* Skip line-length check in cmake-lint
2025-10-10 09:19:18 +02:00
Dantali0n
bfe752433f Fixes #30, Add CMake steps to install python bindings (#31) 2025-09-17 20:03:28 +02:00
Bram Veenboer
8fe8314905 Update GPU backend (#29)
* Update GPU memory management
* Add allocate_memory and free_memory
2025-09-03 09:16:28 +02:00
Wiebe van Breukelen
9d3af8c202 Fixed broken pybind11 target check (#28)
* Fixed broken pybind11 target check

* Update python/CMakeLists.txt

Co-authored-by: Bram Veenboer <bram.veenboer@gmail.com>

---------

Co-authored-by: Bram Veenboer <bram.veenboer@gmail.com>
2025-08-27 17:07:30 +02:00
Bram Veenboer
0f7fd06be8 Extend CI (#27)
* Add build with Intel compiler
* Switch to upload/download-artifacts that retain permissions
2025-08-21 15:15:25 +02:00
11 changed files with 144 additions and 81 deletions

View File

@@ -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 {} \;

View File

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

View File

@@ -12,6 +12,11 @@ option(TRIGDX_BUILD_TESTS "Build tests" ON)
option(TRIGDX_BUILD_BENCHMARKS "Build tests" ON)
option(TRIGDX_BUILD_PYTHON "Build Python interface" ON)
# Add compiler flags
set(CMAKE_CXX_FLAGS
"${CMAKE_CXX_FLAGS} -Wall -Wnon-virtual-dtor -Wduplicated-branches -Wvla -Wpointer-arith -Wextra -Wno-unused-parameter"
)
list(APPEND CMAKE_MODULE_PATH "${PROJECT_SOURCE_DIR}/cmake")
configure_file(
${CMAKE_CURRENT_SOURCE_DIR}/cmake/trigdx_config.hpp.in

View File

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

View File

@@ -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,

View File

@@ -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;

View File

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

16
python/__init__.py Normal file
View File

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

View File

@@ -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);
@@ -91,4 +93,4 @@ PYBIND11_MODULE(pytrigdx, m) {
bind_backend<LookupXSIMDBackend<16384>>(m, "LookupXSIMD16K");
bind_backend<LookupXSIMDBackend<32768>>(m, "LookupXSIMD32K");
#endif
}
}

View File

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

View File

@@ -20,8 +20,8 @@ template <std::size_t NR_SAMPLES> struct lookup_table {
cos_values[i] = cosf(i * PI_FRAC);
}
}
std::array<float, NR_SAMPLES> cos_values;
std::array<float, NR_SAMPLES> sin_values;
std::array<float, NR_SAMPLES> cos_values;
};
template <std::size_t NR_SAMPLES> struct cosf_dispatcher {
@@ -33,7 +33,6 @@ template <std::size_t NR_SAMPLES> struct cosf_dispatcher {
constexpr uint_fast32_t VL = b_type::size;
const uint_fast32_t VS = n - n % VL;
const uint_fast32_t Q_PI = NR_SAMPLES / 4U;
const b_type scale = b_type::broadcast(lookup_table_.SCALE);
const b_type pi_frac = b_type::broadcast(lookup_table_.PI_FRAC);
const m_type mask = m_type::broadcast(lookup_table_.MASK);
@@ -42,7 +41,7 @@ template <std::size_t NR_SAMPLES> struct cosf_dispatcher {
const b_type term2 = b_type::broadcast(lookup_table_.TERM2); // 1/2!
const b_type term3 = b_type::broadcast(lookup_table_.TERM3); // 1/3!
const b_type term4 = b_type::broadcast(lookup_table_.TERM4); // 1/4!
const m_type quarter_pi = m_type::broadcast(Q_PI);
uint_fast32_t i;
for (i = 0; i < VS; i += VL) {
const b_type vx = b_type::load(a + i, Tag());
@@ -60,7 +59,7 @@ template <std::size_t NR_SAMPLES> struct cosf_dispatcher {
const b_type dx4 = xsimd::mul(dx2, dx);
const b_type t2 = xsimd::mul(dx2, term2);
const b_type t3 = xsimd::mul(dx3, term3);
const b_type t4 = xsimd::mul(dx4, term3);
const b_type t4 = xsimd::mul(dx4, term4);
const b_type cosdx = xsimd::add(xsimd::sub(term1, t2), t4);
@@ -98,7 +97,6 @@ template <std::size_t NR_SAMPLES> struct sinf_dispatcher {
constexpr uint_fast32_t VL = b_type::size;
const uint_fast32_t VS = n - n % VL;
const uint_fast32_t Q_PI = NR_SAMPLES / 4U;
const b_type scale = b_type::broadcast(lookup_table_.SCALE);
const b_type pi_frac = b_type::broadcast(lookup_table_.PI_FRAC);
const m_type mask = m_type::broadcast(lookup_table_.MASK);
@@ -107,7 +105,7 @@ template <std::size_t NR_SAMPLES> struct sinf_dispatcher {
const b_type term2 = b_type::broadcast(lookup_table_.TERM2); // 1/2!
const b_type term3 = b_type::broadcast(lookup_table_.TERM3); // 1/3!
const b_type term4 = b_type::broadcast(lookup_table_.TERM4); // 1/4!
const m_type quarter_pi = m_type::broadcast(Q_PI);
uint_fast32_t i;
for (i = 0; i < VS; i += VL) {
const b_type vx = b_type::load(a + i, Tag());
@@ -120,7 +118,7 @@ template <std::size_t NR_SAMPLES> struct sinf_dispatcher {
const b_type dx4 = xsimd::mul(dx2, dx);
const b_type t2 = xsimd::mul(dx2, term2);
const b_type t3 = xsimd::mul(dx3, term3);
const b_type t4 = xsimd::mul(dx4, term3);
const b_type t4 = xsimd::mul(dx4, term4);
const b_type cosdx = xsimd::add(xsimd::sub(term1, t2), t4);
const b_type sindx = xsimd::sub(dx, t3);
@@ -160,7 +158,6 @@ template <std::size_t NR_SAMPLES> struct sin_cosf_dispatcher {
constexpr uint_fast32_t VL = b_type::size;
const uint_fast32_t VS = n - n % VL;
const uint_fast32_t Q_PI = NR_SAMPLES / 4U;
const b_type scale = b_type::broadcast(lookup_table_.SCALE);
const m_type mask = m_type::broadcast(lookup_table_.MASK);
const b_type pi_frac = b_type::broadcast(lookup_table_.PI_FRAC);
@@ -170,7 +167,6 @@ template <std::size_t NR_SAMPLES> struct sin_cosf_dispatcher {
const b_type term3 = b_type::broadcast(lookup_table_.TERM3); // 1/3!
const b_type term4 = b_type::broadcast(lookup_table_.TERM4); // 1/4!
const m_type quarter_pi = m_type::broadcast(Q_PI);
uint_fast32_t i;
for (i = 0; i < VS; i += VL) {
const b_type vx = b_type::load(a + i, Tag());
@@ -183,7 +179,7 @@ template <std::size_t NR_SAMPLES> struct sin_cosf_dispatcher {
const b_type dx4 = xsimd::mul(dx2, dx);
const b_type t2 = xsimd::mul(dx2, term2);
const b_type t3 = xsimd::mul(dx3, term3);
const b_type t4 = xsimd::mul(dx4, term3);
const b_type t4 = xsimd::mul(dx4, term4);
idx = xsimd::bitwise_and(idx, mask);
b_type sinv = b_type::gather(lookup_table_.sin_values.data(), idx);