Skip to content
Open
Show file tree
Hide file tree
Changes from 32 commits
Commits
Show all changes
631 commits
Select commit Hold shift + click to select a range
30576a3
unit test to detect inconsistent access modes in stackable ctx
caugonnet Sep 9, 2025
eead229
Detect erroneous accesses
caugonnet Sep 9, 2025
d16ee51
check that virtual nodes are using data appropriately too
caugonnet Sep 9, 2025
6be49a7
fix programming error
caugonnet Sep 9, 2025
721cbd1
Add a new test that could simulate how we implement while nodes in st…
caugonnet Sep 11, 2025
e46997a
Define virtual methods to detect if we have a graph context, and star…
caugonnet Sep 11, 2025
579baf5
Extract the insert_graph in a dedicated header
caugonnet Sep 11, 2025
9276df9
clang-format
caugonnet Sep 11, 2025
2455266
example as a preparatory work to use reductions to create condition h…
caugonnet Sep 12, 2025
50b319a
define wait and reduce methods on stackable_logical_data (WIP)
caugonnet Sep 12, 2025
b59ceb3
Allow access_mode::reduce{_no_init} in stackable_ctx
caugonnet Sep 12, 2025
0a06b8b
Make sure that the reduce access mode in parallel_for works with grap…
caugonnet Sep 12, 2025
140a373
Bug fix ! there was a confusion between block and grid sizes
caugonnet Sep 12, 2025
a0da9e7
Do use a stackable_ctx in this test
caugonnet Sep 12, 2025
99d6bcc
add missing cuda_safe_call
caugonnet Sep 12, 2025
8c67453
Replace explicit allocations by calls to the uncached allocator
caugonnet Sep 12, 2025
d124172
Add missing file
caugonnet Sep 12, 2025
aab5241
- Use a different implementation of internal ctx nodes in stackable b…
caugonnet Sep 14, 2025
c4346c7
use push/pop rather than scope
caugonnet Sep 14, 2025
70f781e
Add const qualifiers and expose the backend ctx
caugonnet Sep 15, 2025
f1c9063
add missing header
caugonnet Sep 15, 2025
6ed5ca8
refactor for nested ctx again
caugonnet Sep 15, 2025
4922653
Move file to a better location
caugonnet Sep 15, 2025
f1c5537
remove an inappropriate assertion, and minor code cleanup
caugonnet Sep 15, 2025
a006410
First prototype of push_while
caugonnet Sep 15, 2025
08a2fba
more work with while loops
caugonnet Sep 15, 2025
8743bc1
Experiments with while loops
caugonnet Sep 15, 2025
389fff3
Merge branch 'main' into stackable_ctx_data
caugonnet Sep 15, 2025
9246f12
Merge branch 'main' into stackable_ctx_data
caugonnet Sep 16, 2025
ee8b6a7
Add tests to make sure task_deps are movable, copyable...
caugonnet Sep 16, 2025
8e4270d
clang-format
caugonnet Sep 16, 2025
9e24884
Save WIP with update cond helper
caugonnet Sep 16, 2025
41d8aa1
better implementation of update_cond
caugonnet Sep 16, 2025
1521ea5
change default value to 1, since a 0 value means the loop cannot exec…
caugonnet Sep 16, 2025
c72f88b
Merge branch 'main' into stackable_ctx_data
caugonnet Sep 17, 2025
333fe12
clang-format
caugonnet Sep 17, 2025
cad53e6
Make sure the push_while API is only enabled with CUDA 12.4 and above
caugonnet Sep 17, 2025
3d03c4a
Implement an RAII construct to perform loops
caugonnet Sep 17, 2025
9a54090
fix implementation of the repeat guard construct
caugonnet Sep 17, 2025
2689116
Add an FDTD example with repeat (temporary)
caugonnet Sep 17, 2025
1daadb3
cleanups and formatting
caugonnet Sep 17, 2025
35b45b8
Merge branch 'main' into stackable_ctx_data
caugonnet Sep 17, 2025
f08e92b
We cannot put extended lambdas in unit tests, and properly define whe…
caugonnet Sep 17, 2025
254cc07
misplaced ifdef
caugonnet Sep 17, 2025
c6d7569
fix unit tests
caugonnet Sep 17, 2025
9480518
add missing header
caugonnet Sep 17, 2025
c54b498
fix unused variable
caugonnet Sep 17, 2025
39b1f56
Merge branch 'main' into stackable_ctx_data
caugonnet Sep 17, 2025
3ab836d
fix unused warning
caugonnet Sep 17, 2025
114de09
fix unused warning again
caugonnet Sep 17, 2025
a81a269
remove experiment with extract graph
caugonnet Sep 17, 2025
48fa248
fix unused argc argv if waived
caugonnet Sep 17, 2025
7e1fc57
Comment to explain why it's complicated to do ctx.task in stackable_ctx
caugonnet Sep 17, 2025
1bf7825
clang-format
caugonnet Sep 17, 2025
950743a
- Move insert_graph utility
caugonnet Sep 18, 2025
8c0ac6c
Reimplement insert_graph_node which was erased
caugonnet Sep 18, 2025
2f5e496
Start to refactor how we implement nested contexts
caugonnet Sep 18, 2025
6db1fb2
do no try to use conditional nodes unless CUDA 12.4+
caugonnet Sep 18, 2025
7360828
Remove previous virtual_node approach
caugonnet Sep 18, 2025
855f7b8
defer the creation of child graphs until later
caugonnet Sep 18, 2025
1721570
Introduce a new test with nested while loops
caugonnet Sep 18, 2025
8d31202
Fix example with nested while (still WIP)
caugonnet Sep 18, 2025
963409b
remove remaining bits of virtual ctx nodes
caugonnet Sep 18, 2025
c0e7d23
Do not clone a graph to a child graph, directly fill the child graph
caugonnet Sep 18, 2025
8a888b6
Fix the implementation of conditional nodes to reset them after they …
caugonnet Sep 18, 2025
772f708
Merge branch 'main' into stackable_ctx_data
caugonnet Sep 18, 2025
68cfc74
Fix bounds
caugonnet Sep 19, 2025
a0c4ac6
Fix input dependencies when there are more than 1
caugonnet Sep 19, 2025
d6fbd18
Waive tests with conditional nodes before 12.4
caugonnet Sep 19, 2025
9ca5a09
Remove some tests which are broken for now
caugonnet Sep 19, 2025
2d89324
clang-format
caugonnet Sep 19, 2025
c6a0abf
waive pagerank test if cuda <12.4
caugonnet Sep 19, 2025
69d08d6
wait and fence can't be done in a nested context
caugonnet Sep 19, 2025
b997a89
Fix bug with missing __device__ annotations, and add a test where we …
caugonnet Sep 19, 2025
79678fe
fix unused variable warnings
caugonnet Sep 19, 2025
b5fb189
Stream adapters are not cleared directly when we have nested contexts…
caugonnet Sep 20, 2025
1ae04d4
Merge branch 'main' into stackable_ctx_data
caugonnet Sep 20, 2025
22c6072
Add a unit test that uses an host exec place on a task with stackable…
caugonnet Sep 21, 2025
809c449
Deal with the execution place in deferred tasks
caugonnet Sep 21, 2025
5220a03
Fix how we get the default exec place in task
caugonnet Sep 21, 2025
2848aa6
Merge branch 'main' into stackable_ctx_data
caugonnet Sep 21, 2025
fd3dca2
Remove graph_utilities.cuh and reimplement the freeze_while_graphs_2.…
caugonnet Sep 22, 2025
6a52ea7
Remove dead code
caugonnet Sep 22, 2025
47786f6
fix compilation errors
caugonnet Sep 22, 2025
f194afd
while and repeat guards need nvcc
caugonnet Sep 22, 2025
ae24701
Remove a reference on a const data
caugonnet Sep 22, 2025
d44a52b
remove a test on exception because we don't handle them
caugonnet Sep 22, 2025
f04dec4
Merge branch 'main' into stackable_ctx_data
caugonnet Sep 22, 2025
48d6336
Remove debug prints
caugonnet Sep 22, 2025
a3d0a8e
Rename defaultLaunchValue to default_launch_value
caugonnet Sep 22, 2025
a174d0b
Rely on finalize again to instantiate the root graph (and use a cache…
caugonnet Sep 22, 2025
748f7b3
Merge branch 'main' into stackable_ctx_data
caugonnet Sep 22, 2025
8b6c697
Fix the synchronization mechanism used in freeze_while_graph_2 test
caugonnet Sep 22, 2025
a284911
provide export/import resources in context
caugonnet Sep 22, 2025
b8fefdc
For non nested contexts, the finalize step did release resources, no …
caugonnet Sep 22, 2025
66e8d10
easier to verify test
caugonnet Sep 22, 2025
e4e957f
Merge branch 'main' into stackable_ctx_data
caugonnet Sep 22, 2025
b994059
minor test cleanup
caugonnet Sep 22, 2025
fb2387b
revert a previous change, we cannot use finalize directly because the…
caugonnet Sep 22, 2025
4a7c56d
Workaround broken tests on NVHPC
caugonnet Sep 22, 2025
4d030bc
Merge branch 'main' into stackable_ctx_data
caugonnet Sep 22, 2025
ebfab1a
Fixes and uniformization of jacobi examples (pfor and update_cond)
caugonnet Sep 23, 2025
366779f
Move stackable_ctx.cuh to a new directory
caugonnet Sep 23, 2025
27b7c35
Start to split stackable_ctx.cuh into multiple files
caugonnet Sep 23, 2025
c892602
clang-format
caugonnet Sep 23, 2025
24d725e
Extract stackable_ctx impl
caugonnet Sep 23, 2025
9eb9326
clang-format
caugonnet Sep 23, 2025
a112667
frozen_logical_data<T> now inherits from frozen_logical_data_untyped
caugonnet Sep 23, 2025
e58ad48
Implement a ctx.freeze method for untyped objects
caugonnet Sep 23, 2025
8e76ebd
clang-format
caugonnet Sep 23, 2025
10632c9
minor cleanups
caugonnet Sep 23, 2025
a8dd2fe
Allow extra link_libs for header test link checks, add driver for stf.
alliepiper Sep 23, 2025
f48bdfb
Merge branch 'stf_frozen_untyped' into stackable_ctx_data
caugonnet Sep 23, 2025
e644a3b
Add missing const in logical_data_untyped::read()
caugonnet Sep 23, 2025
41e3fc7
Implement stackable_logical_data_untyped
caugonnet Sep 23, 2025
0efbfea
overload set_symbol
caugonnet Sep 23, 2025
c4050bb
Revert "overload set_symbol"
caugonnet Sep 23, 2025
b834061
Revert "Implement stackable_logical_data_untyped"
caugonnet Sep 23, 2025
1c3c893
Automatically include stackable_ctx.cuh from stf.cuh
caugonnet Sep 24, 2025
80882c9
More realistic Sparse CG
caugonnet Sep 24, 2025
495ce68
Implement Sparse CG with a while node, this required fixing a bug in …
caugonnet Sep 24, 2025
5995840
clang-format
caugonnet Sep 24, 2025
4413581
Merge branch 'main' into stackable_ctx_data
caugonnet Sep 24, 2025
bdb5c50
Fix wrongly formed constructor
caugonnet Sep 24, 2025
e98c3be
remove useless header
caugonnet Sep 24, 2025
2ca36c3
Remove a test not needed
caugonnet Sep 24, 2025
99f1828
Fix headers
caugonnet Sep 24, 2025
3d0c958
Merge branch 'main' into stackable_ctx_data
caugonnet Sep 24, 2025
5a78826
Waive cg_csr_stackable test for CUDA < 12.4
caugonnet Sep 24, 2025
9f8d954
Merge branch 'main' into stackable_ctx_data
caugonnet Sep 24, 2025
2441e83
more protection against undefined features on CUDA < 12.4
caugonnet Sep 25, 2025
424af80
Select parameters for jacobi that should converge fast for the sake o…
caugonnet Sep 25, 2025
908c8de
Merge branch 'main' into stackable_ctx_data
caugonnet Sep 25, 2025
05443a2
Start to adapt a cuda sample to show how this is done with STF
caugonnet Sep 25, 2025
7a2fdb5
Adapt more of the cuda sample with conditional graphs
caugonnet Sep 25, 2025
eedccc3
extend the example with cuda_kernel constructs
caugonnet Sep 25, 2025
64cda5c
Burger equation
caugonnet Sep 25, 2025
cec337c
Reduce the amount of prints
caugonnet Sep 25, 2025
ca980b0
Add symbol
caugonnet Sep 26, 2025
17c920f
Add a script to cleanup the DOT files generated by CUDA graph API
caugonnet Sep 26, 2025
00979e6
apply cu++filt
caugonnet Sep 26, 2025
ec212a4
clang-format
caugonnet Sep 26, 2025
6381f2b
Save WIP with Burger solver
caugonnet Sep 26, 2025
e3c2313
Use a cudaGraph_t and not a shared_ptr to a cudaGraph_t to query if w…
caugonnet Sep 26, 2025
198de53
Try to use cached graphs
caugonnet Sep 26, 2025
8565b4a
Merge branch 'main' into stackable_ctx_data
caugonnet Sep 26, 2025
587f309
Only enable graphConditionalNodes example for CUDA >= 12.4
caugonnet Sep 30, 2025
723d914
Repeat the same nested context to ensure we have to reinstantiate graphs
caugonnet Sep 30, 2025
f954c14
Temporary work-around to avoid a CUDA graph bug
caugonnet Sep 30, 2025
e16d4b3
add ctx.launch for stackable_ctx
caugonnet Sep 30, 2025
4d92ec8
declare data where the need to be
caugonnet Oct 1, 2025
d134a44
Merge branch 'main' into stackable_ctx_data
caugonnet Oct 2, 2025
7be2ff6
start to restructure newton/burger example
caugonnet Oct 1, 2025
5f0cfd1
generic newton solver
caugonnet Oct 1, 2025
15acd5d
Pagerank batched example
caugonnet Oct 2, 2025
e3f3490
More general burger examples
caugonnet Oct 2, 2025
1da2239
CUDASTF_DEBUG_STACKABLE_DOT only prints the instantiable graph
caugonnet Oct 2, 2025
2914044
Improve examples
caugonnet Oct 7, 2025
8107630
instrument examples
caugonnet Oct 7, 2025
1a00d05
define ctx::logical_data_t<T>
caugonnet Oct 7, 2025
1335a2b
Merge branch 'main' into stackable_ctx_data
caugonnet Oct 7, 2025
2a20564
move CUDA 12.4+ code into ifdef guards
caugonnet Oct 7, 2025
55ab19d
clang-format
caugonnet Oct 7, 2025
e316290
take into account cuda13 changes in graph API
caugonnet Oct 7, 2025
1260861
Merge branch 'main' into stackable_ctx_data
caugonnet Oct 7, 2025
70fc9d6
get_ready_dependencies is dead code
caugonnet Oct 8, 2025
1b38ec6
Merge branch 'main' into stackable_ctx_data
caugonnet Oct 8, 2025
52097f0
Merge branch 'main' into stackable_ctx_data
caugonnet Feb 8, 2026
2ee8717
Merge branch 'main' into stackable_ctx_data
caugonnet Feb 8, 2026
d6ab5c2
clang-format
caugonnet Feb 8, 2026
d78c3e8
No longer use _CUDA_VSTD
caugonnet Feb 8, 2026
9764592
Fix a logic error with CTK filtering in pagerank_batched
caugonnet Feb 8, 2026
50e8dd9
remove unused captured variable
caugonnet Feb 9, 2026
3b80cd7
Merge branch 'main' into stackable_ctx_data
caugonnet Feb 9, 2026
b9bf978
Merge branch 'main' into stackable_ctx_data
caugonnet Feb 10, 2026
d79d621
API to import resources from a context rather than to import(mv)/export
caugonnet Feb 10, 2026
293f992
Remove unnecessary cmake change
caugonnet Feb 10, 2026
8e64583
Merge branch 'main' into stackable_ctx_data
caugonnet Feb 11, 2026
a54033b
Remove task input nodes which have never been used and which are effe…
caugonnet Feb 11, 2026
8e3f74c
Simpler code without debug leftovers
caugonnet Feb 11, 2026
ba83a97
fix a concurrency issue
caugonnet Feb 11, 2026
29dde3f
Remove unused variable
caugonnet Feb 11, 2026
11a30c6
Remove a useless variadic template parameter
caugonnet Feb 11, 2026
c6ef2e2
Merge branch 'main' into stackable_ctx_data
caugonnet Feb 11, 2026
f49cd29
Add some debug for stackable_ctx heads
caugonnet Feb 13, 2026
b4d57ee
Merge branch 'main' into stackable_ctx_data
caugonnet Feb 13, 2026
5227029
Merge branch 'main' into stackable_ctx_data
caugonnet Mar 15, 2026
43c13c6
Merge branch 'main' into stackable_ctx_data
caugonnet Mar 15, 2026
e2ff95a
Merge branch 'main' into stackable_ctx_data
caugonnet Mar 16, 2026
9b0bab1
Fix header path
caugonnet Mar 16, 2026
f0a0970
Merge branch 'main' into stackable_ctx_data
caugonnet Mar 16, 2026
a449392
Merge branch 'main' into stackable_ctx_data
caugonnet Mar 16, 2026
cd47ea9
Merge branch 'main' into stackable_ctx_data
caugonnet Mar 18, 2026
f7ad3c5
Merge branch 'main' into stackable_ctx_data
caugonnet Mar 18, 2026
4d3add9
Merge branch 'main' into stackable_ctx_data
caugonnet Mar 18, 2026
c1f5477
Merge branch 'main' into stackable_ctx_data
caugonnet Mar 18, 2026
a333433
[STF] Refactor stackable_task_dep to lazy resolution
caugonnet Mar 18, 2026
e079ef6
Rename access_mode_is_compatible to access_mode_permits and handle al…
caugonnet Mar 19, 2026
e587dcb
clang-format
caugonnet Mar 19, 2026
9cdb247
Merge branch 'main' into stackable_ctx_data
caugonnet Mar 19, 2026
ea52412
Remove dead code
caugonnet Mar 19, 2026
aab737f
Merge branch 'main' into stackable_ctx_data
caugonnet Mar 19, 2026
8677364
fix access_mode_permits for write-only access
caugonnet Mar 19, 2026
cca428e
Merge branch 'main' into stackable_ctx_data
caugonnet Mar 19, 2026
33c6f34
Fix access_mode_permits and repeat_graph_scope_guard destruction order
caugonnet Mar 19, 2026
da2243e
Fix dummy_graph leak in graph_ctx_node constructor
caugonnet Mar 19, 2026
2fb2393
Fix set_write_back to use data_root_offset instead of hardcoded 0
caugonnet Mar 19, 2026
0655141
Initialize get_cnt to 0 in data_node
caugonnet Mar 19, 2026
cac496d
Enable correctness assertion in stackable_nested_repeat test
caugonnet Mar 19, 2026
0554463
Implement node_hierarchy growth and fix head_map thread safety
caugonnet Mar 19, 2026
8325b22
clang-format
caugonnet Mar 19, 2026
5e4cdab
Remove redundant optional wrapper around unique_ptr in nodes vector
caugonnet Mar 19, 2026
bc601cb
Unify access-mode combining logic with shared helpers
caugonnet Mar 19, 2026
71f6c4f
Document why deferred_task_builder is needed
caugonnet Mar 19, 2026
c573637
Remove superfluous const/mutable from deferred_task_builder
caugonnet Mar 19, 2026
0a30c31
Add comment explaining why deferred_task_builder stores a reference
caugonnet Mar 19, 2026
fa87356
Extract node_hierarchy into its own header
caugonnet Mar 19, 2026
0147c01
Make warning_count thread-safe with std::atomic
caugonnet Mar 19, 2026
144299f
Remove no-op data_node destructor
caugonnet Mar 19, 2026
c9aafee
Remove unused was_destroyed flag from stackable_logical_data state
caugonnet Mar 19, 2026
40a8f60
Remove dead code: unused parent_depth, commented-out debug fprintf
caugonnet Mar 19, 2026
d5452b6
Remove dead print_logical_data_summary and traverse_nodes
caugonnet Mar 19, 2026
459b96b
Remove unused deferred_arg_info struct
caugonnet Mar 19, 2026
c486b99
Remove commented-out offset_depth variable
caugonnet Mar 19, 2026
156e5d4
Fix ~impl destroying wrong data_node via pop_back
caugonnet Mar 19, 2026
7593a91
Replace inline bounds-check + has_value with was_imported()
caugonnet Mar 19, 2026
9f6460c
Merge branch 'main' into stackable_ctx_data
caugonnet Mar 19, 2026
6d82683
Simplify access_mode_permits to a mutating-implies-mutating check
caugonnet Mar 19, 2026
f5755d4
Make get_head_offset() always abort on missing thread context
caugonnet Mar 19, 2026
ba8d456
Remove const/mutable abuse from stackable_logical_data state
caugonnet Mar 19, 2026
de58572
Remove redundant public: label in deferred_task_builder
caugonnet Mar 19, 2026
bdb6385
Use node_hierarchy growth constants in grow_data_nodes
caugonnet Mar 19, 2026
e3ee676
Remove unused manual lock/unlock methods from impl
caugonnet Mar 19, 2026
19719cc
Minor formatting cleanup in node_hierarchy and stackable test
caugonnet Mar 19, 2026
4bcfd6c
Fix copy-paste EXPECT value in stackable_illegal_export
caugonnet Mar 19, 2026
2593305
Clean up stackable tests and examples
caugonnet Mar 19, 2026
8ca6f6e
Fix ~impl: do not destroy root data_node while frozen_ld may be active
caugonnet Mar 20, 2026
69d714b
Merge branch 'main' into stackable_ctx_data
caugonnet Mar 20, 2026
8864e7b
Add result validation to graph_scope_test
caugonnet Mar 20, 2026
c55c3bc
Add error tests for fence() and wait() in nested stackable context
caugonnet Mar 20, 2026
140017d
Add focused test for stackable_logical_data::set_read_only
caugonnet Mar 20, 2026
f7c3c19
Add minimal iterative solver example using while_graph_scope
caugonnet Mar 20, 2026
457cde8
Fix wait-in-nested error test to use scalar_view
caugonnet Mar 20, 2026
a951fb8
Remove stackable DESIGN.md from tracked files
caugonnet Mar 20, 2026
035425a
Extract scope guards from stackable_ctx class body
caugonnet Mar 20, 2026
7f486e5
Fix scope guard ordering: define guards before inline unit tests
caugonnet Mar 20, 2026
30ffa9e
Less compute intensive default burger parameters to avoid 20 minutes …
caugonnet Mar 20, 2026
5c8f6f0
Merge branch 'main' into stackable_ctx_data
caugonnet Mar 20, 2026
47c36a5
Move dot_cleanup.py from benchmarks/scripts/ to cudax/scripts/stf/
caugonnet Mar 20, 2026
6699386
Merge branch 'main' into stackable_ctx_data
caugonnet Mar 23, 2026
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions cudax/examples/stf/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -19,6 +19,7 @@ set(stf_example_codegen_sources
01-axpy-launch.cu
01-axpy-parallel_for.cu
binary_fhe.cu
binary_fhe_stackable.cu
09-dot-reduce.cu
cfd.cu
custom_data_interface.cu
Expand Down
247 changes: 247 additions & 0 deletions cudax/examples/stf/binary_fhe_stackable.cu
Original file line number Diff line number Diff line change
@@ -0,0 +1,247 @@
//===----------------------------------------------------------------------===//
//
// Part of CUDASTF in CUDA C++ Core Libraries,
// under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
// SPDX-FileCopyrightText: Copyright (c) 2024-2024 NVIDIA CORPORATION & AFFILIATES.
//
//===----------------------------------------------------------------------===//

/**
* @file
* @brief A toy example to illustrate how we can compose logical operations over encrypted data
*/

#include "cuda/experimental/__stf/utility/stackable_ctx.cuh"
#include "cuda/experimental/stf.cuh"

using namespace cuda::experimental::stf;

class ciphertext;

class plaintext
{
public:
plaintext(const stackable_ctx& ctx)
: ctx(ctx)
{}

plaintext(stackable_ctx& ctx, ::std::vector<char> v)
: values(mv(v))
, ctx(ctx)
{
l = ctx.logical_data(values.data(), values.size());
}

void set_symbol(std::string s)
{
l.set_symbol(s);
symbol = s;
}

std::string get_symbol() const
{
return symbol;
}

std::string symbol;

const stackable_logical_data<slice<char>>& data() const
{
return l;
}

stackable_logical_data<slice<char>>& data()
{
return l;
}

// This will asynchronously fill string s
void convert_to_vector(std::vector<char>& v)
{
ctx.host_launch(l.read()).set_symbol("to_vector")->*[&](auto dl) {
v.resize(dl.size());
for (size_t i = 0; i < dl.size(); i++)
{
v[i] = dl(i);
}
};
}

ciphertext encrypt() const;

stackable_logical_data<slice<char>> l;

template <typename... Pack>
void push(Pack&&... pack)
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

remove push pop in this example

{
l.push(::std::forward<Pack>(pack)...);
}

void pop()
{
l.pop();
}

private:
std::vector<char> values;
mutable stackable_ctx ctx;
};

class ciphertext
{
public:
ciphertext() = default;
ciphertext(const ciphertext&) = default;

ciphertext(const stackable_ctx& ctx)
: ctx(ctx)
{}

plaintext decrypt() const
{
plaintext p(ctx);
p.l = ctx.logical_data(shape_of<slice<char>>(l.shape().size()));
ctx.parallel_for(l.shape(), l.read(), p.l.write()).set_symbol("decrypt")->*
[] __device__(size_t i, auto dctxt, auto dptxt) {
dptxt(i) = char((dctxt(i) >> 32));
};
return p;
}

// Copy assignment operator
ciphertext& operator=(const ciphertext& other)
{
if (this != &other)
{
assert(l.shape() == other.l.shape());
other.ctx.parallel_for(l.shape(), other.l.read(), l.write()).set_symbol("copy")->*
[] __device__(size_t i, auto other, auto result) {
result(i) = other(i);
};
}
return *this;
}

ciphertext operator|(const ciphertext& other) const
{
ciphertext result(ctx);
result.l = ctx.logical_data(data().shape());

ctx.parallel_for(data().shape(), data().read(), other.data().read(), result.data().write()).set_symbol("OR")->*
[] __device__(size_t i, auto d_c1, auto d_c2, auto d_res) {
d_res(i) = d_c1(i) | d_c2(i);
};

return result;
}

ciphertext operator&(const ciphertext& other) const
{
ciphertext result(ctx);
result.l = ctx.logical_data(data().shape());

ctx.parallel_for(data().shape(), data().read(), other.data().read(), result.data().write()).set_symbol("AND")->*
[] __device__(size_t i, auto d_c1, auto d_c2, auto d_res) {
d_res(i) = d_c1(i) & d_c2(i);
};

return result;
}

ciphertext operator~() const
{
ciphertext result(ctx);
result.l = ctx.logical_data(data().shape());
ctx.parallel_for(data().shape(), data().read(), result.data().write()).set_symbol("NOT")->*
[] __device__(size_t i, auto d_c, auto d_res) {
d_res(i) = ~d_c(i);
};

return result;
}

const stackable_logical_data<slice<uint64_t>>& data() const
{
return l;
}

stackable_logical_data<slice<uint64_t>>& data()
{
return l;
}

stackable_logical_data<slice<uint64_t>> l;

template <typename... Pack>
void push(Pack&&... pack)
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

remove push/pop

{
l.push(::std::forward<Pack>(pack)...);
}

void pop()
{
l.pop();
}

private:
mutable stackable_ctx ctx;
};

ciphertext plaintext::encrypt() const
{
ciphertext c(ctx);
c.l = ctx.logical_data(shape_of<slice<uint64_t>>(l.shape().size()));

ctx.parallel_for(l.shape(), l.read(), c.l.write()).set_symbol("encrypt")->*
[] __device__(size_t i, auto dptxt, auto dctxt) {
// A super safe encryption !
dctxt(i) = ((uint64_t) (dptxt(i)) << 32 | 0x4);
};

return c;
}

template <typename T>
T circuit(const T& a, const T& b)
{
return (~((a | ~b) & (~a | b)));
}

int main()
{
stackable_ctx ctx;

std::vector<char> vA{3, 3, 2, 2, 17};
plaintext pA(ctx, vA);
pA.set_symbol("A");

std::vector<char> vB{1, 7, 7, 7, 49};
plaintext pB(ctx, vB);
pB.set_symbol("B");

auto eA = pA.encrypt();
auto eB = pB.encrypt();

ctx.push();

eA.push(access_mode::read);
eB.push(access_mode::read);

// TODO find a way to get "out" outside of this scope to do decryption in the main ctx
auto out = circuit(eA, eB);

std::vector<char> v_out;
out.decrypt().convert_to_vector(v_out);

ctx.pop();

ctx.finalize();

for (size_t i = 0; i < v_out.size(); i++)
{
char expected = circuit(vA[i], vB[i]);
EXPECT(expected == v_out[i]);
}
}
Loading