From 93c9094272b839ef80f0e177561b5e462da5f520 Mon Sep 17 00:00:00 2001 From: keithlostracco Date: Wed, 27 May 2026 12:27:35 -0700 Subject: [PATCH 1/4] Add release docs, fix ExternalTime mode, add external time example --- README.md | 22 +++-- RELEASING.md | 45 ++++++++++ examples/basics/10_external_time.py | 135 ++++++++++++++++++++++++++++ source/comp.cpp | 38 +++++--- 4 files changed, 221 insertions(+), 19 deletions(-) create mode 100644 RELEASING.md create mode 100644 examples/basics/10_external_time.py diff --git a/README.md b/README.md index a45c2ab..d28258f 100644 --- a/README.md +++ b/README.md @@ -55,24 +55,30 @@ Full API documentation is available at [intentdev.github.io/touchpy](https://int ### Setup ```bash -uv venv --python 3.12 # torch doesn't fill in the metatdata name for 3.13/3.14 +uv venv --python 3.12 uv sync --extra examples ``` -This creates a local `.venv`, builds TouchPy from source, and installs CUDA-enabled PyTorch + numpy for running examples. +This creates a `.venv` with Python 3.12, builds TouchPy from source (via scikit-build-core), and installs CUDA-enabled PyTorch + numpy for running examples. -### Build a wheel +> **Note:** Python 3.12 is recommended. PyTorch wheels for 3.13+ may have packaging issues. + +### Rebuilding after C++ changes ```bash -uv build +uv sync --extra examples --reinstall-package touchpy ``` -### Local CMake build (for debugging / IDE integration) +Or equivalently: ```bash -# From a Visual Studio Developer Command Prompt -cmake --preset x64-release -cmake --build out/build/x64-release +uv pip install -ve . +``` + +### Build a wheel + +```bash +uv build ``` ### Verify diff --git a/RELEASING.md b/RELEASING.md new file mode 100644 index 0000000..5eff8c7 --- /dev/null +++ b/RELEASING.md @@ -0,0 +1,45 @@ +# Release Process + +## Version Convention + +- During development, `pyproject.toml` has version `X.Y.Z.dev0` +- At release time, the `.dev0` suffix is stripped +- After release, version is bumped to the next `X.Y.Z.dev0` + +## Release Checklist + +1. **Prepare the release PR:** + - Update `version` in `pyproject.toml`: remove `.dev0` suffix (e.g., `0.12.0.dev0` → `0.12.0`) + - Update CHANGELOG.md with release notes + - Merge PR to `main` + +2. **Create the release:** + - Create a git tag: `git tag v0.12.0` + - Push the tag: `git push origin v0.12.0` + - Create a GitHub Release from the tag at https://github.com/IntentDev/touchpy-dev/releases/new + - This triggers the `wheels.yml` workflow which: + - Builds wheels for Python 3.9-3.13 (Windows x64) + - Publishes to PyPI via trusted publishing + - Uploads wheels as GitHub Release assets + - This also triggers the `docs.yml` workflow which rebuilds and deploys documentation + +3. **Post-release:** + - Create a follow-up PR bumping version to next `.dev0` (e.g., `0.13.0.dev0`) + - Merge to `main` + +## PyPI Trusted Publishing Setup (one-time) + +Configure at https://pypi.org under your account → Publishing: + +- **Owner**: `IntentDev` +- **Repository**: `touchpy-dev` (update after repo rename to `touchpy`) +- **Workflow name**: `wheels.yml` +- **Environment**: (leave blank) + +## Manual Workflow Triggers + +All workflows can be triggered manually from the GitHub Actions tab: + +- **wheels.yml**: Build wheels without publishing (manual dispatch skips PyPI upload) +- **docs.yml**: Rebuild and deploy documentation +- **ci.yml**: Runs automatically on push/PR diff --git a/examples/basics/10_external_time.py b/examples/basics/10_external_time.py new file mode 100644 index 0000000..3a88407 --- /dev/null +++ b/examples/basics/10_external_time.py @@ -0,0 +1,135 @@ +import touchpy as tp +import time + +import torch +import modules.utils as utils +from modules.image_filter import ImageFilter + +# ── Configuration ────────────────────────────────────────────── +TD_FPS = 60 +HOST_FPS = 5 + +REALTIME = False # (every frame / offline) +# Steps through every 1/TD_FPS second sequentially. +# Host loop runs at HOST_FPS, but each iteration advances TD +# by exactly one frame (1/TD_FPS). No frames are skipped. +# At HOST_FPS=5 and TD_FPS=60 it takes 12 seconds of wall time +# to render 1 second of TD time. +# +# REALTIME = True +# Uses wall-clock elapsed time. Host loop runs at HOST_FPS. +# TD skips ahead to match real elapsed time, so ~12 TD frames +# are skipped per host iteration (60/5). TD time tracks wall time. +# ─────────────────────────────────────────────────────────────── + +class MyComp(tp.Comp): + def __init__(self): + super().__init__( + flags=tp.CompFlags.EXTERNAL_TIME | tp.CompFlags.CUDA_STREAM_DEFAULT, + fps=TD_FPS + ) + self.layout_ready = False + self.set_on_layout_change_callback(self.on_layout_change) + + self.device = torch.device('cuda:0') + self.frame_count = 0 + + self.imag_filter = ImageFilter( + in_channels=3, + out_channels=3, + kernel_size=12, + stride=2, + groups=3 + ).to(self.device) + + self.alpha_tensor = torch.ones(1, 1080, 1920, device=self.device, dtype=torch.float16) + + def on_layout_change(self): + print('layout changed:') + print('out tops:\n', *[f"\t{name}\n" for name in self.out_tops.names]) + print('in tops:\n', *[f"\t{name}\n" for name in self.in_tops.names]) + + if 'Monitortop' in self.par.names: + self.par['Monitortop'].val = 'topIn1' + + self.out_tops[0].set_cuda_flags(tp.CudaFlags.HWC) + self.out_tops[1].set_cuda_flags(tp.CudaFlags.RGB | tp.CudaFlags.CHW) + self.out_tops[2].set_cuda_flags(tp.CudaFlags.RGB) + + self.layout_ready = True + + +comp = MyComp() +comp.load('TopChopDatIO.tox', fps=TD_FPS) +comp.start() + +while not comp.layout_ready: + if comp.frame_did_finish(): + pass + time.sleep(0.001) + +mode_label = "REALTIME" if REALTIME else "EVERY FRAME" +print(f"\nRunning EXTERNAL_TIME | TD @ {TD_FPS} fps | host @ {HOST_FPS} fps | mode: {mode_label}") +print(f"Press 'q' to quit\n") + +frame_step = 1.0 / TD_FPS +host_period = 1.0 / HOST_FPS +current_time = 0.0 +start_wall = time.perf_counter() + +comp.start_next_frame(current_time) + +running = True +while running: + loop_start = time.perf_counter() + + if utils.check_key('q'): + # while not comp.frame_did_finish(): + # time.sleep(0.001) + break + + if comp.frame_did_finish(): + t = comp.time() + elapsed = time.perf_counter() - start_wall + print(f"frame {comp.frame_count:4d} | td time: {t.seconds:8.3f}s | wall: {elapsed:8.3f}s") + + out_top1_tensor = comp.out_tops[0].as_tensor() + out_top2_tensor = comp.out_tops[1].as_tensor() + out_top3_tensor = comp.out_tops[2].as_tensor() + + if REALTIME: + current_time = time.perf_counter() - start_wall + else: + current_time += frame_step + + comp.start_next_frame(current_time) + + modified_tensor = out_top1_tensor.float() + modified_tensor[:, :, :-1] *= .5 + modified_tensor[:, :, :-1] += 100 + modified_tensor = torch.clamp(modified_tensor, 0, 255) + modified_tensor = modified_tensor.to(torch.uint8) + comp.in_tops[0].from_tensor(modified_tensor) + + filtered_tensor = comp.imag_filter(out_top2_tensor.unsqueeze(0)).squeeze(0) + comp.in_tops[1].from_tensor(filtered_tensor) + + rgba_tensor = torch.cat((out_top3_tensor, comp.alpha_tensor), dim=0) + comp.in_tops[2].from_tensor(rgba_tensor) + + comp.frame_count += 1 + + # throttle host loop to HOST_FPS + loop_elapsed = time.perf_counter() - loop_start + sleep_time = host_period - loop_elapsed + if sleep_time > 0: + time.sleep(sleep_time) + + +comp.clear_on_layout_change_callback() +comp.stop() +comp.unload() +del out_top1_tensor, out_top2_tensor, out_top3_tensor +del modified_tensor, filtered_tensor, rgba_tensor +del comp +import gc; gc.collect() # noqa: E702 diff --git a/source/comp.cpp b/source/comp.cpp index 46c7292..526cc26 100644 --- a/source/comp.cpp +++ b/source/comp.cpp @@ -213,21 +213,25 @@ bool Comp::load() throw std::runtime_error("Failed to open tox file"); } - int64_t numerator = static_cast(time_.rate * 1000); - - TEResult result = TEInstanceSetFrameRate(instance_, numerator, 1000); + TEResult result = TEInstanceSetFloatFrameRate(instance_, static_cast(time_.rate)); if (result != TEResultSuccess) { spdlog::error("Failed to set frame rate: {}", TEResultGetDescription(result)); throw std::runtime_error("Failed to set frame rate"); } + time_.scale = static_cast(time_.rate) * 100; spdlog::info("Loading tox: {}", filePath_); auto timeMode = TETimeInternal; - if (compFlags_ & CompFlagBits::ExternalTime) timeMode = TETimeExternal; + auto uiMode = TEUIWindows; + if (compFlags_ & CompFlagBits::ExternalTime) + { + timeMode = TETimeExternal; + uiMode = TEUINone; + } - result = TEInstanceConfigure(instance_, filePath_.c_str(), timeMode, TEUIWindows); + result = TEInstanceConfigure(instance_, filePath_.c_str(), timeMode, uiMode); if (result != TEResultSuccess) { spdlog::error("Failed to configure TEInstance: {}", TEResultGetDescription(result)); @@ -400,8 +404,8 @@ Comp::onEventInstanceDidUnload(TEResult result, Comp* comp) spdlog::default_logger()->flush(); } -void -Comp::onEventFrameDidFinish(TEResult result, int64_t start_time_value, int32_t start_time_scale, +void +Comp::onEventFrameDidFinish(TEResult result, int64_t start_time_value, int32_t start_time_scale, int64_t end_time_value, int32_t end_time_scale, Comp* comp) { //if (result == TEResultSuccess && start_time_value >= 0) @@ -417,7 +421,7 @@ Comp::onEventFrameDidFinish(TEResult result, int64_t start_time_value, int32_t s else { - // need go through all possible results and handle them accordingly... + // need go through all possible results and handle them accordingly... auto severity = TEResultGetSeverity(result); if (severity == TESeverityWarning) @@ -994,7 +998,7 @@ Comp::asyncUpdate() } } -bool +bool Comp::frameDidFinish() { auto state = getState(); @@ -1005,12 +1009,24 @@ Comp::frameDidFinish() applyLayoutChange(); return false; } - return !state.inFrame; + + if (!state.inFrame) + { + if (compFlags_ & CompFlagBits::ExternalTime) applyValueChanges(); + return true; + } + return false; } bool Comp::startNextFrame() { + if (compFlags_ & CompFlagBits::ExternalTime) + { + double nextTime = time_.seconds + (100.0 / time_.scale); + return startNextFrame(nextTime); + } + setInFrame(true); TEResult result = TEInstanceStartFrameAtTime(instance_, 0, 0, false); if (result != TEResultSuccess) @@ -1054,7 +1070,7 @@ Comp::startNextFrame(int64_t timeValue, int32_t timeScale) time_.scale = timeScale; time_.frame = seconds / time_.rate; - TEResult result = TEInstanceStartFrameAtTime(instance_, timeValue, timeScale, false); + TEResult result = TEInstanceStartFrameAtTime(instance_, timeValue, timeScale, discontinuity); if (result != TEResultSuccess) { spdlog::error("Frame did not start successfully: {}", TEResultGetDescription(result)); From cc15229081b7a90e7c7f3d2ed642f64a35464221 Mon Sep 17 00:00:00 2001 From: keithlostracco Date: Wed, 27 May 2026 12:51:08 -0700 Subject: [PATCH 2/4] remove promote basics examples to examples directory --- examples/{basics => }/01_load_comp_basic.py | 0 examples/{basics => }/02_load_comp.py | 0 examples/{basics => }/03_set_pars.py | 0 examples/{basics => }/04_chops.py | 0 examples/{basics => }/05_dats.py | 0 examples/{basics => }/06_tops.py | 0 examples/{basics => }/07_on_loaded.py | 0 examples/{basics => }/08_on_loaded_async.py | 0 examples/{basics => }/09_concurrent_comps.py | 0 examples/{basics => }/10_external_time.py | 0 examples/{basics => }/ExampleComps.toe | Bin examples/{basics => }/ExampleReceive.toe | Bin examples/{basics => }/README.md | 0 examples/{basics => }/TopChopDatIO.tox | Bin examples/{basics => }/example_sync_dat.txt | 0 examples/{basics => }/modules/image_filter.py | 0 examples/{basics => }/modules/utils.py | 0 17 files changed, 0 insertions(+), 0 deletions(-) rename examples/{basics => }/01_load_comp_basic.py (100%) rename examples/{basics => }/02_load_comp.py (100%) rename examples/{basics => }/03_set_pars.py (100%) rename examples/{basics => }/04_chops.py (100%) rename examples/{basics => }/05_dats.py (100%) rename examples/{basics => }/06_tops.py (100%) rename examples/{basics => }/07_on_loaded.py (100%) rename examples/{basics => }/08_on_loaded_async.py (100%) rename examples/{basics => }/09_concurrent_comps.py (100%) rename examples/{basics => }/10_external_time.py (100%) rename examples/{basics => }/ExampleComps.toe (100%) rename examples/{basics => }/ExampleReceive.toe (100%) rename examples/{basics => }/README.md (100%) rename examples/{basics => }/TopChopDatIO.tox (100%) rename examples/{basics => }/example_sync_dat.txt (100%) rename examples/{basics => }/modules/image_filter.py (100%) rename examples/{basics => }/modules/utils.py (100%) diff --git a/examples/basics/01_load_comp_basic.py b/examples/01_load_comp_basic.py similarity index 100% rename from examples/basics/01_load_comp_basic.py rename to examples/01_load_comp_basic.py diff --git a/examples/basics/02_load_comp.py b/examples/02_load_comp.py similarity index 100% rename from examples/basics/02_load_comp.py rename to examples/02_load_comp.py diff --git a/examples/basics/03_set_pars.py b/examples/03_set_pars.py similarity index 100% rename from examples/basics/03_set_pars.py rename to examples/03_set_pars.py diff --git a/examples/basics/04_chops.py b/examples/04_chops.py similarity index 100% rename from examples/basics/04_chops.py rename to examples/04_chops.py diff --git a/examples/basics/05_dats.py b/examples/05_dats.py similarity index 100% rename from examples/basics/05_dats.py rename to examples/05_dats.py diff --git a/examples/basics/06_tops.py b/examples/06_tops.py similarity index 100% rename from examples/basics/06_tops.py rename to examples/06_tops.py diff --git a/examples/basics/07_on_loaded.py b/examples/07_on_loaded.py similarity index 100% rename from examples/basics/07_on_loaded.py rename to examples/07_on_loaded.py diff --git a/examples/basics/08_on_loaded_async.py b/examples/08_on_loaded_async.py similarity index 100% rename from examples/basics/08_on_loaded_async.py rename to examples/08_on_loaded_async.py diff --git a/examples/basics/09_concurrent_comps.py b/examples/09_concurrent_comps.py similarity index 100% rename from examples/basics/09_concurrent_comps.py rename to examples/09_concurrent_comps.py diff --git a/examples/basics/10_external_time.py b/examples/10_external_time.py similarity index 100% rename from examples/basics/10_external_time.py rename to examples/10_external_time.py diff --git a/examples/basics/ExampleComps.toe b/examples/ExampleComps.toe similarity index 100% rename from examples/basics/ExampleComps.toe rename to examples/ExampleComps.toe diff --git a/examples/basics/ExampleReceive.toe b/examples/ExampleReceive.toe similarity index 100% rename from examples/basics/ExampleReceive.toe rename to examples/ExampleReceive.toe diff --git a/examples/basics/README.md b/examples/README.md similarity index 100% rename from examples/basics/README.md rename to examples/README.md diff --git a/examples/basics/TopChopDatIO.tox b/examples/TopChopDatIO.tox similarity index 100% rename from examples/basics/TopChopDatIO.tox rename to examples/TopChopDatIO.tox diff --git a/examples/basics/example_sync_dat.txt b/examples/example_sync_dat.txt similarity index 100% rename from examples/basics/example_sync_dat.txt rename to examples/example_sync_dat.txt diff --git a/examples/basics/modules/image_filter.py b/examples/modules/image_filter.py similarity index 100% rename from examples/basics/modules/image_filter.py rename to examples/modules/image_filter.py diff --git a/examples/basics/modules/utils.py b/examples/modules/utils.py similarity index 100% rename from examples/basics/modules/utils.py rename to examples/modules/utils.py From 7260c43f4518acc16c8fc4b500b823eabc247a44 Mon Sep 17 00:00:00 2001 From: keithlostracco Date: Wed, 27 May 2026 14:16:51 -0700 Subject: [PATCH 3/4] Update Python version requirements and remove obsolete release process documentation --- README.md | 2 +- RELEASING.md | 45 --------------------------------------------- pyproject.toml | 1 + 3 files changed, 2 insertions(+), 46 deletions(-) delete mode 100644 RELEASING.md diff --git a/README.md b/README.md index d28258f..6326454 100644 --- a/README.md +++ b/README.md @@ -7,7 +7,7 @@ TouchPy is a high-performance Python toolset for working with TouchDesigner comp - Windows 10/11 (x64) - [TouchDesigner](https://derivative.ca/download) installed (runtime dependency) - NVIDIA GPU with CUDA support -- Python 3.9 - 3.13 +- Python 3.9 - 3.14 (3.12 recommended for development) ## Installation diff --git a/RELEASING.md b/RELEASING.md deleted file mode 100644 index 5eff8c7..0000000 --- a/RELEASING.md +++ /dev/null @@ -1,45 +0,0 @@ -# Release Process - -## Version Convention - -- During development, `pyproject.toml` has version `X.Y.Z.dev0` -- At release time, the `.dev0` suffix is stripped -- After release, version is bumped to the next `X.Y.Z.dev0` - -## Release Checklist - -1. **Prepare the release PR:** - - Update `version` in `pyproject.toml`: remove `.dev0` suffix (e.g., `0.12.0.dev0` → `0.12.0`) - - Update CHANGELOG.md with release notes - - Merge PR to `main` - -2. **Create the release:** - - Create a git tag: `git tag v0.12.0` - - Push the tag: `git push origin v0.12.0` - - Create a GitHub Release from the tag at https://github.com/IntentDev/touchpy-dev/releases/new - - This triggers the `wheels.yml` workflow which: - - Builds wheels for Python 3.9-3.13 (Windows x64) - - Publishes to PyPI via trusted publishing - - Uploads wheels as GitHub Release assets - - This also triggers the `docs.yml` workflow which rebuilds and deploys documentation - -3. **Post-release:** - - Create a follow-up PR bumping version to next `.dev0` (e.g., `0.13.0.dev0`) - - Merge to `main` - -## PyPI Trusted Publishing Setup (one-time) - -Configure at https://pypi.org under your account → Publishing: - -- **Owner**: `IntentDev` -- **Repository**: `touchpy-dev` (update after repo rename to `touchpy`) -- **Workflow name**: `wheels.yml` -- **Environment**: (leave blank) - -## Manual Workflow Triggers - -All workflows can be triggered manually from the GitHub Actions tab: - -- **wheels.yml**: Build wheels without publishing (manual dispatch skips PyPI upload) -- **docs.yml**: Rebuild and deploy documentation -- **ci.yml**: Runs automatically on push/PR diff --git a/pyproject.toml b/pyproject.toml index dc3f1ba..b30475e 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -18,6 +18,7 @@ classifiers = [ "Programming Language :: Python :: 3.11", "Programming Language :: Python :: 3.12", "Programming Language :: Python :: 3.13", + "Programming Language :: Python :: 3.14", "Operating System :: Microsoft :: Windows", ] From d42912b5426189e575ffcf617f164820bc3a6ab3 Mon Sep 17 00:00:00 2001 From: keithlostracco Date: Wed, 27 May 2026 14:34:28 -0700 Subject: [PATCH 4/4] Refactor: Update .gitignore, remove unused files, and clean up code structure - Updated .gitignore to exclude backup files and logs. - Deleted unused launch configuration file (launch.vs.json). - Removed obsolete argument parsing script (parse_args.bat). - Deleted component mask header and associated test files (componentmask.h, componentmask_gtest.cpp). - Removed texture copy implementation files (copyTeVkTextureToTeVkTexture.cpp, copykernels.cu, copykernels.cuh). - Deleted main touchpy implementation file (touchpy.cpp). - Removed vcpkg configuration file (vcpkg.json). --- .gitignore | 3 +- launch.vs.json | 17 -- parse_args.bat | 30 --- reference/componentmask.h | 154 ----------- reference/componentmask_gtest.cpp | 87 ------- reference/copyTeVkTextureToTeVkTexture.cpp | 201 -------------- reference/copykernels.cu | 164 ------------ reference/copykernels.cuh | 288 --------------------- reference/touchpy.cpp | 161 ------------ reference/vcpkg.json | 10 - 10 files changed, 2 insertions(+), 1113 deletions(-) delete mode 100644 launch.vs.json delete mode 100644 parse_args.bat delete mode 100644 reference/componentmask.h delete mode 100644 reference/componentmask_gtest.cpp delete mode 100644 reference/copyTeVkTextureToTeVkTexture.cpp delete mode 100644 reference/copykernels.cu delete mode 100644 reference/copykernels.cuh delete mode 100644 reference/touchpy.cpp delete mode 100644 reference/vcpkg.json diff --git a/.gitignore b/.gitignore index 45b69dc..ba37bed 100644 --- a/.gitignore +++ b/.gitignore @@ -42,7 +42,6 @@ /.vscode # Misc -*Backup *.dmp __pycache__ @@ -54,8 +53,10 @@ docs/source/reference/* CMakeUserPresets.json # TouchDesigner +*Backup/ *.*.toe */models +logs/ *log.txt # Bypass general rules for vendored SDK diff --git a/launch.vs.json b/launch.vs.json deleted file mode 100644 index 97f9765..0000000 --- a/launch.vs.json +++ /dev/null @@ -1,17 +0,0 @@ -{ - "version": "0.1.1", - "defaults": {}, - "configurations": [ - { - "type": "python", - "name": "Python: Run example_run_comp.py", - "projectTarget": "ExampleRunComp", - "projectName": "TouchPy", - "pythonPath": "${command:python.path}", - "program": "${workspaceRoot}/test/example_run_comp.py", - "cwd": "${workspaceRoot}/test", - "args": [], - "request": "launch" - } - ] -} \ No newline at end of file diff --git a/parse_args.bat b/parse_args.bat deleted file mode 100644 index 3fe523d..0000000 --- a/parse_args.bat +++ /dev/null @@ -1,30 +0,0 @@ -@echo off - -setlocal enabledelayedexpansion - -set "args=%*" - -:parse_args -if "%args%" == "" goto end_parse_args -for /f "tokens=1,2,* delims= " %%a in ("%args%") do ( - set "arg=%%a" - set "args=%%c" ; Update args to skip the current value - - if "!arg:~0,1!" == "-" ( - set "key=!arg:~1!" - set "value_!key!=%%b" - ) -) -goto parse_args - -:end_parse_args - -if defined value_n ( - echo -n: !value_n! -) - -if defined value_m ( - echo -m: !value_m! -) - -endlocal diff --git a/reference/componentmask.h b/reference/componentmask.h deleted file mode 100644 index 8325baf..0000000 --- a/reference/componentmask.h +++ /dev/null @@ -1,154 +0,0 @@ -#pragma once -#include -#include - -#include - - -enum class ComponentMask : uint8_t -{ - None = 0, - R = 1 << 0, - G = 1 << 1, - B = 1 << 2, - A = 1 << 3, - RG = R | G, - RGB = R | G | B, - RGBA = R | G | B | A -}; - -constexpr ComponentMask operator|(ComponentMask a, ComponentMask b) -{ - return static_cast( - static_cast>(a) | - static_cast>(b)); -} - -constexpr ComponentMask operator&(ComponentMask a, ComponentMask b) -{ - return static_cast( - static_cast>(a) & - static_cast>(b)); -} - -constexpr ComponentMask operator^(ComponentMask a, ComponentMask b) -{ - return static_cast( - static_cast>(a) ^ - static_cast>(b)); -} - -constexpr ComponentMask operator~(ComponentMask a) -{ - return static_cast(~static_cast>(a)); -} - -constexpr ComponentMask& operator|=(ComponentMask& a, ComponentMask b) -{ - a = a | b; - return a; -} - -constexpr ComponentMask& operator&=(ComponentMask& a, ComponentMask b) -{ - a = a & b; - return a; -} - -constexpr ComponentMask& operator^=(ComponentMask& a, ComponentMask b) -{ - a = a ^ b; - return a; -} - -constexpr bool operator!(ComponentMask a) -{ - return !static_cast>(a); -} - -constexpr bool operator==(ComponentMask a, ComponentMask b) -{ - return static_cast>(a) == static_cast>(b); -} - -constexpr bool operator!=(ComponentMask a, ComponentMask b) -{ - return static_cast>(a) != static_cast>(b); -} - -constexpr bool operator&&(ComponentMask a, ComponentMask b) -{ - return static_cast(a & b); -} - -constexpr bool operator||(ComponentMask a, ComponentMask b) -{ - return static_cast(a | b); -} - -namespace componentMask -{ - inline uint8_t numActiveComponents(ComponentMask flags) - { - uint8_t result = 0; - if (flags && ComponentMask::R) - { - ++result; - } - if (flags && ComponentMask::G) - { - ++result; - } - if (flags && ComponentMask::B) - { - ++result; - } - if (flags && ComponentMask::A) - { - ++result; - } - return result; - } - - inline std::string toString(ComponentMask flags) - { - std::string result; - if (flags == ComponentMask::None) - { - result = "None"; - } - else - { - if (flags && ComponentMask::R) - { - result += "Red"; - } - if (flags && ComponentMask::G) - { - if (!result.empty()) - { - result += " | "; - } - result += "Green"; - } - if (flags && ComponentMask::B) - { - if (!result.empty()) - { - result += " | "; - } - result += "Blue"; - } - if (flags && ComponentMask::A) - { - if (!result.empty()) - { - result += " | "; - } - result += "Alpha"; - } - } - return result; - } -} - diff --git a/reference/componentmask_gtest.cpp b/reference/componentmask_gtest.cpp deleted file mode 100644 index 7fdbabb..0000000 --- a/reference/componentmask_gtest.cpp +++ /dev/null @@ -1,87 +0,0 @@ -#include - -#include "componentmask.h" - -// Notes to use GTest: - -// Create a file with the same name as the header file you want to test, but with _gtest appended to the end -// include and the header file you want to test. -// -// Create a test suite with TEST(TestSuiteName, testName) { /* test code here */ } -// all tests in the same file should have the same TestSuiteName. -// -// Best practice is to have one test per function if possible (or at least one test per logical unit) unless the -// function is very simple. The test "operators" belows could be split into separate tests for each operator but -// in this case they are so simple that it's not really necessary. -// -// Run tests by selecting touchpygtest.exe as the startup project and pressing F5 -// existing tests will be run and the results will be displayed in the output window. -// -// Note: when making a new file, cmake has to rebuild the project before the new file will be included in the build -// the fastest way to do this in VS I've found is to have cmake file open and press ctrl-s to save it, -// then press F5 to build and run the tests. -// -// Also note: if you add a new test file, don't add it to the CMakeLists.text file (if popup appears asking if you -// want to reload, say no, it can be disable in options). CMake automatically includes all files in the project. - - -TEST(ComponentMask, operators) -{ - ComponentMask mask1 = ComponentMask::R; - ComponentMask mask2 = ComponentMask::G; - ComponentMask mask3 = ComponentMask::B; - - EXPECT_EQ(mask1 & mask2, ComponentMask::None); - EXPECT_EQ(mask1 | mask2 | mask3, ComponentMask::RGB); - EXPECT_EQ(mask1 ^ mask1, ComponentMask::None); - EXPECT_EQ(mask1 ^ mask2, ComponentMask::RG); - - EXPECT_NE(mask1, ~mask1); - - auto tmp = mask1; - tmp |= mask2; - EXPECT_EQ(mask1 | mask2, tmp); - - tmp = mask1; - tmp &= mask2; - EXPECT_EQ(mask1 & mask2, tmp); - - EXPECT_EQ(mask1 & mask1, mask1); - EXPECT_EQ(mask1 | mask1, mask1); - - EXPECT_EQ(mask1 & mask2, ComponentMask::None); - EXPECT_EQ(mask1 | mask2, ComponentMask::RG); - - tmp = ComponentMask::R; - EXPECT_TRUE(mask1 == tmp); - EXPECT_TRUE(mask1 != mask2); - EXPECT_FALSE(mask1 == mask2); - - - auto flags = ComponentMask::R | ComponentMask::G | ComponentMask::B; - EXPECT_TRUE(flags && ComponentMask::R); - EXPECT_TRUE(flags && ComponentMask::G); - EXPECT_TRUE(flags && ComponentMask::B); - EXPECT_FALSE(flags && ComponentMask::A); -} - -TEST(ComponentMask, numActiveComponents) -{ - auto flags = ComponentMask::R | ComponentMask::G; - - EXPECT_EQ(componentMask::numActiveComponents(flags), 2); - EXPECT_EQ(componentMask::numActiveComponents(ComponentMask::R), 1); - EXPECT_EQ(componentMask::numActiveComponents(ComponentMask::None), 0); -} - -TEST(ComponentMask, toString) -{ - EXPECT_EQ(componentMask::toString(ComponentMask::R), "Red"); - EXPECT_EQ(componentMask::toString(ComponentMask::G), "Green"); - EXPECT_EQ(componentMask::toString(ComponentMask::B), "Blue"); - EXPECT_EQ(componentMask::toString(ComponentMask::A), "Alpha"); - EXPECT_EQ(componentMask::toString(ComponentMask::RG), "Red | Green"); - EXPECT_EQ(componentMask::toString(ComponentMask::RGB), "Red | Green | Blue"); - EXPECT_EQ(componentMask::toString(ComponentMask::RGBA), "Red | Green | Blue | Alpha"); - EXPECT_EQ(componentMask::toString(ComponentMask::None), "None"); -} \ No newline at end of file diff --git a/reference/copyTeVkTextureToTeVkTexture.cpp b/reference/copyTeVkTextureToTeVkTexture.cpp deleted file mode 100644 index 19adbdc..0000000 --- a/reference/copyTeVkTextureToTeVkTexture.cpp +++ /dev/null @@ -1,201 +0,0 @@ -bool -Comp::applyOutputTextureChange() -{ - for (const auto& identifier : changedOutputTextures_) - { - TouchObject teTex; - TEResult result = TEInstanceLinkGetTextureValue( - instance_, identifier.c_str(), TELinkValueCurrent, teTex.take()); - - if (result == TEResultSuccess && TEInstanceHasTextureTransfer(instance_, teTex)) - { - - HANDLE handle = TEVulkanTextureGetHandle(static_cast(teTex.get())); - //std::cout << "Has Texture Transfer: " << identifier << ", Texture Handle: " << handle << std::endl; - - auto it = texturesExternal_.find(handle); - if (it == texturesExternal_.end()) - { - texturesExternal_[handle] = std::make_unique( - renderer_->vContext().physicalDevice, - renderer_->vContext().device, - instance_, - static_cast(teTex.get()) - ); - - if (texToTE_.get() == nullptr && texturesExternal_[handle]) - { - texToTE_ = std::make_unique( - physicalDevice_, - device_, - texturesExternal_[handle]->extent(), - texturesExternal_[handle]->format() - ); - } - - std::cout << "Texture created for handle: " << handle << std::endl; - - return true; // we need to wait for the texture to be ready - } - - auto texExternal = texturesExternal_[handle].get(); - - TouchObject teSemaphore; - //teSemaphore.set(texFromTE_->teVkSemaphore()); - teSemaphore.set(texExternal->teVkSemaphore()); - - uint64_t waitValue = 0; - result = TEInstanceGetTextureTransfer( - instance_, - teTex, - teSemaphore.take(), - &waitValue); - - - // use if using the texture after transfer is complete, such as for display but need - // use queue index that supports graphics... - //VkImageLayout srcLayout; - //VkImageLayout dstLayout; - - //result = TEInstanceGetVulkanTextureTransfer( - // instance_, - // teTex, - // &srcLayout, - // &dstLayout, - // teSemaphore.take(), - // &waitValue - //); - - //std::cout << "srcLayout: " << string_VkImageLayout(srcLayout) - // << " dstLayout: " << string_VkImageLayout(dstLayout) << std::endl; - - //std::cout << "TESemaphore: " << teSemaphore.get() << " waitValue: " << waitValue << std::endl; - - if (result == TEResultSuccess) - { - //std::cout << "Texture transfer: " << identifier << " : " << waitValue << std::endl; - if (TESemaphoreGetType(teSemaphore) == TESemaphoreTypeVulkan) - { - - // wait for semaphore - VkSemaphoreWaitInfoKHR waitInfo = {}; - waitInfo.sType = VK_STRUCTURE_TYPE_SEMAPHORE_WAIT_INFO_KHR; - waitInfo.pNext = nullptr; - waitInfo.flags = 0; - waitInfo.semaphoreCount = 1; - //VkSemaphore importSemaphore = texFromTE_->semaphore(); - VkSemaphore importSemaphore = texExternal->semaphore(); - waitInfo.pSemaphores = &importSemaphore; - waitInfo.pValues = &waitValue; - - VK_CHECK(vkWaitSemaphores(device_, &waitInfo, UINT64_MAX)); - - // reset command buffer - VK_CHECK(vkResetCommandBuffer(commandBuffer_, 0)); - - // copy to texToTE_ - VkCommandBufferBeginInfo beginInfo = {}; - beginInfo.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO; - beginInfo.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT; - beginInfo.pInheritanceInfo = nullptr; - - - VK_CHECK(vkBeginCommandBuffer(commandBuffer_, &beginInfo)); - - //if (texFromTE_ && !srcInitialized_) - if (texExternal && texExternal->imageLayout() != VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL) - { - // transition to transfer src optimal - - //texFromTE_->cmdTransitionImageLayout( - texExternal->cmdTransitionImageLayout( - commandBuffer_, - VK_IMAGE_LAYOUT_UNDEFINED, - VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL - ); - - //srcInitialized_ = true; - //std::cout << "srcInitialized_" << std::endl; - } - - if (texToTE_->imageLayout() != VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL) - //if (texInternal && texInternal->imageLayout() != VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL) - { - // transition to transfer dst optimal - texToTE_->cmdTransitionImageLayout( - commandBuffer_, - VK_IMAGE_LAYOUT_UNDEFINED, - VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL - ); - //dstInitialized_ = true; - //std::cout << "dstInitialized_" << std::endl; - } - - VkImageSubresourceLayers subresourceLayers = {}; - subresourceLayers.aspectMask = VK_IMAGE_ASPECT_COLOR_BIT; - subresourceLayers.mipLevel = 0; - subresourceLayers.baseArrayLayer = 0; - subresourceLayers.layerCount = 1; - - VkImageCopy imageCopy = {}; - imageCopy.srcSubresource = subresourceLayers; - imageCopy.dstSubresource = subresourceLayers; - imageCopy.extent = { texToTE_->extent().width, texToTE_->extent().height, 1 }; - - - vkCmdCopyImage( - commandBuffer_, - //texFromTE_->image(), - texExternal->image(), - VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL, - texToTE_->image(), - //texInternal->image(), - VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL, - 1, - &imageCopy - ); - - VK_CHECK(vkEndCommandBuffer(commandBuffer_)); - - uint64_t signalValue; - VK_CHECK(vkGetSemaphoreCounterValue(device_, texToTE_->semaphore(), &signalValue)); - texToTE_->setSignalValue(++signalValue); - //texInternal->setSignalValue(++signalValue); - - - VkTimelineSemaphoreSubmitInfo timelineSemaphoreSubmitInfo = {}; - timelineSemaphoreSubmitInfo.sType = VK_STRUCTURE_TYPE_TIMELINE_SEMAPHORE_SUBMIT_INFO; - timelineSemaphoreSubmitInfo.pNext = nullptr; - timelineSemaphoreSubmitInfo.waitSemaphoreValueCount = 1; - timelineSemaphoreSubmitInfo.pWaitSemaphoreValues = &waitValue; - timelineSemaphoreSubmitInfo.signalSemaphoreValueCount = 1; - timelineSemaphoreSubmitInfo.pSignalSemaphoreValues = &signalValue; - - VkSubmitInfo submitInfo = {}; - submitInfo.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; - submitInfo.pNext = &timelineSemaphoreSubmitInfo; - submitInfo.commandBufferCount = 1; - submitInfo.pCommandBuffers = &commandBuffer_; - - submitInfo.waitSemaphoreCount = 1; - submitInfo.pWaitSemaphores = &importSemaphore; - VkPipelineStageFlags waitStages[] = { VK_PIPELINE_STAGE_TRANSFER_BIT }; - submitInfo.pWaitDstStageMask = waitStages; - - // - submitInfo.signalSemaphoreCount = 1; - VkSemaphore signalSemaphores[] = { texToTE_->semaphore() }; - //VkSemaphore signalSemaphores[] = { texInternal->semaphore() }; - submitInfo.pSignalSemaphores = signalSemaphores; - - VK_CHECK(vkQueueSubmit(queue_, 1, &submitInfo, submitFence_)); - VK_CHECK(vkWaitForFences(device_, 1, &submitFence_, VK_TRUE, UINT64_MAX)); - VK_CHECK(vkResetFences(device_, 1, &submitFence_)); - - std::cout << " copied texture "; - - } - } - } - } -} \ No newline at end of file diff --git a/reference/copykernels.cu b/reference/copykernels.cu deleted file mode 100644 index 228d117..0000000 --- a/reference/copykernels.cu +++ /dev/null @@ -1,164 +0,0 @@ - -#include "copykernels.cuh" - -cudaError_t -memCopyBRGA8USurfaceToRGBA8U(void* dst, int width, int height, cudaSurfaceObject_t src, cudaStream_t stream) -{ - dim3 blockSize(16, 16, 1); - dim3 gridSize(divUp(width, blockSize.x), divUp(height, blockSize.y), 1); - fromSurfaceBRGA8UToRGBA8U <<>> (dst, width, height, src); - - CHECK_CUDA_ERROR_AND_RETURN_STATUS(cudaDeviceSynchronize()); -} - -cudaError_t -memCopyBRGA8USurfaceToPlanarRGBA8U(void* dst, int width, int height, cudaSurfaceObject_t src, cudaStream_t stream) -{ - dim3 blockSize(16, 16, 1); - dim3 gridSize(divUp(width, blockSize.x), divUp(height, blockSize.y), 1); - fromSurfaceBRGA8UToPlanarRGBA8U <<>> (dst, width, height, src); - - CHECK_CUDA_ERROR_AND_RETURN_STATUS(cudaDeviceSynchronize()); -} - -template cudaError_t -memCopyFromSurface(void* dst, int width, int height, cudaSurfaceObject_t src, cudaStream_t stream) -{ - dim3 blockSize(16, 16, 1); - dim3 gridSize(divUp(width, blockSize.x), divUp(height, blockSize.y), 1); - fromSurface <<>> (dst, width, height, src); - - CHECK_CUDA_ERROR_AND_RETURN_STATUS(cudaDeviceSynchronize()); -} - -// instantiate the template for the types we need, so the compiler can generate the code. -// must ensure that the template definitions are visible to any translation unit that instantiates those templates -// (i.e. they must be in a header file, or included in the translation unit before the instantiation but since this is -// cuda code, we can't include the .cu or .cuh file in the .h file) -template cudaError_t -memCopyFromSurface(void* dst, int width, int height, cudaSurfaceObject_t src, cudaStream_t stream); - -template cudaError_t -memCopyFromSurface(void* dst, int width, int height, cudaSurfaceObject_t src, cudaStream_t stream); - -template cudaError_t -memCopyFromSurface(void* dst, int width, int height, cudaSurfaceObject_t src, cudaStream_t stream); - -template cudaError_t -memCopyFromSurface(void* dst, int width, int height, cudaSurfaceObject_t src, cudaStream_t stream); - - -template cudaError_t -memCopyFromSurfaceToPlanar(void* dst, int width, int height, cudaSurfaceObject_t src, cudaStream_t stream) -{ - dim3 blockSize(16, 16, 1); - dim3 gridSize(divUp(width, blockSize.x), divUp(height, blockSize.y), 1); - fromSurfaceToPlanar << > > (dst, width, height, src); - - CHECK_CUDA_ERROR_AND_RETURN_STATUS(cudaDeviceSynchronize()); -} - -template cudaError_t -memCopyFromSurfaceToPlanar(void* dst, int width, int height, cudaSurfaceObject_t src, cudaStream_t stream); - -template cudaError_t -memCopyFromSurfaceToPlanar(void* dst, int width, int height, cudaSurfaceObject_t src, cudaStream_t stream); - -template cudaError_t -memCopyFromSurfaceToPlanar(void* dst, int width, int height, cudaSurfaceObject_t src, cudaStream_t stream); - -cudaError_t -memCopyRGBA8UToBGRA8USurface(cudaSurfaceObject_t dst, int width, int height, const void* src, cudaStream_t stream) -{ - dim3 blockSize(16, 16, 1); - dim3 gridSize(divUp(width, blockSize.x), divUp(height, blockSize.y), 1); - toSurfaceBRGA8UFromRGBA8U <<>> (dst, width, height, src); - - CHECK_CUDA_ERROR_AND_RETURN_STATUS(cudaDeviceSynchronize()); -} - -cudaError_t -memCopyPlanarRGBA8UToBGRA8USurface(cudaSurfaceObject_t dst, int width, int height, const void* src, cudaStream_t stream) -{ - dim3 blockSize(16, 16, 1); - dim3 gridSize(divUp(width, blockSize.x), divUp(height, blockSize.y), 1); - toSurfaceBRGA8UFromPlanarRGBA8U << > > (dst, width, height, src); - - CHECK_CUDA_ERROR_AND_RETURN_STATUS(cudaDeviceSynchronize()); -} - -cudaError_t -memCopyRGB8UToBGRA8USurface(cudaSurfaceObject_t dst, int width, int height, const void* src, cudaStream_t stream) -{ - dim3 blockSize(16, 16, 1); - dim3 gridSize(divUp(width, blockSize.x), divUp(height, blockSize.y), 1); - toSurfaceBRGA8UFromRGB8U <<>> (dst, width, height, src); - - CHECK_CUDA_ERROR_AND_RETURN_STATUS(cudaDeviceSynchronize()); -} - -cudaError_t -memCopyPlanarRGB8UToBGRA8USurface(cudaSurfaceObject_t dst, int width, int height, const void* src, cudaStream_t stream) -{ - dim3 blockSize(16, 16, 1); - dim3 gridSize(divUp(width, blockSize.x), divUp(height, blockSize.y), 1); - toSurfaceBRGA8UFromPlanarRGB8U <<>> (dst, width, height, src); - - CHECK_CUDA_ERROR_AND_RETURN_STATUS(cudaDeviceSynchronize()); -} - -template cudaError_t -memCopyToSurface(cudaSurfaceObject_t dst, int width, int height, const void* src, cudaStream_t stream) -{ - dim3 blockSize(16, 16, 1); - dim3 gridSize(divUp(width, blockSize.x), divUp(height, blockSize.y), 1); - toSurface <<>> (dst, width, height, src); - - CHECK_CUDA_ERROR_AND_RETURN_STATUS(cudaDeviceSynchronize()); -} - -// instantiate the template for the types we need, so the compiler can generate -// the code (needed because the template is in a .cu file) -template cudaError_t -memCopyToSurface(cudaSurfaceObject_t dst, int width, int height, const void* src, cudaStream_t stream); - -template cudaError_t -memCopyToSurface(cudaSurfaceObject_t dst, int width, int height, const void* src, cudaStream_t stream); - -template cudaError_t -memCopyToSurface(cudaSurfaceObject_t dst, int width, int height, const void* src, cudaStream_t stream); - -template cudaError_t -memCopyToSurface(cudaSurfaceObject_t dst, int width, int height, const void* src, cudaStream_t stream); - -template cudaError_t -memCopyToSurface(cudaSurfaceObject_t dst, int width, int height, const void* src, cudaStream_t stream) -{ - dim3 blockSize(16, 16, 1); - dim3 gridSize(divUp(width, blockSize.x), divUp(height, blockSize.y), 1); - toSurface << > > (dst, width, height, src); - - CHECK_CUDA_ERROR_AND_RETURN_STATUS(cudaDeviceSynchronize()); -} - -template cudaError_t -memCopyToSurface(cudaSurfaceObject_t dst, int width, int height, const void* src, cudaStream_t stream); - - - - -template cudaError_t -memCopyToSurface2(cudaSurfaceObject_t dst, int width, int height, const void* src, cudaStream_t stream) -{ - dim3 blockSize(16, 16, 1); - dim3 gridSize(divUp(width, blockSize.x), divUp(height, blockSize.y), 1); - toSurface2 << > > (dst, width, height, src); - - CHECK_CUDA_ERROR_AND_RETURN_STATUS(cudaDeviceSynchronize()); -} - -template cudaError_t -memCopyToSurface2(cudaSurfaceObject_t dst, int width, int height, const void* src, cudaStream_t stream); - -template cudaError_t -memCopyToSurface2(cudaSurfaceObject_t dst, int width, int height, const void* src, cudaStream_t stream); \ No newline at end of file diff --git a/reference/copykernels.cuh b/reference/copykernels.cuh deleted file mode 100644 index 5bd68ab..0000000 --- a/reference/copykernels.cuh +++ /dev/null @@ -1,288 +0,0 @@ -#pragma once - -#include "cuda_runtime.h" -#include "device_launch_parameters.h" -#include -#include - -#ifdef _DEBUG -#define CHECK_CUDA_ERROR_AND_RETURN_STATUS(call) do { \ - cudaError_t cudaStatus = (call); \ - if (cudaStatus != cudaSuccess) { \ - fprintf(stderr, "CUDA Error at %s:%d: %s\n", __FILE__, __LINE__, cudaGetErrorString(cudaStatus)); \ - } \ - return cudaStatus; \ -} while(0) -#else -#define CHECK_CUDA_ERROR_AND_RETURN_STATUS(call) return cudaSuccess; -#endif - -inline int divUp(int a, int b) -{ - return ((a % b) != 0) ? (a / b + 1) : (a / b); -} - -__global__ void -fromSurfaceBRGA8UToRGBA8U(void* dst, int width, int height, cudaSurfaceObject_t src) -{ - unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; - unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; - - if (x >= width || y >= height) - return; - - uchar4 color; - size_t size = sizeof(uchar4); - surf2Dread(&color, src, x * size, y, cudaBoundaryModeZero); - - uint8_t* dstPtr = (uint8_t*)dst + x * size + y * width * size; - dstPtr[0] = color.z; - dstPtr[1] = color.y; - dstPtr[2] = color.x; - dstPtr[3] = color.w; -} - -__global__ void -fromSurfaceBRGA8UToPlanarRGBA8U(void* dst, int width, int height, cudaSurfaceObject_t src) -{ - unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; - unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; - - if (x >= width || y >= height) - return; - - uchar4 color; - size_t size = sizeof(uchar4); - surf2Dread(&color, src, x * size, y, cudaBoundaryModeZero); - - size_t stride = width * height; - - uint8_t* dstPtr = (uint8_t*)dst; - - //uint32_t row = (height - y - 1) * width; - uint32_t row = y * width; - - dstPtr[x + row] = color.z; - dstPtr[x + row + stride] = color.y; - dstPtr[x + row + 2 * stride] = color.x; - dstPtr[x + row + 3 * stride] = color.w; - -} - -__global__ void -fromSurfaceBRGA8UToRGB8U(void* dst, int width, int height, cudaSurfaceObject_t src) -{ - unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; - unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; - - if (x >= width || y >= height) - return; - - uchar4 color; - surf2Dread(&color, src, x * sizeof(uchar4), y, cudaBoundaryModeZero); - - size_t size = sizeof(uchar3); - uint8_t* dstPtr = (uint8_t*)dst + x * size + y * width * size; - dstPtr[0] = color.z; - dstPtr[1] = color.y; - dstPtr[2] = color.x; -} - -__global__ void -fromSurfaceBRGA8UToPlanarRGB8U(void* dst, int width, int height, cudaSurfaceObject_t src) -{ - unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; - unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; - - if (x >= width || y >= height) - return; - - uchar4 color; - surf2Dread(&color, src, x * sizeof(uchar4), y, cudaBoundaryModeZero); - - size_t stride = width * height; - - uint8_t* dstPtr = (uint8_t*)dst; - - dstPtr[x + y * width] = color.z; - dstPtr[x + y * width + stride] = color.y; - dstPtr[x + y * width + 2 * stride] = color.x; -} - -// interleave stride RGBA, RGBA, RGBA, ... -// RGBAFS32, RGFS32, RFS32, RGBAFS16, RGFS16, RFS16, RGU8, RU8 -template __global__ void -fromSurface(void* dst, int width, int height, cudaSurfaceObject_t src) -{ - unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; - unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; - - if (x >= width || y >= height) - return; - - T color; - size_t size = sizeof(T); - surf2Dread(&color, src, x * size, y, cudaBoundaryModeZero); - T* dstPtr = (T*)((T*)dst + y * width); - dstPtr[x] = color; -} - -// Planar stride RRRR..., GGGG..., BBBB... -// RGBAFS32, RGFS32, RFS32, RGBAFS16, RGFS16, RFS16, RGU8, RU8 -template __global__ void -fromSurfaceToPlanar(void* dst, int width, int height, cudaSurfaceObject_t src) -{ - unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; - unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; - - if (x >= width || y >= height) - return; - - T color; - size_t size = sizeof(T); - surf2Dread(&color, src, x * size, y, cudaBoundaryModeZero); - - size_t stride = width * height; - - for (size_t i = 0; i < numComps; i++) - { - CompType* dstPtr = (CompType*)((CompType*)dst + i * stride); - dstPtr[x + y * width] = ((CompType*)&color)[i]; - } -} - -__global__ void -toSurfaceBRGA8UFromRGBA8U(cudaSurfaceObject_t dst, int width, int height, const void* src) -{ - unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; - unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; - - if (x >= width || y >= height) - return; - - size_t size = sizeof(uchar4); - uint8_t* srcPtr = (uint8_t*)src + x * size + y * width * size; - uchar4 color; - color.z = srcPtr[0]; - color.y = srcPtr[1]; - color.x = srcPtr[2]; - color.w = srcPtr[3]; - surf2Dwrite(color, dst, x * size, y, cudaBoundaryModeZero); -} - -__global__ void -toSurfaceBRGA8UFromPlanarRGBA8U(cudaSurfaceObject_t dst, int width, int height, const void* src) -{ - unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; - unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; - - if (x >= width || y >= height) - return; - - size_t stride = width * height; - - uint8_t* srcPtr = (uint8_t*)src; - uchar4 color; - color.z = srcPtr[x + y * width]; - color.y = srcPtr[x + y * width + stride]; - color.x = srcPtr[x + y * width + 2 * stride]; - color.w = srcPtr[x + y * width + 3 * stride]; - surf2Dwrite(color, dst, x * sizeof(uchar4), y, cudaBoundaryModeZero); -} - - -__global__ void -toSurfaceBRGA8UFromRGB8U(cudaSurfaceObject_t dst, int width, int height, const void* src) -{ - unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; - unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; - - if (x >= width || y >= height) - return; - - size_t size = sizeof(uchar3); - uint8_t* srcPtr = (uint8_t*)src + x * size + y * width * size; - uchar4 color; - color.z = srcPtr[0]; - color.y = srcPtr[1]; - color.x = srcPtr[2]; - color.w = 255; - surf2Dwrite(color, dst, x * sizeof(uchar4), y, cudaBoundaryModeZero); -} - -__global__ void -toSurfaceBRGA8UFromPlanarRGB8U(cudaSurfaceObject_t dst, int width, int height, const void* src) -{ - unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; - unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; - - if (x >= width || y >= height) - return; - - size_t stride = width * height; - - uint8_t* srcPtr = (uint8_t*)src; - uchar4 color; - color.z = srcPtr[x + y * width]; - color.y = srcPtr[x + y * width + stride]; - color.x = srcPtr[x + y * width + 2 * stride]; - color.w = 255; - surf2Dwrite(color, dst, x * sizeof(uchar4), y, cudaBoundaryModeZero); -} - -// RGBAFS32, RGFS32, RFS32, RGBAFS16, RGFS16, RFS16, RGU8, RU8 -template __global__ void -toSurface(cudaSurfaceObject_t dst, int width, int height, const void* src) -{ - unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; - unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; - - if (x >= width || y >= height) - return; - - size_t size = sizeof(T); - T color = *(T*)((T*)src + x + y * width); - surf2Dwrite(color, dst, x * size, y, cudaBoundaryModeZero); -} - -// RGBAFS32, RGFS32, RFS32, RGBAFS16, RGFS16, RFS16, RGU8, RU8 -template __global__ void -toSurface(cudaSurfaceObject_t dst, int width, int height, const void* src) -{ - unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; - unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; - - if (x >= width || y >= height) - return; - - size_t size = sizeof(T); - size_t numComps = size / sizeof(CompType); - size_t stride = width * height; - - T color; - for (size_t i = 0; i < numComps; i++) - { - CompType* srcPtr = (CompType*)((CompType*)src + i * stride); - ((CompType*)&color)[i] = srcPtr[x + y * width]; - } - - surf2Dwrite(color, dst, x * size, y, cudaBoundaryModeZero); -} - -template __global__ void -toSurface2(cudaSurfaceObject_t dst, int width, int height, const void* src) -{ - unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; - unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; - - if (x >= width || y >= height) - return; - - SrcT srcColor = *(SrcT*)((SrcT*)src + x + y * width); - DstT color; - color.x = srcColor.x; - color.y = srcColor.y; - color.z = srcColor.z; - color.w = 255; - surf2Dwrite(color, dst, x * sizeof(DstT), y, cudaBoundaryModeZero); -} \ No newline at end of file diff --git a/reference/touchpy.cpp b/reference/touchpy.cpp deleted file mode 100644 index c7df078..0000000 --- a/reference/touchpy.cpp +++ /dev/null @@ -1,161 +0,0 @@ -#include -#include -#include -#include -#include - - -namespace nb = nanobind; - -// Buffer that writes to Python instead of C++ -class pythonbuf : public std::streambuf { -private: - using traits_type = std::streambuf::traits_type; - - size_t buf_size; - std::unique_ptr d_buffer; - nb::object pywrite; - nb::object pyflush; - mutable std::mutex buffer_mutex; - - int overflow(int c) override { - //std::lock_guard lock(buffer_mutex); - nb::gil_scoped_acquire acquire; - if (!traits_type::eq_int_type(c, traits_type::eof())) { - *pptr() = traits_type::to_char_type(c); - pbump(1); - } - return sync() == 0 ? traits_type::not_eof(c) : traits_type::eof(); - } - - // Computes how many bytes at the end of the buffer are part of an - // incomplete sequence of UTF-8 bytes. - // Precondition: pbase() < pptr() - size_t utf8_remainder() const { - const auto rbase = std::reverse_iterator(pbase()); - const auto rpptr = std::reverse_iterator(pptr()); - auto is_ascii = [](char c) { return (static_cast(c) & 0x80) == 0x00; }; - auto is_leading = [](char c) { return (static_cast(c) & 0xC0) == 0xC0; }; - auto is_leading_2b = [](char c) { return static_cast(c) <= 0xDF; }; - auto is_leading_3b = [](char c) { return static_cast(c) <= 0xEF; }; - // If the last character is ASCII, there are no incomplete code points - if (is_ascii(*rpptr)) { - return 0; - } - // Otherwise, work back from the end of the buffer and find the first - // UTF-8 leading byte - const auto rpend = rbase - rpptr >= 3 ? rpptr + 3 : rbase; - const auto leading = std::find_if(rpptr, rpend, is_leading); - if (leading == rbase) { - return 0; - } - const auto dist = static_cast(leading - rpptr); - size_t remainder = 0; - - if (dist == 0) { - remainder = 1; // 1-byte code point is impossible - } - else if (dist == 1) { - remainder = is_leading_2b(*leading) ? 0 : dist + 1; - } - else if (dist == 2) { - remainder = is_leading_3b(*leading) ? 0 : dist + 1; - } - // else if (dist >= 3), at least 4 bytes before encountering an UTF-8 - // leading byte, either no remainder or invalid UTF-8. - // Invalid UTF-8 will cause an exception later when converting - // to a Python string, so that's not handled here. - return remainder; - } - - // This function must be non-virtual to be called in a destructor. - int _sync() { - //std::lock_guard lock(buffer_mutex); - nb::gil_scoped_acquire acquire; - if (pbase() != pptr()) { - - // If buffer is not empty - //nb::gil_scoped_acquire tmp; - // This subtraction cannot be negative, so dropping the sign. - auto size = static_cast(pptr() - pbase()); - size_t remainder = utf8_remainder(); - - if (size > remainder) { - nb::str line(pbase(), size - remainder); - pywrite(std::move(line)); - pyflush(); - } - - // Copy the remainder at the end of the buffer to the beginning: - if (remainder > 0) { - std::memmove(pbase(), pptr() - remainder, remainder); - } - setp(pbase(), epptr()); - pbump(static_cast(remainder)); - } - return 0; - } - - int sync() override { return _sync(); } - -public: - explicit pythonbuf(const nb::object& pyostream, size_t buffer_size = 1024) - : buf_size(buffer_size), d_buffer(new char[buf_size]), pywrite(pyostream.attr("write")), - pyflush(pyostream.attr("flush")) { - setp(d_buffer.get(), d_buffer.get() + buf_size - 1); - } - - pythonbuf(pythonbuf&&) = default; - - /// Sync before destroy - ~pythonbuf() override { _sync(); } -}; - -class scoped_ostream_redirect -{ -protected: - std::streambuf* old; - std::ostream& costream; - pythonbuf buffer; - -public: - explicit scoped_ostream_redirect( - std::ostream& costream = std::cout, - const nb::object& pyostream = nb::module_::import_("sys").attr("stdout")) - noexcept - : costream(costream), buffer(pyostream) - { - old = costream.rdbuf(&buffer); - } - ~scoped_ostream_redirect() { - costream.rdbuf(old); - } - scoped_ostream_redirect(const scoped_ostream_redirect&) = delete; - scoped_ostream_redirect(scoped_ostream_redirect&& other) = default; - scoped_ostream_redirect& operator=(const scoped_ostream_redirect&) = delete; - scoped_ostream_redirect& operator=(scoped_ostream_redirect&&) = delete; -}; - - - - -extern void initCompBindings(nb::module_& m); -extern void initTopLinkBindings(nb::module_& m); -extern void initChopLinkBindings(nb::module_& m); -extern void initDatLinkBindings(nb::module_& m); -extern void initParLinkBindings(nb::module_& m); - -NB_MODULE(touchpy, m) -{ - //nb::set_leak_warnings(false); - - nb::class_(m, "ScopedOstreamRedirect") - .def(nb::init<>()); - - - initCompBindings(m); - initTopLinkBindings(m); - initChopLinkBindings(m); - initDatLinkBindings(m); - initParLinkBindings(m); -} \ No newline at end of file diff --git a/reference/vcpkg.json b/reference/vcpkg.json deleted file mode 100644 index 07e2d73..0000000 --- a/reference/vcpkg.json +++ /dev/null @@ -1,10 +0,0 @@ -{ - "dependencies": [ - "vulkan-memory-allocator", - "gtest", - { - "name": "spdlog", - "version>=": "1.13.0" - } - ] - } \ No newline at end of file