From 1b8cb1ddbef3f858b64afeeb109b1c34448d8b7a Mon Sep 17 00:00:00 2001 From: feiy Date: Sat, 9 May 2020 09:40:57 +0800 Subject: [PATCH] Switching to CFFI. --- DVTuple.cpp | 5 - DVTuple.h | 3 - DVVector.cpp | 19 +- DVVector.h | 7 +- DeviceViewable.h | 58 +++++- docs/QuickStartGuide.md | 4 +- fake_vectors/DVConstant.cpp | 12 +- fake_vectors/DVConstant.h | 1 - fake_vectors/DVCounter.cpp | 10 +- fake_vectors/DVCounter.h | 1 - fake_vectors/DVCustomVector.cpp | 6 - fake_vectors/DVCustomVector.h | 2 - fake_vectors/DVDiscard.cpp | 6 +- fake_vectors/DVDiscard.h | 1 - fake_vectors/DVPermutation.cpp | 8 +- fake_vectors/DVPermutation.h | 1 - fake_vectors/DVRange.cpp | 8 +- fake_vectors/DVRange.h | 1 - fake_vectors/DVReverse.cpp | 7 +- fake_vectors/DVReverse.h | 1 - fake_vectors/DVTransform.cpp | 7 +- fake_vectors/DVTransform.h | 3 - fake_vectors/DVZipped.cpp | 10 +- fake_vectors/DVZipped.h | 2 - functor.cpp | 5 - functor.h | 2 - python/CMakeLists.txt | 85 +++------ python/Context.hpp | 262 ---------------------------- python/DVConstant.hpp | 12 -- python/DVCounter.hpp | 12 -- python/DVCustomVector.hpp | 23 --- python/DVDiscard.hpp | 12 -- python/DVPermutation.hpp | 12 -- python/DVRange.hpp | 27 --- python/DVReverse.hpp | 11 -- python/DVTransform.hpp | 13 -- python/DVTuple.hpp | 18 -- python/DVVector.hpp | 77 -------- python/DVZipped.hpp | 26 --- python/DeviceViewable.hpp | 54 ------ python/PyThrustRTC.cpp | 177 ------------------- python/ThrustRTC/Context.py | 48 +++-- python/ThrustRTC/Copying.py | 9 +- python/ThrustRTC/DVTuple.py | 9 +- python/ThrustRTC/DVVector.py | 18 +- python/ThrustRTC/DeviceViewable.py | 80 +++++++-- python/ThrustRTC/FakeVectors.py | 30 +++- python/ThrustRTC/Functor.py | 46 +++-- python/ThrustRTC/Merging.py | 6 +- python/ThrustRTC/Native.py | 18 ++ python/ThrustRTC/PrefixSums.py | 18 +- python/ThrustRTC/Reductions.py | 87 ++++++--- python/ThrustRTC/Reordering.py | 10 +- python/ThrustRTC/Searching.py | 58 ++++-- python/ThrustRTC/Sorting.py | 8 +- python/ThrustRTC/Transformations.py | 17 +- python/ThrustRTC/__init__.py | 6 - python/ThrustRTC/cffi.py | 8 + python/ThrustRTC/cffi_build.py | 181 +++++++++++++++++++ python/ThrustRTC/utils.py | 34 ++++ python/adjacent_difference.hpp | 30 ---- python/api.h | 184 +++++++++++++++++++ python/api_Context.cpp | 200 +++++++++++++++++++++ python/api_Copying.cpp | 96 ++++++++++ python/api_DVTuple.cpp | 21 +++ python/api_DVVector.cpp | 71 ++++++++ python/api_DeviceViewable.cpp | 135 ++++++++++++++ python/api_FakeVectors.cpp | 89 ++++++++++ python/api_Functor.cpp | 34 ++++ python/api_Merging.cpp | 54 ++++++ python/api_PrefixSums.cpp | 113 ++++++++++++ python/api_Reductions.cpp | 256 +++++++++++++++++++++++++++ python/api_Reordering.cpp | 165 ++++++++++++++++++ python/api_Searching.cpp | 232 ++++++++++++++++++++++++ python/api_Sorting.cpp | 47 +++++ python/api_Transformations.cpp | 200 +++++++++++++++++++++ python/api_utils.cpp | 50 ++++++ python/binary_search.hpp | 160 ----------------- python/copy.hpp | 38 ---- python/count.hpp | 26 --- python/equal.hpp | 25 --- python/extrema.hpp | 73 -------- python/fill.hpp | 13 -- python/find.hpp | 38 ---- python/for_each.hpp | 13 -- python/functor.hpp | 30 ---- python/gather.hpp | 42 ----- python/inner_product.hpp | 36 ---- python/logical.hpp | 36 ---- python/merge.hpp | 60 ------- python/mismatch.hpp | 28 --- python/partition.hpp | 77 -------- python/reduce.hpp | 79 --------- python/remove.hpp | 76 -------- python/replace.hpp | 49 ------ python/scan.hpp | 138 --------------- python/scatter.hpp | 42 ----- python/sequence.hpp | 38 ---- python/setup.py | 90 +--------- python/sort.hpp | 102 ----------- python/swap.hpp | 13 -- python/tabulate.hpp | 13 -- python/test/CMakeLists.txt | 7 +- python/transform.hpp | 64 ------- python/transform_reduce.hpp | 16 -- python/transform_scan.hpp | 31 ---- python/unique.hpp | 110 ------------ python/viewbuf_to_python.hpp | 38 ---- 108 files changed, 2594 insertions(+), 2686 deletions(-) delete mode 100644 python/Context.hpp delete mode 100644 python/DVConstant.hpp delete mode 100644 python/DVCounter.hpp delete mode 100644 python/DVCustomVector.hpp delete mode 100644 python/DVDiscard.hpp delete mode 100644 python/DVPermutation.hpp delete mode 100644 python/DVRange.hpp delete mode 100644 python/DVReverse.hpp delete mode 100644 python/DVTransform.hpp delete mode 100644 python/DVTuple.hpp delete mode 100644 python/DVVector.hpp delete mode 100644 python/DVZipped.hpp delete mode 100644 python/DeviceViewable.hpp delete mode 100644 python/PyThrustRTC.cpp create mode 100644 python/ThrustRTC/Native.py create mode 100644 python/ThrustRTC/cffi.py create mode 100644 python/ThrustRTC/cffi_build.py create mode 100644 python/ThrustRTC/utils.py delete mode 100644 python/adjacent_difference.hpp create mode 100644 python/api.h create mode 100644 python/api_Context.cpp create mode 100644 python/api_Copying.cpp create mode 100644 python/api_DVTuple.cpp create mode 100644 python/api_DVVector.cpp create mode 100644 python/api_DeviceViewable.cpp create mode 100644 python/api_FakeVectors.cpp create mode 100644 python/api_Functor.cpp create mode 100644 python/api_Merging.cpp create mode 100644 python/api_PrefixSums.cpp create mode 100644 python/api_Reductions.cpp create mode 100644 python/api_Reordering.cpp create mode 100644 python/api_Searching.cpp create mode 100644 python/api_Sorting.cpp create mode 100644 python/api_Transformations.cpp create mode 100644 python/api_utils.cpp delete mode 100644 python/binary_search.hpp delete mode 100644 python/copy.hpp delete mode 100644 python/count.hpp delete mode 100644 python/equal.hpp delete mode 100644 python/extrema.hpp delete mode 100644 python/fill.hpp delete mode 100644 python/find.hpp delete mode 100644 python/for_each.hpp delete mode 100644 python/functor.hpp delete mode 100644 python/gather.hpp delete mode 100644 python/inner_product.hpp delete mode 100644 python/logical.hpp delete mode 100644 python/merge.hpp delete mode 100644 python/mismatch.hpp delete mode 100644 python/partition.hpp delete mode 100644 python/reduce.hpp delete mode 100644 python/remove.hpp delete mode 100644 python/replace.hpp delete mode 100644 python/scan.hpp delete mode 100644 python/scatter.hpp delete mode 100644 python/sequence.hpp delete mode 100644 python/sort.hpp delete mode 100644 python/swap.hpp delete mode 100644 python/tabulate.hpp delete mode 100644 python/transform.hpp delete mode 100644 python/transform_reduce.hpp delete mode 100644 python/transform_scan.hpp delete mode 100644 python/unique.hpp delete mode 100644 python/viewbuf_to_python.hpp diff --git a/DVTuple.cpp b/DVTuple.cpp index d95e6b2..d3e151b 100644 --- a/DVTuple.cpp +++ b/DVTuple.cpp @@ -18,11 +18,6 @@ DVTuple::DVTuple(const std::vector& elem_map) TRTC_Query_Struct(m_name_view_cls.c_str(), name_elems, m_offsets.data()); } -std::string DVTuple::name_view_cls() const -{ - return m_name_view_cls; -} - ViewBuf DVTuple::view() const { ViewBuf ret(m_offsets[m_offsets.size() - 1]); diff --git a/DVTuple.h b/DVTuple.h index 6937431..39850a4 100644 --- a/DVTuple.h +++ b/DVTuple.h @@ -9,12 +9,9 @@ class THRUST_RTC_API DVTuple : public DeviceViewable { public: DVTuple(const std::vector& elem_map); - - virtual std::string name_view_cls() const; virtual ViewBuf view() const; private: - std::string m_name_view_cls; std::vector m_view_elems; std::vector m_offsets; diff --git a/DVVector.cpp b/DVVector.cpp index 361b07e..310be41 100644 --- a/DVVector.cpp +++ b/DVVector.cpp @@ -20,6 +20,8 @@ DVVector::DVVector(const char* elem_cls, size_t size, void* hdata) cuMemcpyHtoD(dptr, hdata, m_elem_size*m_size); else cuMemsetD8(dptr, 0, m_elem_size*m_size); + + m_name_view_cls = std::string("VectorView<") + m_elem_cls + ">"; } DVVector::~DVVector() @@ -34,11 +36,6 @@ void DVVector::to_host(void* hdata, size_t begin, size_t end) const cuMemcpyDtoH(hdata, (CUdeviceptr)((char*)m_data + begin* m_elem_size), m_elem_size*n); } -std::string DVVector::name_view_cls() const -{ - return std::string("VectorView<") + m_elem_cls + ">"; -} - ViewBuf DVVector::view() const { ViewBuf buf(sizeof(VectorView)); @@ -49,26 +46,26 @@ ViewBuf DVVector::view() const } DVVectorAdaptor::DVVectorAdaptor(const char* elem_cls, size_t size, void* ddata) - : DVVectorLike(elem_cls, (std::string(elem_cls) + "&").c_str(), size), m_data(ddata){} + : DVVectorLike(elem_cls, (std::string(elem_cls) + "&").c_str(), size), m_data(ddata) +{ + m_name_view_cls = std::string("VectorView<") + m_elem_cls + ">"; +} DVVectorAdaptor::DVVectorAdaptor(const DVVector& vec, size_t begin, size_t end) : DVVectorLike(vec.name_elem_cls().c_str(), vec.name_ref_type().c_str(), (end == (size_t)(-1)? vec.size() : end) - begin) { m_data = (char*)vec.data() + begin * vec.elem_size(); + m_name_view_cls = std::string("VectorView<") + m_elem_cls + ">"; } DVVectorAdaptor::DVVectorAdaptor(const DVVectorAdaptor& vec, size_t begin, size_t end) : DVVectorLike(vec.name_elem_cls().c_str(), vec.name_ref_type().c_str(), (end == (size_t)(-1) ? vec.size() : end) - begin) { m_data = (char*)vec.data() + begin * vec.elem_size(); + m_name_view_cls = std::string("VectorView<") + m_elem_cls + ">"; } -std::string DVVectorAdaptor::name_view_cls() const -{ - return std::string("VectorView<") + m_elem_cls + ">"; -} - ViewBuf DVVectorAdaptor::view() const { ViewBuf buf(sizeof(VectorView)); diff --git a/DVVector.h b/DVVector.h index 8ce3829..f60faa1 100644 --- a/DVVector.h +++ b/DVVector.h @@ -8,8 +8,8 @@ class THRUST_RTC_API DVVectorLike : public DeviceViewable { public: - std::string name_elem_cls() const { return m_elem_cls; } - std::string name_ref_type() const { return m_ref_type; } + const std::string& name_elem_cls() const { return m_elem_cls; } + const std::string& name_ref_type() const { return m_ref_type; } size_t elem_size() const { return m_elem_size; } size_t size() const { return m_size; } @@ -35,7 +35,6 @@ class THRUST_RTC_API DVVector : public DVVectorLike virtual bool is_writable() const { return true; } void to_host(void* hdata, size_t begin=0, size_t end = (size_t)(-1)) const; - virtual std::string name_view_cls() const; virtual ViewBuf view() const; private: @@ -53,8 +52,6 @@ class THRUST_RTC_API DVVectorAdaptor : public DVVectorLike DVVectorAdaptor(const DVVectorAdaptor& vec, size_t begin = 0, size_t end = (size_t)(-1)); virtual bool is_writable() const { return true; } - - virtual std::string name_view_cls() const; virtual ViewBuf view() const; private: diff --git a/DeviceViewable.h b/DeviceViewable.h index be73839..8c7632a 100644 --- a/DeviceViewable.h +++ b/DeviceViewable.h @@ -14,8 +14,11 @@ class DeviceViewable public: DeviceViewable(){} virtual ~DeviceViewable(){} - virtual std::string name_view_cls() const = 0; virtual ViewBuf view() const = 0; + const std::string& name_view_cls() const { return m_name_view_cls; } + +protected: + std::string m_name_view_cls; }; @@ -36,18 +39,12 @@ class BuiltIn : public DeviceViewable memcpy(m_view_buf.data(), data_view, size_view); } - virtual std::string name_view_cls() const - { - return m_name_view_cls; - } - virtual ViewBuf view() const { return m_view_buf; } private: - std::string m_name_view_cls; ViewBuf m_view_buf; }; @@ -84,4 +81,51 @@ DECLAR_DV_BASIC(DVUInt64, uint64_t) DECLAR_DV_BASIC(DVSizeT, size_t) + +class SomeDeviceViewable : public DeviceViewable +{ +public: + SomeDeviceViewable(const ViewBuf& buf, const char* type) + { + m_buf = buf; + m_name_view_cls = type; + + } + virtual ViewBuf view() const + { + return m_buf; + } +private: + ViewBuf m_buf; +}; + +inline DeviceViewable* dv_from_viewbuf(const ViewBuf& buf, const char* type) +{ + std::string s_type = type; + if (s_type == "int8_t") + return new DVInt8(*(int8_t*)buf.data()); + else if (s_type == "uint8_t") + return new DVUInt8(*(uint8_t*)buf.data()); + else if (s_type == "int16_t") + return new DVInt16(*(int16_t*)buf.data()); + else if (s_type == "uint16_t") + return new DVUInt16(*(uint16_t*)buf.data()); + else if (s_type == "int32_t") + return new DVInt32(*(int32_t*)buf.data()); + else if (s_type == "uint32_t") + return new DVUInt32(*(uint32_t*)buf.data()); + else if (s_type == "int64_t") + return new DVInt64(*(int64_t*)buf.data()); + else if (s_type == "uint64_t") + return new DVUInt64(*(uint64_t*)buf.data()); + else if (s_type == "float") + return new DVFloat(*(float*)buf.data()); + else if (s_type == "double") + return new DVDouble(*(double*)buf.data()); + else if (s_type == "bool") + return new DVBool(*(bool*)buf.data()); + else + return new SomeDeviceViewable(buf, type); +} + #endif diff --git a/docs/QuickStartGuide.md b/docs/QuickStartGuide.md index 904b285..de15c73 100644 --- a/docs/QuickStartGuide.md +++ b/docs/QuickStartGuide.md @@ -29,7 +29,6 @@ At build time, you will only need: * UnQLite source code, as submodule: thirdparty/unqlite * CMake 3.x -* C libraries of Python 3 is required to build the Python binding part of the code. After cloning the repo from github and resolving the submodules, you can build it with CMake. Libraries for different lauguages are now built separately. The C++ library is used as a @@ -57,7 +56,7 @@ You will get the library headers, binaries and examples in the "install" directo #### Install ThrustRTC for Python from Pypi -Builds for Win64/Linux64 + Python 3.7 are available from [Pypi](https://pypi.org/project/ThrustRTC/) +Builds for Win64/Linux64 + Python 3.x are available from [Pypi](https://pypi.org/project/ThrustRTC/) If your environment matches, you can try: ``` @@ -90,6 +89,7 @@ Zip packages are available at: For Python * Python 3 +* cffi * numpy * numba (optional) diff --git a/fake_vectors/DVConstant.cpp b/fake_vectors/DVConstant.cpp index 8397a28..be6a479 100644 --- a/fake_vectors/DVConstant.cpp +++ b/fake_vectors/DVConstant.cpp @@ -4,15 +4,9 @@ DVConstant::DVConstant(const DeviceViewable& dvobj, size_t size) : DVVectorLike(dvobj.name_view_cls().c_str(), (std::string("const ") + dvobj.name_view_cls() + "&").c_str(), size) { - m_value = dvobj.view(); - - std::string name_struct = name_view_cls(); - TRTC_Query_Struct(name_struct.c_str(), { "_size", "_value" }, m_offsets); -} - -std::string DVConstant::name_view_cls() const -{ - return std::string("ConstantView<") + m_elem_cls + ">"; + m_value = dvobj.view(); + m_name_view_cls = std::string("ConstantView<") + m_elem_cls + ">"; + TRTC_Query_Struct(m_name_view_cls.c_str(), { "_size", "_value" }, m_offsets); } ViewBuf DVConstant::view() const diff --git a/fake_vectors/DVConstant.h b/fake_vectors/DVConstant.h index 84b232f..d311ec4 100644 --- a/fake_vectors/DVConstant.h +++ b/fake_vectors/DVConstant.h @@ -8,7 +8,6 @@ class THRUST_RTC_API DVConstant : public DVVectorLike public: ViewBuf value() const { return m_value; } DVConstant(const DeviceViewable& dvobj, size_t size = (size_t)(-1)); - virtual std::string name_view_cls() const; virtual ViewBuf view() const; private: diff --git a/fake_vectors/DVCounter.cpp b/fake_vectors/DVCounter.cpp index b267428..b91b4fa 100644 --- a/fake_vectors/DVCounter.cpp +++ b/fake_vectors/DVCounter.cpp @@ -5,14 +5,8 @@ DVCounter::DVCounter(const DeviceViewable& dvobj_init, size_t size) : DVVectorLike(dvobj_init.name_view_cls().c_str(), dvobj_init.name_view_cls().c_str(), size) { m_value_init = dvobj_init.view(); - - std::string name_struct = name_view_cls(); - TRTC_Query_Struct(name_struct.c_str(), { "_size", "_value_init" }, m_offsets); -} - -std::string DVCounter::name_view_cls() const -{ - return std::string("CounterView<") + m_elem_cls + ">"; + m_name_view_cls = std::string("CounterView<") + m_elem_cls + ">"; + TRTC_Query_Struct(m_name_view_cls.c_str(), { "_size", "_value_init" }, m_offsets); } ViewBuf DVCounter::view() const diff --git a/fake_vectors/DVCounter.h b/fake_vectors/DVCounter.h index 6432869..200e5a3 100644 --- a/fake_vectors/DVCounter.h +++ b/fake_vectors/DVCounter.h @@ -8,7 +8,6 @@ class THRUST_RTC_API DVCounter : public DVVectorLike public: ViewBuf value_init() const { return m_value_init; } DVCounter(const DeviceViewable& dvobj_init, size_t size = (size_t)(-1)); - virtual std::string name_view_cls() const; virtual ViewBuf view() const; private: diff --git a/fake_vectors/DVCustomVector.cpp b/fake_vectors/DVCustomVector.cpp index 8a4ba41..f734d13 100644 --- a/fake_vectors/DVCustomVector.cpp +++ b/fake_vectors/DVCustomVector.cpp @@ -38,12 +38,6 @@ DVCustomVector::DVCustomVector(const std::vector& arg_ma TRTC_Query_Struct(m_name_view_cls.c_str(), { "_size", "_op" }, m_offsets); } - -std::string DVCustomVector::name_view_cls() const -{ - return m_name_view_cls; -} - ViewBuf DVCustomVector::view() const { ViewBuf view_op(m_arg_offsets[m_arg_offsets.size() - 1]); diff --git a/fake_vectors/DVCustomVector.h b/fake_vectors/DVCustomVector.h index 376f497..2e11753 100644 --- a/fake_vectors/DVCustomVector.h +++ b/fake_vectors/DVCustomVector.h @@ -11,13 +11,11 @@ class THRUST_RTC_API DVCustomVector : public DVVectorLike DVCustomVector(const std::vector& arg_map, const char* name_idx, const char* code_body, const char* elem_cls, size_t size = (size_t)(-1), bool read_only = true); - virtual std::string name_view_cls() const; virtual ViewBuf view() const; virtual bool is_writable() const { return !m_read_only; } private: - std::string m_name_view_cls; size_t m_size; bool m_read_only; std::vector m_view_args; diff --git a/fake_vectors/DVDiscard.cpp b/fake_vectors/DVDiscard.cpp index b2d7be5..d5ea893 100644 --- a/fake_vectors/DVDiscard.cpp +++ b/fake_vectors/DVDiscard.cpp @@ -1,11 +1,9 @@ #include "DVDiscard.h" DVDiscard::DVDiscard(const char* elem_cls, size_t size) - :DVVectorLike(elem_cls, (std::string("_Sink<")+ elem_cls +">&").c_str(), size){} - -std::string DVDiscard::name_view_cls() const + :DVVectorLike(elem_cls, (std::string("_Sink<")+ elem_cls +">&").c_str(), size) { - return std::string("DiscardView<") + m_elem_cls + ">"; + m_name_view_cls = std::string("DiscardView<") + m_elem_cls + ">"; } ViewBuf DVDiscard::view() const diff --git a/fake_vectors/DVDiscard.h b/fake_vectors/DVDiscard.h index 653f7a0..98570c5 100644 --- a/fake_vectors/DVDiscard.h +++ b/fake_vectors/DVDiscard.h @@ -7,7 +7,6 @@ class THRUST_RTC_API DVDiscard : public DVVectorLike { public: DVDiscard(const char* elem_cls, size_t size = (size_t)(-1)); - virtual std::string name_view_cls() const; virtual ViewBuf view() const; virtual bool is_readable() const { return false; } virtual bool is_writable() const { return true; } diff --git a/fake_vectors/DVPermutation.cpp b/fake_vectors/DVPermutation.cpp index 08b7cc9..2a6ebe3 100644 --- a/fake_vectors/DVPermutation.cpp +++ b/fake_vectors/DVPermutation.cpp @@ -11,13 +11,9 @@ DVPermutation::DVPermutation(const DVVectorLike& vec_value, const DVVectorLike& m_cls_index = vec_index.name_view_cls(); m_view_index = vec_index.view(); - std::string name_struct = name_view_cls(); - TRTC_Query_Struct(name_struct.c_str(), { "_view_vec_value", "_view_vec_index" }, m_offsets); -} + m_name_view_cls = std::string("PermutationView<") + m_cls_value + ", " + m_cls_index + ">"; + TRTC_Query_Struct(m_name_view_cls.c_str(), { "_view_vec_value", "_view_vec_index" }, m_offsets); -std::string DVPermutation::name_view_cls() const -{ - return std::string("PermutationView<") + m_cls_value + ", "+ m_cls_index+">"; } ViewBuf DVPermutation::view() const diff --git a/fake_vectors/DVPermutation.h b/fake_vectors/DVPermutation.h index 70fca20..6dc7fe2 100644 --- a/fake_vectors/DVPermutation.h +++ b/fake_vectors/DVPermutation.h @@ -12,7 +12,6 @@ class THRUST_RTC_API DVPermutation : public DVVectorLike ViewBuf view_index() const { return m_view_index; } DVPermutation(const DVVectorLike& vec_value, const DVVectorLike& vec_index ); - virtual std::string name_view_cls() const; virtual ViewBuf view() const; virtual bool is_readable() const { return m_readable; } virtual bool is_writable() const { return m_writable; } diff --git a/fake_vectors/DVRange.cpp b/fake_vectors/DVRange.cpp index 45057b3..feb87a9 100644 --- a/fake_vectors/DVRange.cpp +++ b/fake_vectors/DVRange.cpp @@ -12,14 +12,10 @@ DVRange::DVRange(const DVVectorLike& vec_value, size_t begin, size_t end) m_begin = begin; m_end = end; - std::string name_struct = name_view_cls(); - TRTC_Query_Struct(name_struct.c_str(), { "_view_vec_value", "_begin", "_end" }, m_offsets); + m_name_view_cls = std::string("RangeView<") + m_cls_value + ">"; + TRTC_Query_Struct(m_name_view_cls.c_str(), { "_view_vec_value", "_begin", "_end" }, m_offsets); } -std::string DVRange::name_view_cls() const -{ - return std::string("RangeView<") + m_cls_value + ">"; -} ViewBuf DVRange::view() const { diff --git a/fake_vectors/DVRange.h b/fake_vectors/DVRange.h index 86e4df0..239fa12 100644 --- a/fake_vectors/DVRange.h +++ b/fake_vectors/DVRange.h @@ -10,7 +10,6 @@ class THRUST_RTC_API DVRange : public DVVectorLike ViewBuf view_value() const { return m_view_value; } DVRange(const DVVectorLike& vec_value, size_t begin = 0, size_t end = (size_t)(-1)); - virtual std::string name_view_cls() const; virtual ViewBuf view() const; virtual bool is_readable() const { return m_readable; } virtual bool is_writable() const { return m_writable; } diff --git a/fake_vectors/DVReverse.cpp b/fake_vectors/DVReverse.cpp index 29ddc6f..2caebec 100644 --- a/fake_vectors/DVReverse.cpp +++ b/fake_vectors/DVReverse.cpp @@ -8,15 +8,10 @@ DVReverse::DVReverse(const DVVectorLike& vec_value) m_writable = vec_value.is_writable(); m_cls_value = vec_value.name_view_cls(); m_view_value = vec_value.view(); -} - -std::string DVReverse::name_view_cls() const -{ - return std::string("ReverseView<") + m_cls_value + ">"; + m_name_view_cls = std::string("ReverseView<") + m_cls_value + ">"; } - ViewBuf DVReverse::view() const { ViewBuf buf(m_view_value.size()); diff --git a/fake_vectors/DVReverse.h b/fake_vectors/DVReverse.h index fc1f6dd..5caabd5 100644 --- a/fake_vectors/DVReverse.h +++ b/fake_vectors/DVReverse.h @@ -10,7 +10,6 @@ class THRUST_RTC_API DVReverse : public DVVectorLike ViewBuf view_value() const { return m_view_value; } DVReverse(const DVVectorLike& vec_value); - virtual std::string name_view_cls() const; virtual ViewBuf view() const; virtual bool is_readable() const { return m_readable; } virtual bool is_writable() const { return m_writable; } diff --git a/fake_vectors/DVTransform.cpp b/fake_vectors/DVTransform.cpp index a88a64c..e2cdd26 100644 --- a/fake_vectors/DVTransform.cpp +++ b/fake_vectors/DVTransform.cpp @@ -5,17 +5,12 @@ DVTransform::DVTransform(const DVVectorLike& vec_in, const char* elem_cls, const Functor& op) : DVVectorLike(elem_cls, elem_cls, vec_in.size()) { - m_name_view_cls = std::string("TransformView<") + elem_cls + "," + vec_in.name_view_cls() + "," + op.name_view_cls() + ">"; m_view_in = vec_in.view(); m_view_op = op.view(); + m_name_view_cls = std::string("TransformView<") + elem_cls + "," + vec_in.name_view_cls() + "," + op.name_view_cls() + ">"; TRTC_Query_Struct(m_name_view_cls.c_str(), { "_view_vec_in", "_view_op" }, m_offsets); } -std::string DVTransform::name_view_cls() const -{ - return m_name_view_cls; -} - ViewBuf DVTransform::view() const { ViewBuf buf(m_offsets[2]); diff --git a/fake_vectors/DVTransform.h b/fake_vectors/DVTransform.h index ee9ae11..2cb0757 100644 --- a/fake_vectors/DVTransform.h +++ b/fake_vectors/DVTransform.h @@ -8,12 +8,9 @@ class THRUST_RTC_API DVTransform : public DVVectorLike { public: DVTransform(const DVVectorLike& vec_in, const char* elem_cls, const Functor& op); - - virtual std::string name_view_cls() const; virtual ViewBuf view() const; private: - std::string m_name_view_cls; ViewBuf m_view_in; ViewBuf m_view_op; size_t m_offsets[3]; diff --git a/fake_vectors/DVZipped.cpp b/fake_vectors/DVZipped.cpp index feac093..3875aa2 100644 --- a/fake_vectors/DVZipped.cpp +++ b/fake_vectors/DVZipped.cpp @@ -74,16 +74,8 @@ DVZipped::DVZipped(const std::vector& vecs, const std::vector& vecs, const std::vector& elem_names); - virtual std::string name_view_cls() const; virtual ViewBuf view() const; virtual bool is_readable() const { return m_readable; } virtual bool is_writable() const { return m_writable; } @@ -15,7 +14,6 @@ class THRUST_RTC_API DVZipped : public DVVectorLike private: bool m_readable; bool m_writable; - std::string m_name_view_cls; std::vector m_view_elems; std::vector m_offsets; }; diff --git a/functor.cpp b/functor.cpp index ee1b045..4f39c41 100644 --- a/functor.cpp +++ b/functor.cpp @@ -51,11 +51,6 @@ Functor::Functor(const char* name_built_in_view_cls) m_offsets[0] = 1; } -std::string Functor::name_view_cls() const -{ - return m_name_view_cls; -} - ViewBuf Functor::view() const { ViewBuf ret(m_offsets[m_offsets.size()-1]); diff --git a/functor.h b/functor.h index 055ac69..f1a13ab 100644 --- a/functor.h +++ b/functor.h @@ -9,11 +9,9 @@ class THRUST_RTC_API Functor : public DeviceViewable Functor(const std::vector& arg_map, const std::vector& functor_params, const char* code_body); Functor(const char* name_built_in_view_cls); - virtual std::string name_view_cls() const; virtual ViewBuf view() const; private: - std::string m_name_view_cls; std::vector m_view_args; std::vector m_offsets; }; diff --git a/python/CMakeLists.txt b/python/CMakeLists.txt index 3798a7a..116c1fd 100644 --- a/python/CMakeLists.txt +++ b/python/CMakeLists.txt @@ -1,7 +1,6 @@ cmake_minimum_required (VERSION 3.0) project(PyThrustRTC) -find_package(PythonLibs 3 REQUIRED) add_executable(PackHeaders ../internal/pack_headers.cpp) add_custom_target(Run_PackHeaders @@ -11,6 +10,10 @@ COMMENT "Running PackHeaders in ${CMAKE_CURRENT_SOURCE_DIR}/.." SOURCES ../internal/pack_headers.cpp ) +add_custom_target(Run_CFFIBuild +COMMAND python ThrustRTC/cffi_build.py +WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR} +COMMENT "Running cffi_build") set(LIB_SOURCES ../thirdparty/crc64/crc64.cpp @@ -136,55 +139,26 @@ set(INTERNAL_HEADERS ) set (SRC -PyThrustRTC.cpp -Context.hpp -viewbuf_to_python.hpp -DeviceViewable.hpp -DVVector.hpp -DVTuple.hpp -DVRange.hpp -DVConstant.hpp -DVCounter.hpp -DVDiscard.hpp -DVPermutation.hpp -DVReverse.hpp -DVTransform.hpp -DVZipped.hpp -DVCustomVector.hpp -fill.hpp -functor.hpp -replace.hpp -for_each.hpp -adjacent_difference.hpp -sequence.hpp -tabulate.hpp -transform.hpp -gather.hpp -scatter.hpp -copy.hpp -swap.hpp -count.hpp -reduce.hpp -equal.hpp -extrema.hpp -inner_product.hpp -transform_reduce.hpp -logical.hpp -scan.hpp -transform_scan.hpp -remove.hpp -unique.hpp -partition.hpp -find.hpp -mismatch.hpp -binary_search.hpp -merge.hpp -sort.hpp +api.h +api_utils.cpp +api_Context.cpp +api_DeviceViewable.cpp +api_DVVector.cpp +api_DVTuple.cpp +api_FakeVectors.cpp +api_Functor.cpp +api_Transformations.cpp +api_Copying.cpp +api_Reductions.cpp +api_PrefixSums.cpp +api_Reordering.cpp +api_Searching.cpp +api_Merging.cpp +api_Sorting.cpp ) set (INCLUDE_DIR -${PYTHON_INCLUDE_DIRS} . ../thirdparty/crc64 ../thirdparty/unqlite @@ -210,12 +184,7 @@ add_definitions(${DEFINES}) add_library(unqlite STATIC ../thirdparty/unqlite/unqlite.h ../thirdparty/unqlite/unqlite.c) add_library(PyThrustRTC SHARED ${LIB_SOURCES} ${LIB_HEADERS} ${LIB_HEADERS_FAKEVECTORS} ${INTERNAL_HEADERS} ${SRC}) -target_link_libraries(PyThrustRTC ${PYTHON_LIBRARIES} unqlite) - -if (WIN32) -else() -target_link_libraries(PyThrustRTC dl) -endif() +target_link_libraries(PyThrustRTC unqlite) add_dependencies(Run_PackHeaders PackHeaders) add_dependencies(PyThrustRTC Run_PackHeaders) @@ -224,13 +193,6 @@ if (WIN32) target_compile_definitions(PyThrustRTC PRIVATE THRUST_RTC_DLL_EXPORT) endif() -if (WIN32) -set_target_properties(PyThrustRTC PROPERTIES SUFFIX ".pyd") -else() -set_target_properties(PyThrustRTC PROPERTIES SUFFIX ".so") -set_target_properties(PyThrustRTC PROPERTIES PREFIX "") -endif() - IF(CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT) SET(CMAKE_INSTALL_PREFIX ../install CACHE PATH "Install path" FORCE) ENDIF(CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT) @@ -255,6 +217,9 @@ endif() set(PYTHON ThrustRTC/__init__.py + ThrustRTC/cffi.py + ThrustRTC/Native.py + ThrustRTC/utils.py ThrustRTC/Context.py ThrustRTC/DeviceViewable.py ThrustRTC/DVVector.py @@ -272,7 +237,7 @@ set(PYTHON ) install(FILES ${PYTHON} DESTINATION test_python/ThrustRTC) - +install(FILES setup.py README.md DESTINATION test_python) set(INCLUDE_TESTS false CACHE BOOL "Include tests") diff --git a/python/Context.hpp b/python/Context.hpp deleted file mode 100644 index dda1b47..0000000 --- a/python/Context.hpp +++ /dev/null @@ -1,262 +0,0 @@ -#include -#include "TRTCContext.h" - -static PyObject* n_set_verbose(PyObject* self, PyObject* args) -{ - bool verbose = PyObject_IsTrue(PyTuple_GetItem(args, 0)) != 0; - TRTC_Set_Verbose(verbose); - return PyLong_FromLong(0); -} - -static PyObject* n_add_include_dir(PyObject* self, PyObject* args) -{ - const char* dir = PyUnicode_AsUTF8(PyTuple_GetItem(args, 0)); - TRTC_Add_Include_Dir(dir); - return PyLong_FromLong(0); -} - -static PyObject* n_add_built_in_header(PyObject* self, PyObject* args) -{ - const char* filename = PyUnicode_AsUTF8(PyTuple_GetItem(args, 0)); - const char* filecontent = PyUnicode_AsUTF8(PyTuple_GetItem(args, 1)); - TRTC_Add_Built_In_Header(filename, filecontent); - return PyLong_FromLong(0); -} - -static PyObject* n_add_inlcude_filename(PyObject* self, PyObject* args) -{ - const char* fn = PyUnicode_AsUTF8(PyTuple_GetItem(args, 0)); - TRTC_Add_Inlcude_Filename(fn); - return PyLong_FromLong(0); -} - -static PyObject* n_add_code_block(PyObject* self, PyObject* args) -{ - const char* line = PyUnicode_AsUTF8(PyTuple_GetItem(args, 0)); - TRTC_Add_Code_Block(line); - return PyLong_FromLong(0); -} - -static PyObject* n_add_constant_object(PyObject* self, PyObject* args) -{ - const char* name = PyUnicode_AsUTF8(PyTuple_GetItem(args, 0)); - DeviceViewable* dv = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - TRTC_Add_Constant_Object(name, *dv); - return PyLong_FromLong(0); -} - -static PyObject* n_wait(PyObject* self, PyObject* args) -{ - TRTC_Wait(); - return PyLong_FromLong(0); -} - -static PyObject* n_kernel_create(PyObject* self, PyObject* args) -{ - PyObject* pyParamList = PyTuple_GetItem(args, 0); - ssize_t num_params = PyList_Size(pyParamList); - std::vector params(num_params); - for (ssize_t i = 0; i < num_params; i++) - params[i] = PyUnicode_AsUTF8(PyList_GetItem(pyParamList, i)); - const char* body = PyUnicode_AsUTF8(PyTuple_GetItem(args, 1)); - TRTC_Kernel* cptr = new TRTC_Kernel(params, body); - - return PyLong_FromVoidPtr(cptr); -} - -static PyObject* n_kernel_destroy(PyObject* self, PyObject* args) -{ - TRTC_Kernel* cptr = (TRTC_Kernel*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - delete cptr; - return PyLong_FromLong(0); -} - -static PyObject* n_kernel_num_params(PyObject* self, PyObject* args) -{ - TRTC_Kernel* cptr = (TRTC_Kernel*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - return PyLong_FromLong((long)cptr->num_params()); -} - -static PyObject* n_kernel_calc_optimal_block_size(PyObject* self, PyObject* args) -{ - TRTC_Kernel* cptr = (TRTC_Kernel*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - size_t num_params = cptr->num_params(); - - PyObject* arg2 = PyTuple_GetItem(args, 1); - std::vector params; - ssize_t size = PyList_Size(arg2); - if ((ssize_t)num_params != size) - { - PyErr_Format(PyExc_ValueError, "Wrong number of arguments received. %d required, %d received.", num_params, size); - Py_RETURN_NONE; - } - params.resize(size); - for (ssize_t i = 0; i < size; i++) - params[i] = (DeviceViewable*)PyLong_AsVoidPtr(PyList_GetItem(arg2, i)); - - unsigned sharedMemBytes = (unsigned)PyLong_AsUnsignedLong(PyTuple_GetItem(args, 2)); - int sizeBlock; - if (cptr->calc_optimal_block_size(params.data(), sizeBlock, sharedMemBytes)) - return PyLong_FromLong((long)sizeBlock); - else - Py_RETURN_NONE; -} - -static PyObject* n_kernel_calc_number_blocks(PyObject* self, PyObject* args) -{ - TRTC_Kernel* cptr = (TRTC_Kernel*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - size_t num_params = cptr->num_params(); - - PyObject* arg2 = PyTuple_GetItem(args, 1); - std::vector params; - ssize_t size = PyList_Size(arg2); - if ((ssize_t)num_params != size) - { - PyErr_Format(PyExc_ValueError, "Wrong number of arguments received. %d required, %d received.", num_params, size); - Py_RETURN_NONE; - } - params.resize(size); - for (ssize_t i = 0; i < size; i++) - params[i] = (DeviceViewable*)PyLong_AsVoidPtr(PyList_GetItem(arg2, i)); - - int sizeBlock = (int)PyLong_AsLongLong(PyTuple_GetItem(args, 2)); - unsigned sharedMemBytes = (unsigned)PyLong_AsUnsignedLong(PyTuple_GetItem(args, 3)); - int numBlocks; - if (cptr->calc_number_blocks(params.data(), sizeBlock, numBlocks, sharedMemBytes)) - return PyLong_FromLong((long)numBlocks); - else - Py_RETURN_NONE; -} - -static PyObject* n_kernel_launch(PyObject* self, PyObject* args) -{ - TRTC_Kernel* cptr = (TRTC_Kernel*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - size_t num_params = cptr->num_params(); - - dim_type gridDim; - PyObject* arg1 = PyTuple_GetItem(args, 1); - if (PyObject_TypeCheck(arg1, &PyTuple_Type)) - { - ssize_t size = PyTuple_Size(arg1); - gridDim.x = size > 0 ? (unsigned)PyLong_AsUnsignedLong(PyTuple_GetItem(arg1, 0)) : 1; - gridDim.y = size > 1 ? (unsigned)PyLong_AsUnsignedLong(PyTuple_GetItem(arg1, 1)) : 1; - gridDim.z = size > 2 ? (unsigned)PyLong_AsUnsignedLong(PyTuple_GetItem(arg1, 2)) : 1; - } - else - { - gridDim.x = (unsigned)PyLong_AsUnsignedLong(arg1); - gridDim.y = 1; - gridDim.z = 1; - } - - dim_type blockDim; - PyObject* arg2 = PyTuple_GetItem(args, 2); - if (PyObject_TypeCheck(arg2, &PyTuple_Type)) - { - ssize_t size = PyTuple_Size(arg2); - blockDim.x = size > 0 ? (unsigned)PyLong_AsUnsignedLong(PyTuple_GetItem(arg2, 0)) : 1; - blockDim.y = size > 1 ? (unsigned)PyLong_AsUnsignedLong(PyTuple_GetItem(arg2, 1)) : 1; - blockDim.z = size > 2 ? (unsigned)PyLong_AsUnsignedLong(PyTuple_GetItem(arg2, 2)) : 1; - } - else - { - blockDim.x = (unsigned)PyLong_AsUnsignedLong(arg2); - blockDim.y = 1; - blockDim.z = 1; - } - - PyObject* arg3 = PyTuple_GetItem(args, 3); - std::vector params; - ssize_t size = PyList_Size(arg3); - if ((ssize_t)num_params != size) - { - PyErr_Format(PyExc_ValueError, "Wrong number of arguments received. %d required, %d received.", num_params, size); - Py_RETURN_NONE; - } - params.resize(size); - for (ssize_t i = 0; i < size; i++) - params[i] = (DeviceViewable*)PyLong_AsVoidPtr(PyList_GetItem(arg3, i)); - - unsigned sharedMemBytes = (unsigned)PyLong_AsUnsignedLong(PyTuple_GetItem(args, 4)); - if(cptr->launch(gridDim, blockDim, params.data(), sharedMemBytes)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; -} - - -static PyObject* n_for_create(PyObject* self, PyObject* args) -{ - PyObject* pyParamList = PyTuple_GetItem(args, 0); - ssize_t num_params = PyList_Size(pyParamList); - std::vector params(num_params); - for (ssize_t i = 0; i < num_params; i++) - params[i] = PyUnicode_AsUTF8(PyList_GetItem(pyParamList, i)); - const char* name_iter = PyUnicode_AsUTF8(PyTuple_GetItem(args, 1)); - const char* body = PyUnicode_AsUTF8(PyTuple_GetItem(args, 2)); - TRTC_For* cptr = new TRTC_For(params, name_iter, body); - - return PyLong_FromVoidPtr(cptr); -} - -static PyObject* n_for_destroy(PyObject* self, PyObject* args) -{ - TRTC_For* cptr = (TRTC_For*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - delete cptr; - return PyLong_FromLong(0); -} - -static PyObject* n_for_num_params(PyObject* self, PyObject* args) -{ - TRTC_For* cptr = (TRTC_For*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - return PyLong_FromLong((long)cptr->num_params()); -} - -static PyObject* n_for_launch(PyObject* self, PyObject* args) -{ - TRTC_For* cptr = (TRTC_For*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - size_t num_params = cptr->num_params(); - size_t begin = (size_t)PyLong_AsUnsignedLong(PyTuple_GetItem(args, 1)); - size_t end = (size_t)PyLong_AsUnsignedLong(PyTuple_GetItem(args, 2)); - - PyObject* arg3 = PyTuple_GetItem(args, 3); - std::vector params; - ssize_t size = PyList_Size(arg3); - if ((ssize_t)num_params != size) - { - PyErr_Format(PyExc_ValueError, "Wrong number of arguments received. %d required, %d received.", num_params, size); - Py_RETURN_NONE; - } - params.resize(size); - for (ssize_t i = 0; i < size; i++) - params[i] = (DeviceViewable*)PyLong_AsVoidPtr(PyList_GetItem(arg3, i)); - - if(cptr->launch(begin, end, params.data())) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; -} - - -static PyObject* n_for_launch_n(PyObject* self, PyObject* args) -{ - TRTC_For* cptr = (TRTC_For*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - size_t num_params = cptr->num_params(); - size_t n = (size_t)PyLong_AsUnsignedLong(PyTuple_GetItem(args, 1)); - - PyObject* arg3 = PyTuple_GetItem(args, 2); - std::vector params; - ssize_t size = PyList_Size(arg3); - if ((ssize_t)num_params != size) - { - PyErr_Format(PyExc_ValueError, "Wrong number of arguments received. %d required, %d received.", num_params, size); - Py_RETURN_NONE; - } - params.resize(size); - for (ssize_t i = 0; i < size; i++) - params[i] = (DeviceViewable*)PyLong_AsVoidPtr(PyList_GetItem(arg3, i)); - if (cptr->launch_n(n, params.data())) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; -} diff --git a/python/DVConstant.hpp b/python/DVConstant.hpp deleted file mode 100644 index c8ff164..0000000 --- a/python/DVConstant.hpp +++ /dev/null @@ -1,12 +0,0 @@ -#include -#include "TRTCContext.h" -#include "fake_vectors/DVConstant.h" - -static PyObject* n_dvconstant_create(PyObject* self, PyObject* args) -{ - DeviceViewable* dvobj = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - size_t size = (size_t)PyLong_AsLongLong(PyTuple_GetItem(args, 1)); - DVConstant* ret = new DVConstant(*dvobj, size); - return PyLong_FromVoidPtr(ret); -} - diff --git a/python/DVCounter.hpp b/python/DVCounter.hpp deleted file mode 100644 index 4f1807e..0000000 --- a/python/DVCounter.hpp +++ /dev/null @@ -1,12 +0,0 @@ -#include -#include "TRTCContext.h" -#include "fake_vectors/DVCounter.h" - -static PyObject* n_dvcounter_create(PyObject* self, PyObject* args) -{ - DeviceViewable* dvobj_init = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - size_t size = (size_t)PyLong_AsLongLong(PyTuple_GetItem(args, 1)); - DVCounter* ret = new DVCounter(*dvobj_init, size); - return PyLong_FromVoidPtr(ret); -} - diff --git a/python/DVCustomVector.hpp b/python/DVCustomVector.hpp deleted file mode 100644 index 4186ebc..0000000 --- a/python/DVCustomVector.hpp +++ /dev/null @@ -1,23 +0,0 @@ -#include -#include "TRTCContext.h" -#include "fake_vectors/DVCustomVector.h" - -static PyObject* n_dvcustomvector_create(PyObject* self, PyObject* args) -{ - PyObject* pyArgMap = PyTuple_GetItem(args, 0); - ssize_t num_params = PyList_Size(pyArgMap); - std::vector arg_map(num_params); - for (ssize_t i = 0; i < num_params; i++) - { - PyObject* pyCapturedDeviceViewable = PyList_GetItem(pyArgMap, i); - arg_map[i].obj_name = PyUnicode_AsUTF8(PyTuple_GetItem(pyCapturedDeviceViewable, 0)); - arg_map[i].obj = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(pyCapturedDeviceViewable, 1)); - } - const char* name_idx = PyUnicode_AsUTF8(PyTuple_GetItem(args, 1)); - const char* body = PyUnicode_AsUTF8(PyTuple_GetItem(args, 2)); - const char* elem_cls = PyUnicode_AsUTF8(PyTuple_GetItem(args, 3)); - size_t size = (size_t)PyLong_AsLongLong(PyTuple_GetItem(args, 4)); - bool read_only = PyObject_IsTrue(PyTuple_GetItem(args, 5)) != 0; - DVCustomVector* ret = new DVCustomVector(arg_map, name_idx, body, elem_cls, size, read_only); - return PyLong_FromVoidPtr(ret); -} diff --git a/python/DVDiscard.hpp b/python/DVDiscard.hpp deleted file mode 100644 index 96a9b5c..0000000 --- a/python/DVDiscard.hpp +++ /dev/null @@ -1,12 +0,0 @@ -#include -#include "TRTCContext.h" -#include "fake_vectors/DVDiscard.h" - -static PyObject* n_dvdiscard_create(PyObject* self, PyObject* args) -{ - const char* elem_cls = PyUnicode_AsUTF8(PyTuple_GetItem(args, 0)); - size_t size = (size_t)PyLong_AsLongLong(PyTuple_GetItem(args, 1)); - DVDiscard* ret = new DVDiscard(elem_cls, size); - return PyLong_FromVoidPtr(ret); -} - diff --git a/python/DVPermutation.hpp b/python/DVPermutation.hpp deleted file mode 100644 index ea38684..0000000 --- a/python/DVPermutation.hpp +++ /dev/null @@ -1,12 +0,0 @@ -#include -#include "TRTCContext.h" -#include "fake_vectors/DVPermutation.h" - -static PyObject* n_dvpermutation_create(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_value = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_index = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DVPermutation* ret = new DVPermutation(*vec_value, *vec_index); - return PyLong_FromVoidPtr(ret); -} - diff --git a/python/DVRange.hpp b/python/DVRange.hpp deleted file mode 100644 index c98f59e..0000000 --- a/python/DVRange.hpp +++ /dev/null @@ -1,27 +0,0 @@ -#include -#include "TRTCContext.h" -#include "fake_vectors/DVRange.h" - -static PyObject* n_dvrange_create(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_value = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - size_t begin = (size_t)PyLong_AsUnsignedLong(PyTuple_GetItem(args, 1)); - size_t end = (size_t)PyLong_AsUnsignedLong(PyTuple_GetItem(args, 2)); - - DVVector* p_vec = dynamic_cast(vec_value); - if (p_vec) - { - DVVectorAdaptor* ret = new DVVectorAdaptor(*p_vec, begin, end); - return PyLong_FromVoidPtr(ret); - } - - DVVectorAdaptor* p_vec_adpt = dynamic_cast(vec_value); - if (p_vec_adpt) - { - DVVectorAdaptor* ret = new DVVectorAdaptor(*p_vec_adpt, begin, end); - return PyLong_FromVoidPtr(ret); - } - - DVRange* ret = new DVRange(*vec_value, begin, end); - return PyLong_FromVoidPtr(ret); -} diff --git a/python/DVReverse.hpp b/python/DVReverse.hpp deleted file mode 100644 index 37a68a1..0000000 --- a/python/DVReverse.hpp +++ /dev/null @@ -1,11 +0,0 @@ -#include -#include "TRTCContext.h" -#include "fake_vectors/DVReverse.h" - -static PyObject* n_dvreverse_create(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_value = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVReverse* ret = new DVReverse(*vec_value); - return PyLong_FromVoidPtr(ret); -} - diff --git a/python/DVTransform.hpp b/python/DVTransform.hpp deleted file mode 100644 index 60bca21..0000000 --- a/python/DVTransform.hpp +++ /dev/null @@ -1,13 +0,0 @@ -#include -#include "TRTCContext.h" -#include "fake_vectors/DVTransform.h" - -static PyObject* n_dvtransform_create(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - const char* elem_cls = PyUnicode_AsUTF8(PyTuple_GetItem(args, 1)); - Functor* op = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - DVTransform* ret = new DVTransform(*vec_in, elem_cls, *op); - return PyLong_FromVoidPtr(ret); -} - diff --git a/python/DVTuple.hpp b/python/DVTuple.hpp deleted file mode 100644 index 2bff07f..0000000 --- a/python/DVTuple.hpp +++ /dev/null @@ -1,18 +0,0 @@ -#include -#include "TRTCContext.h" -#include "DVTuple.h" - -static PyObject* n_dvtuple_create(PyObject* self, PyObject* args) -{ - PyObject* pyArgMap = PyTuple_GetItem(args, 0); - ssize_t num_params = PyList_Size(pyArgMap); - std::vector arg_map(num_params); - for (ssize_t i = 0; i < num_params; i++) - { - PyObject* pyCapturedDeviceViewable = PyList_GetItem(pyArgMap, i); - arg_map[i].obj_name = PyUnicode_AsUTF8(PyTuple_GetItem(pyCapturedDeviceViewable, 0)); - arg_map[i].obj = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(pyCapturedDeviceViewable, 1)); - } - return PyLong_FromVoidPtr(new DVTuple(arg_map)); -} - diff --git a/python/DVVector.hpp b/python/DVVector.hpp deleted file mode 100644 index a4f7e3b..0000000 --- a/python/DVVector.hpp +++ /dev/null @@ -1,77 +0,0 @@ -#include -#include "TRTCContext.h" -#include "DVVector.h" - -static PyObject* n_dvvectorlike_name_elem_cls(PyObject* self, PyObject* args) -{ - DVVectorLike* dvvec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - return PyUnicode_FromString(dvvec->name_elem_cls().c_str()); -} - -static PyObject* n_dvvectorlike_size(PyObject* self, PyObject* args) -{ - DVVectorLike* dvvec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - return PyLong_FromLongLong((long long)dvvec->size()); -} - -static PyObject* n_dvvector_create(PyObject* self, PyObject* args) -{ - const char* elem_cls = PyUnicode_AsUTF8(PyTuple_GetItem(args, 0)); - size_t size = (size_t)PyLong_AsUnsignedLongLong(PyTuple_GetItem(args, 1)); - PyObject* py_data = PyTuple_GetItem(args, 2); - DVVector* ret = nullptr; - if (py_data == Py_None) - ret = new DVVector(elem_cls, size); - else - { - void* data = PyLong_AsVoidPtr(py_data); - ret = new DVVector(elem_cls, size, data); - } - return PyLong_FromVoidPtr(ret); -} - -static PyObject* n_dvvector_to_host(PyObject* self, PyObject* args) -{ - DVVector* dvvec = (DVVector*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - void* ptr = PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - size_t begin = (size_t)PyLong_AsLong(PyTuple_GetItem(args, 2)); - size_t end = (size_t)PyLong_AsLong(PyTuple_GetItem(args, 3)); - dvvec->to_host(ptr, begin, end); - return PyLong_FromLong(0); -} - -static PyObject* n_dvvector_from_dvs(PyObject* self, PyObject* args) -{ - PyObject* lst = PyTuple_GetItem(args, 0); - ssize_t num_items = PyList_Size(lst); - if (num_items < 1) Py_RETURN_NONE; - PyObject* py_viewable = PyList_GetItem(lst, 0); - DeviceViewable* dv = (DeviceViewable*)PyLong_AsVoidPtr(py_viewable); - std::string elem_cls = dv->name_view_cls(); - for (ssize_t i = 1; i < num_items; i++) - { - py_viewable = PyList_GetItem(lst, i); - dv = (DeviceViewable*)PyLong_AsVoidPtr(py_viewable); - if (elem_cls != dv->name_view_cls()) - Py_RETURN_NONE; - } - size_t elem_size = TRTC_Size_Of(elem_cls.c_str()); - std::vector buf(elem_size*num_items); - for (ssize_t i = 0; i < num_items; i++) - { - py_viewable = PyList_GetItem(lst, i); - dv = (DeviceViewable*)PyLong_AsVoidPtr(py_viewable); - memcpy(buf.data() + elem_size*i, dv->view().data(), elem_size); - } - DVVector* ret = new DVVector(elem_cls.data(), num_items, buf.data()); - return PyLong_FromVoidPtr(ret); -} - -static PyObject* n_dvvectoradaptor_create(PyObject* self, PyObject* args) -{ - const char* elem_cls = PyUnicode_AsUTF8(PyTuple_GetItem(args, 0)); - size_t size = (size_t)PyLong_AsUnsignedLongLong(PyTuple_GetItem(args, 1)); - void* data = PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - DVVectorAdaptor* ret = new DVVectorAdaptor(elem_cls, size, data); - return PyLong_FromVoidPtr(ret); -} \ No newline at end of file diff --git a/python/DVZipped.hpp b/python/DVZipped.hpp deleted file mode 100644 index 0888765..0000000 --- a/python/DVZipped.hpp +++ /dev/null @@ -1,26 +0,0 @@ -#include -#include "TRTCContext.h" -#include "fake_vectors/DVZipped.h" - -static PyObject* n_dvzipped_create(PyObject* self, PyObject* args) -{ - PyObject* pyvecs = PyTuple_GetItem(args, 0); - ssize_t num_vecs = PyList_Size(pyvecs); - std::vector vecs(num_vecs); - for (ssize_t i = 0; i < num_vecs; i++) - vecs[i] = (DVVectorLike*)PyLong_AsVoidPtr(PyList_GetItem(pyvecs, i)); - - PyObject* py_elem_names = PyTuple_GetItem(args, 1); - ssize_t num_elems = PyList_Size(py_elem_names); - if (num_elems!= num_vecs) - { - PyErr_Format(PyExc_ValueError, "Number of vectors %d mismatch with number of element names %d.", num_vecs, num_elems); - Py_RETURN_NONE; - } - std::vector elem_names(num_elems); - for (ssize_t i = 0; i < num_elems; i++) - elem_names[i] = PyUnicode_AsUTF8(PyList_GetItem(py_elem_names, i)); - return PyLong_FromVoidPtr(new DVZipped(vecs, elem_names)); - -} - diff --git a/python/DeviceViewable.hpp b/python/DeviceViewable.hpp deleted file mode 100644 index a813220..0000000 --- a/python/DeviceViewable.hpp +++ /dev/null @@ -1,54 +0,0 @@ -#include -#include "TRTCContext.h" -#include "DeviceViewable.h" -#include "viewbuf_to_python.hpp" - -static PyObject* n_dv_name_view_cls(PyObject* self, PyObject* args) -{ - DeviceViewable* dv = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - return PyUnicode_FromString(dv->name_view_cls().c_str()); -} - -static PyObject* n_dv_destroy(PyObject* self, PyObject* args) -{ - DeviceViewable* dv = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - delete dv; - return PyLong_FromLong(0); -} - -static PyObject* n_dv_create_basic(PyObject* self, PyObject* args) -{ - std::string type = PyUnicode_AsUTF8(PyTuple_GetItem(args, 0)); - PyObject* value = PyTuple_GetItem(args, 1); - DeviceViewable* ret = nullptr; - if (type == "int8_t") - ret = new DVInt8((int8_t)PyLong_AsLong(value)); - else if (type == "uint8_t") - ret = new DVUInt8((uint8_t)PyLong_AsLong(value)); - else if (type == "int16_t") - ret = new DVInt16((int16_t)PyLong_AsLong(value)); - else if (type == "uint16_t") - ret = new DVUInt16((uint16_t)PyLong_AsUnsignedLong(value)); - else if (type == "int32_t") - ret = new DVInt32((int32_t)PyLong_AsLong(value)); - else if (type == "uint32_t") - ret = new DVUInt32((uint32_t)PyLong_AsUnsignedLong(value)); - else if (type == "int64_t") - ret = new DVInt64((int64_t)PyLong_AsLong(value)); - else if (type == "uint64_t") - ret = new DVUInt64((uint64_t)PyLong_AsUnsignedLong(value)); - else if (type == "float") - ret = new DVFloat((float)PyFloat_AsDouble(value)); - else if (type == "double") - ret = new DVDouble((double)PyFloat_AsDouble(value)); - else if (type == "bool") - ret = new DVBool(PyObject_IsTrue(value) != 0); - return PyLong_FromVoidPtr(ret); -} - -static PyObject* n_dv_value(PyObject* self, PyObject* args) -{ - DeviceViewable* dv = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - return PyValue_FromViewBuf(dv->view(), dv->name_view_cls().c_str()); -} - diff --git a/python/PyThrustRTC.cpp b/python/PyThrustRTC.cpp deleted file mode 100644 index f521acd..0000000 --- a/python/PyThrustRTC.cpp +++ /dev/null @@ -1,177 +0,0 @@ -#include -#include "TRTCContext.h" - -static PyObject* n_set_libnvrtc_path(PyObject* self, PyObject* args) -{ - const char* path = PyUnicode_AsUTF8(PyTuple_GetItem(args, 0)); - set_libnvrtc_path(path); - return PyLong_FromLong(0); -} - -#include "Context.hpp" -#include "DeviceViewable.hpp" -#include "DVVector.hpp" -#include "DVTuple.hpp" -#include "DVRange.hpp" -#include "DVConstant.hpp" -#include "DVCounter.hpp" -#include "DVDiscard.hpp" -#include "DVPermutation.hpp" -#include "DVReverse.hpp" -#include "DVTransform.hpp" -#include "DVZipped.hpp" -#include "DVCustomVector.hpp" -#include "functor.hpp" -#include "fill.hpp" -#include "replace.hpp" -#include "for_each.hpp" -#include "adjacent_difference.hpp" -#include "sequence.hpp" -#include "tabulate.hpp" -#include "transform.hpp" -#include "gather.hpp" -#include "scatter.hpp" -#include "copy.hpp" -#include "swap.hpp" -#include "count.hpp" -#include "reduce.hpp" -#include "equal.hpp" -#include "extrema.hpp" -#include "inner_product.hpp" -#include "transform_reduce.hpp" -#include "logical.hpp" -#include "scan.hpp" -#include "transform_scan.hpp" -#include "remove.hpp" -#include "unique.hpp" -#include "partition.hpp" -#include "find.hpp" -#include "mismatch.hpp" -#include "binary_search.hpp" -#include "merge.hpp" -#include "sort.hpp" - -static PyMethodDef s_Methods[] = { - { "n_set_libnvrtc_path", n_set_libnvrtc_path, METH_VARARGS, "" }, - { "n_set_verbose", n_set_verbose, METH_VARARGS, "" }, - { "n_add_include_dir", n_add_include_dir, METH_VARARGS, "" }, - { "n_add_built_in_header", n_add_built_in_header, METH_VARARGS, "" }, - { "n_add_inlcude_filename", n_add_inlcude_filename, METH_VARARGS, "" }, - { "n_add_code_block", n_add_code_block, METH_VARARGS, "" }, - { "n_add_constant_object", n_add_constant_object, METH_VARARGS, "" }, - { "n_wait", n_wait, METH_VARARGS, "" }, - { "n_kernel_create", n_kernel_create, METH_VARARGS, "" }, - { "n_kernel_destroy", n_kernel_destroy, METH_VARARGS, "" }, - { "n_kernel_num_params", n_kernel_num_params, METH_VARARGS, "" }, - { "n_kernel_calc_optimal_block_size", n_kernel_calc_optimal_block_size, METH_VARARGS, "" }, - { "n_kernel_calc_number_blocks", n_kernel_calc_number_blocks, METH_VARARGS, "" }, - { "n_kernel_launch", n_kernel_launch, METH_VARARGS, "" }, - { "n_for_create", n_for_create, METH_VARARGS, "" }, - { "n_for_destroy", n_for_destroy, METH_VARARGS, "" }, - { "n_for_num_params", n_for_num_params, METH_VARARGS, "" }, - { "n_for_launch", n_for_launch, METH_VARARGS, "" }, - { "n_for_launch_n", n_for_launch_n, METH_VARARGS, "" }, - { "n_dv_name_view_cls", n_dv_name_view_cls, METH_VARARGS, "" }, - { "n_dv_destroy", n_dv_destroy, METH_VARARGS, "" }, - { "n_dv_create_basic", n_dv_create_basic, METH_VARARGS, "" }, - { "n_dv_value", n_dv_value, METH_VARARGS, "" }, - { "n_dvvectorlike_name_elem_cls", n_dvvectorlike_name_elem_cls, METH_VARARGS, "" }, - { "n_dvvectorlike_size", n_dvvectorlike_size, METH_VARARGS, "" }, - { "n_dvvector_create", n_dvvector_create, METH_VARARGS, "" }, - { "n_dvvector_to_host", n_dvvector_to_host, METH_VARARGS, "" }, - { "n_dvvector_from_dvs", n_dvvector_from_dvs, METH_VARARGS, "" }, - { "n_dvvectoradaptor_create", n_dvvectoradaptor_create, METH_VARARGS, "" }, - { "n_dvrange_create", n_dvrange_create, METH_VARARGS, "" }, - { "n_dvtuple_create", n_dvtuple_create, METH_VARARGS, "" }, - { "n_dvconstant_create", n_dvconstant_create, METH_VARARGS, "" }, - { "n_dvcounter_create", n_dvcounter_create, METH_VARARGS, "" }, - { "n_dvdiscard_create", n_dvdiscard_create, METH_VARARGS, "" }, - { "n_dvpermutation_create", n_dvpermutation_create, METH_VARARGS, "" }, - { "n_dvreverse_create", n_dvreverse_create, METH_VARARGS, "" }, - { "n_dvtransform_create", n_dvtransform_create, METH_VARARGS, "" }, - { "n_dvzipped_create", n_dvzipped_create, METH_VARARGS, "" }, - { "n_dvcustomvector_create", n_dvcustomvector_create, METH_VARARGS, "" }, - { "n_functor_create", n_functor_create, METH_VARARGS, "" }, - { "n_built_in_functor_create", n_built_in_functor_create, METH_VARARGS, "" }, - { "n_fill", n_fill, METH_VARARGS, "" }, - { "n_replace", n_replace, METH_VARARGS, "" }, - { "n_replace_if", n_replace_if, METH_VARARGS, "" }, - { "n_replace_copy", n_replace_copy, METH_VARARGS, "" }, - { "n_replace_copy_if", n_replace_copy_if, METH_VARARGS, "" }, - { "n_for_each", n_for_each, METH_VARARGS, "" }, - { "n_adjacent_difference", n_adjacent_difference, METH_VARARGS, "" }, - { "n_sequence", n_sequence, METH_VARARGS, "" }, - { "n_tabulate", n_tabulate, METH_VARARGS, "" }, - { "n_transform", n_transform, METH_VARARGS, "" }, - { "n_transform_binary", n_transform_binary, METH_VARARGS, "" }, - { "n_transform_if", n_transform_if, METH_VARARGS, "" }, - { "n_transform_if_stencil", n_transform_if_stencil, METH_VARARGS, "" }, - { "n_transform_binary_if_stencil", n_transform_binary_if_stencil, METH_VARARGS, "" }, - { "n_gather", n_gather, METH_VARARGS, "" }, - { "n_gather_if", n_gather_if, METH_VARARGS, "" }, - { "n_scatter", n_scatter, METH_VARARGS, "" }, - { "n_scatter_if", n_scatter_if, METH_VARARGS, "" }, - { "n_copy", n_copy, METH_VARARGS, "" }, - { "n_swap", n_swap, METH_VARARGS, "" }, - { "n_count", n_count, METH_VARARGS, "" }, - { "n_count_if", n_count_if, METH_VARARGS, "" }, - { "n_reduce", n_reduce, METH_VARARGS, "" }, - { "n_equal", n_equal, METH_VARARGS, "" }, - { "n_min_element", n_min_element, METH_VARARGS, "" }, - { "n_max_element", n_max_element, METH_VARARGS, "" }, - { "n_minmax_element", n_minmax_element, METH_VARARGS, "" }, - { "n_inner_product", n_inner_product, METH_VARARGS, "" }, - { "n_transform_reduce", n_transform_reduce, METH_VARARGS, "" }, - { "n_all_of", n_all_of, METH_VARARGS, "" }, - { "n_any_of", n_any_of, METH_VARARGS, "" }, - { "n_none_of", n_none_of, METH_VARARGS, "" }, - { "n_inclusive_scan", n_inclusive_scan, METH_VARARGS, "" }, - { "n_exclusive_scan", n_exclusive_scan, METH_VARARGS, "" }, - { "n_transform_inclusive_scan", n_transform_inclusive_scan, METH_VARARGS, "" }, - { "n_transform_exclusive_scan", n_transform_exclusive_scan, METH_VARARGS, "" }, - { "n_inclusive_scan_by_key", n_inclusive_scan_by_key, METH_VARARGS, "" }, - { "n_exclusive_scan_by_key", n_exclusive_scan_by_key, METH_VARARGS, "" }, - { "n_copy_if", n_copy_if, METH_VARARGS, "" }, - { "n_copy_if_stencil", n_copy_if_stencil, METH_VARARGS, "" }, - { "n_remove", n_remove, METH_VARARGS, "" }, - { "n_remove_copy", n_remove_copy, METH_VARARGS, "" }, - { "n_remove_if", n_remove_if, METH_VARARGS, "" }, - { "n_remove_copy_if", n_remove_copy_if, METH_VARARGS, "" }, - { "n_remove_if_stencil", n_remove_if_stencil, METH_VARARGS, "" }, - { "n_remove_copy_if_stencil", n_remove_copy_if_stencil, METH_VARARGS, "" }, - { "n_unique", n_unique, METH_VARARGS, "" }, - { "n_unique_copy", n_unique_copy, METH_VARARGS, "" }, - { "n_unique_by_key", n_unique_by_key, METH_VARARGS, "" }, - { "n_unique_by_key_copy", n_unique_by_key_copy, METH_VARARGS, "" }, - { "n_partition", n_partition, METH_VARARGS, "" }, - { "n_partition_stencil", n_partition_stencil, METH_VARARGS, "" }, - { "n_partition_copy", n_partition_copy, METH_VARARGS, "" }, - { "n_partition_copy_stencil", n_partition_copy_stencil, METH_VARARGS, "" }, - { "n_find", n_find, METH_VARARGS, "" }, - { "n_find_if", n_find_if, METH_VARARGS, "" }, - { "n_find_if_not", n_find_if_not, METH_VARARGS, "" }, - { "n_mismatch", n_mismatch, METH_VARARGS, "" }, - { "n_lower_bound", n_lower_bound, METH_VARARGS, "" }, - { "n_upper_bound", n_upper_bound, METH_VARARGS, "" }, - { "n_binary_search", n_binary_search, METH_VARARGS, "" }, - { "n_partition_point", n_partition_point, METH_VARARGS, "" }, - { "n_reduce_by_key", n_reduce_by_key, METH_VARARGS, "" }, - { "n_merge", n_merge, METH_VARARGS, "" }, - { "n_merge_by_key", n_merge_by_key, METH_VARARGS, "" }, - { "n_sort", n_sort, METH_VARARGS, "" }, - { "n_sort_by_key", n_sort_by_key, METH_VARARGS, "" }, - { "n_is_partitioned", n_is_partitioned, METH_VARARGS, "" }, - { "n_is_sorted", n_is_sorted, METH_VARARGS, "" }, - { "n_is_sorted_until", n_is_sorted_until, METH_VARARGS, "" }, - { "n_lower_bound_v", n_lower_bound_v, METH_VARARGS, "" }, - { "n_upper_bound_v", n_upper_bound_v, METH_VARARGS, "" }, - { "n_binary_search_v", n_binary_search_v, METH_VARARGS, "" }, - 0 -}; - -static struct PyModuleDef cModPyDem = { PyModuleDef_HEAD_INIT, "ThrustRTC_module", "", -1, s_Methods }; - -PyMODINIT_FUNC PyInit_PyThrustRTC(void) -{ - return PyModule_Create(&cModPyDem); -} diff --git a/python/ThrustRTC/Context.py b/python/ThrustRTC/Context.py index 9bc0897..55e3297 100644 --- a/python/ThrustRTC/Context.py +++ b/python/ThrustRTC/Context.py @@ -1,29 +1,34 @@ -import PyThrustRTC as native +from .Native import native +from .utils import * + +def set_libnvrtc_path(path): + native.n_set_libnvrtc_path(path.encode('utf-8')) def Set_Verbose(verbose=True): native.n_set_verbose(verbose) def Add_Include_Dir(path): - native.n_add_include_dir(path) + native.n_add_include_dir(path.encode('utf-8')) def Add_Built_In_Header(filename, filecontent): - native.n_add_built_in_header(filename, filecontent) + native.n_add_built_in_header(filename.encode('utf-8'), filecontent.encode('utf-8')) def Add_Inlcude_Filename(filename): - native.n_add_inlcude_filename(filename) + native.n_add_inlcude_filename(filename.encode('utf-8')) def Add_Code_Block(code): - native.n_add_code_block(code) + native.n_add_code_block(code.encode('utf-8')) def Add_Constant_Object(name, dv): - native.n_add_constant_object(name, dv.m_cptr) + native.n_add_constant_object(name.encode('utf-8'), dv.m_cptr) def Wait(): native.n_wait() class Kernel: def __init__(self, param_names, body): - self.m_cptr = native.n_kernel_create(param_names, body) + o_param_names = StrArray(param_names) + self.m_cptr = native.n_kernel_create(o_param_names.m_cptr, body.encode('utf-8')) def __del__(self): native.n_kernel_destroy(self.m_cptr) @@ -32,29 +37,34 @@ def num_params(self): return native.n_kernel_num_params(self.m_cptr) def calc_optimal_block_size(self, args, sharedMemBytes=0): + arg_list = ObjArray(args) return native.n_kernel_calc_optimal_block_size( - self.m_cptr, - [item.m_cptr for item in args], - sharedMemBytes) + self.m_cptr, arg_list.m_cptr, sharedMemBytes) def calc_number_blocks(self, args, size_block, sharedMemBytes=0): + arg_list = ObjArray(args) return native.n_kernel_calc_number_blocks( self.m_cptr, - [item.m_cptr for item in args], + arg_list.m_cptr, size_block, sharedMemBytes) def launch(self, gridDim, blockDim, args, sharedMemBytes=0): + d_gridDim = Dim3(gridDim) + d_blockDim = Dim3(blockDim) + arg_list = ObjArray(args) native.n_kernel_launch( self.m_cptr, - gridDim, - blockDim, - [item.m_cptr for item in args], + d_gridDim.m_cptr, + d_blockDim.m_cptr, + arg_list.m_cptr, sharedMemBytes) + class For: - def __init__(self, param_descs, name_iter, body): - self.m_cptr = native.n_for_create(param_descs, name_iter, body) + def __init__(self, param_names, name_iter, body): + o_param_names = StrArray(param_names) + self.m_cptr = native.n_for_create(o_param_names.m_cptr, name_iter.encode('utf-8'), body.encode('utf-8')) def __del__(self): native.n_for_destroy(self.m_cptr) @@ -63,7 +73,9 @@ def num_params(self): return native.n_for_num_params(self.m_cptr) def launch(self, begin, end, args): - native.n_for_launch(self.m_cptr, begin, end, [item.m_cptr for item in args]) + arg_list = ObjArray(args) + native.n_for_launch(self.m_cptr, begin, end, arg_list.m_cptr) def launch_n(self, n, args): - native.n_for_launch_n(self.m_cptr, n, [item.m_cptr for item in args]) + arg_list = ObjArray(args) + native.n_for_launch_n(self.m_cptr, n, arg_list.m_cptr) diff --git a/python/ThrustRTC/Copying.py b/python/ThrustRTC/Copying.py index 4baeef0..96012dd 100644 --- a/python/ThrustRTC/Copying.py +++ b/python/ThrustRTC/Copying.py @@ -1,10 +1,10 @@ -import PyThrustRTC as native +from .Native import ffi, native def Gather(vec_map, vec_in, vec_out): native.n_gather(vec_map.m_cptr, vec_in.m_cptr, vec_out.m_cptr) def Gather_If(vec_map, vec_stencil, vec_in, vec_out, pred = None): - cptr_pred = None + cptr_pred = ffi.NULL if pred!=None: cptr_pred = pred.m_cptr native.n_gather_if(vec_map.m_cptr, vec_stencil.m_cptr, vec_in.m_cptr, vec_out.m_cptr, cptr_pred) @@ -13,7 +13,7 @@ def Scatter(vec_in, vec_map, vec_out): native.n_scatter(vec_in.m_cptr, vec_map.m_cptr, vec_out.m_cptr) def Scatter_If(vec_in, vec_map, vec_stencil, vec_out, pred = None): - cptr_pred = None + cptr_pred = ffi.NULL if pred!=None: cptr_pred = pred.m_cptr native.n_scatter_if(vec_in.m_cptr, vec_map.m_cptr, vec_stencil.m_cptr, vec_out.m_cptr, cptr_pred) @@ -22,4 +22,5 @@ def Copy(vec_in, vec_out): native.n_copy(vec_in.m_cptr, vec_out.m_cptr) def Swap(vec1, vec2): - native.n_swap(vec1.m_cptr, vec2.m_cptr) + native.n_swap(vec1.m_cptr, vec2.m_cptr) + diff --git a/python/ThrustRTC/DVTuple.py b/python/ThrustRTC/DVTuple.py index b098a23..d5273db 100644 --- a/python/ThrustRTC/DVTuple.py +++ b/python/ThrustRTC/DVTuple.py @@ -1,8 +1,13 @@ -import PyThrustRTC as native +from .Native import ffi, native from .DeviceViewable import DeviceViewable +from .utils import * class DVTuple(DeviceViewable): def __init__(self, elem_map): self.m_elem_map = elem_map - self.m_cptr = native.n_dvtuple_create([ (param_name, elem.m_cptr) for param_name, elem in elem_map.items()]) + param_names = [param_name for param_name, elem in elem_map.items()] + o_param_names = StrArray(param_names) + elems = [elem for param_name, elem in elem_map.items()] + o_elems = ObjArray(elems) + self.m_cptr = native.n_dvtuple_create(o_elems.m_cptr, o_param_names.m_cptr) diff --git a/python/ThrustRTC/DVVector.py b/python/ThrustRTC/DVVector.py index 079e2b6..282096a 100644 --- a/python/ThrustRTC/DVVector.py +++ b/python/ThrustRTC/DVVector.py @@ -1,10 +1,11 @@ -import PyThrustRTC as native +from .Native import ffi, native import numpy as np from .DeviceViewable import DeviceViewable +from .utils import * class DVVectorLike(DeviceViewable): def name_elem_cls(self): - return native.n_dvvectorlike_name_elem_cls(self.m_cptr) + return ffi.string(native.n_dvvectorlike_name_elem_cls(self.m_cptr)).decode('utf-8') def size(self): return native.n_dvvectorlike_size(self.m_cptr) @@ -17,7 +18,6 @@ def __init__(self, src, begin = 0, end = -1): self.m_src = src self.m_cptr = native.n_dvrange_create(src.m_cptr, begin, end) - class DVVector(DVVectorLike): def __init__(self, cptr): self.m_cptr = cptr @@ -49,11 +49,14 @@ def to_host(self, begin = 0, end = -1): if end == -1: end = self.size() ret = np.empty(end - begin, dtype=nptype) - native.n_dvvector_to_host(self.m_cptr, ret.__array_interface__['data'][0], begin, end) + native.n_dvvector_to_host(self.m_cptr, ffi.cast("void *", ret.__array_interface__['data'][0]), begin, end) return ret def device_vector(elem_cls, size, ptr_host_data=None): - return DVVector(native.n_dvvector_create(elem_cls, size, ptr_host_data)) + ffiptr = ffi.NULL + if ptr_host_data!=None: + ffiptr = ffi.cast("void *", ptr_host_data) + return DVVector(native.n_dvvector_create(elem_cls.encode('utf-8'), size, ffiptr)) def device_vector_from_numpy(nparr): if nparr.dtype == np.int8: @@ -115,7 +118,8 @@ def device_vector_from_list(lst, elem_cls): return device_vector(elem_cls, size, ptr_host_data) def device_vector_from_dvs(lst_dv): - return DVVector(native.n_dvvector_from_dvs([item.m_cptr for item in lst_dv])) + dvarr = ObjArray(lst_dv) + return DVVector(native.n_dvvector_from_dvs(dvarr.m_cptr)) class DVNumbaVector(DVVectorLike): def __init__(self, nbarr): @@ -148,5 +152,5 @@ def __init__(self, nbarr): elem_cls = 'bool' size = nbarr.size ptr_device_data = nbarr.device_ctypes_pointer.value - self.m_cptr = native.n_dvvectoradaptor_create(elem_cls, size, ptr_device_data) + self.m_cptr = native.n_dvvectoradaptor_create(elem_cls.encode('utf-8'), size, ffi.cast("void *", ptr_device_data)) diff --git a/python/ThrustRTC/DeviceViewable.py b/python/ThrustRTC/DeviceViewable.py index 90d8dd4..13550ab 100644 --- a/python/ThrustRTC/DeviceViewable.py +++ b/python/ThrustRTC/DeviceViewable.py @@ -1,53 +1,105 @@ -import PyThrustRTC as native +from .Native import ffi, native + +def ValueFromDVPtr(cptr): + name_view_cls = ffi.string(native.n_dv_name_view_cls(cptr)).decode('utf-8') + if name_view_cls=='int8_t': + return native.n_dvint8_value(cptr) + elif name_view_cls=='uint8_t': + return native.n_dvuint8_value(cptr) + elif name_view_cls=='int16_t': + return native.n_dvint16_value(cptr) + elif name_view_cls=='uint16_t': + return native.n_dvuint16_value(cptr) + elif name_view_cls=='int32_t': + return native.n_dvint32_value(cptr) + elif name_view_cls=='uint32_t': + return native.n_dvuint32_value(cptr) + elif name_view_cls=='int64_t': + return native.n_dvuint64_value(cptr) + elif name_view_cls=='uint64_t': + return native.n_dvuint64_value(cptr) + elif name_view_cls=='float': + return native.n_dvfloat_value(cptr) + elif name_view_cls=='double': + return native.n_dvdouble_value(cptr) + elif name_view_cls=='bool': + return native.n_dvbool_value(cptr)!=0 + else: + return '[Device-viewable object, type: %s]'%name_view_cls class DeviceViewable: def name_view_cls(self): - return native.n_dv_name_view_cls(self.m_cptr) + return ffi.string(native.n_dv_name_view_cls(self.m_cptr)).decode('utf-8') def __del__(self): native.n_dv_destroy(self.m_cptr) def value(self): - return native.n_dv_value(self.m_cptr) + s_type = self.name_view_cls() + return '[Device-viewable object, type: %s]'%s_type class DVInt8(DeviceViewable): def __init__(self, value): - self.m_cptr = native.n_dv_create_basic('int8_t', value) + self.m_cptr = native.n_dvint8_create(value) + def value(self): + return native.n_dvint8_value(self.m_cptr) class DVUInt8(DeviceViewable): def __init__(self, value): - self.m_cptr = native.n_dv_create_basic('uint8_t', value) + self.m_cptr = native.n_dvuint8_create(value) + def value(self): + return native.n_dvuint8_value(self.m_cptr) class DVInt16(DeviceViewable): def __init__(self, value): - self.m_cptr = native.n_dv_create_basic('int16_t', value) + self.m_cptr = native.n_dvint16_create(value) + def value(self): + return native.n_dvint16_value(self.m_cptr) class DVUInt16(DeviceViewable): def __init__(self, value): - self.m_cptr = native.n_dv_create_basic('uint16_t', value) + self.m_cptr = native.n_dvuint16_create(value) + def value(self): + return native.n_dvuint16_value(self.m_cptr) class DVInt32(DeviceViewable): def __init__(self, value): - self.m_cptr = native.n_dv_create_basic('int32_t', value) + self.m_cptr = native.n_dvint32_create(value) + def value(self): + return native.n_dvint32_value(self.m_cptr) class DVUInt32(DeviceViewable): def __init__(self, value): - self.m_cptr = native.n_dv_create_basic('uint32_t', value) + self.m_cptr = native.n_dvuint32_create(value) + def value(self): + return native.n_dvuint32_value(self.m_cptr) class DVInt64(DeviceViewable): def __init__(self, value): - self.m_cptr = native.n_dv_create_basic('int64_t', value) + self.m_cptr = native.n_dvint64_create(value) + def value(self): + return native.n_dvint64_value(self.m_cptr) class DVUInt64(DeviceViewable): def __init__(self, value): - self.m_cptr = native.n_dv_create_basic('uint64_t', value) + self.m_cptr = native.n_dvuint64_create(value) + def value(self): + return native.n_dvuint64_value(self.m_cptr) class DVFloat(DeviceViewable): def __init__(self, value): - self.m_cptr = native.n_dv_create_basic('float', value) + self.m_cptr = native.n_dvfloat_create(value) + def value(self): + return native.n_dvfloat_value(self.m_cptr) class DVDouble(DeviceViewable): def __init__(self, value): - self.m_cptr = native.n_dv_create_basic('double', value) + self.m_cptr = native.n_dvdouble_create(value) + def value(self): + return native.n_dvdouble_value(self.m_cptr) class DVBool(DeviceViewable): def __init__(self, value): - self.m_cptr = native.n_dv_create_basic('bool', value) \ No newline at end of file + self.m_cptr = native.n_dvbool_create(value) + def value(self): + return native.n_dvbool_value(self.m_cptr)!=0 + + diff --git a/python/ThrustRTC/FakeVectors.py b/python/ThrustRTC/FakeVectors.py index 7ff2125..adcf224 100644 --- a/python/ThrustRTC/FakeVectors.py +++ b/python/ThrustRTC/FakeVectors.py @@ -1,19 +1,21 @@ -import PyThrustRTC as native +from .Native import native from .DVVector import DVVectorLike +from .utils import * +import ctypes class DVConstant(DVVectorLike): def __init__(self, dvobj, size = -1): self.m_dvobj = dvobj - self.m_cptr = native.n_dvconstant_create(dvobj.m_cptr, size) + self.m_cptr = native.n_dvconstant_create(dvobj.m_cptr, ctypes.c_ulonglong(size).value) class DVCounter(DVVectorLike): def __init__(self, dvobj_init, size = -1): self.m_dvobj_init = dvobj_init - self.m_cptr = native.n_dvcounter_create(dvobj_init.m_cptr, size) + self.m_cptr = native.n_dvcounter_create(dvobj_init.m_cptr, ctypes.c_ulonglong(size).value) class DVDiscard(DVVectorLike): def __init__(self, elem_cls, size = -1): - self.m_cptr = native.n_dvdiscard_create(elem_cls, size) + self.m_cptr = native.n_dvdiscard_create(elem_cls.encode('utf-8'), ctypes.c_ulonglong(size).value) class DVPermutation(DVVectorLike): def __init__(self, vec_value, vec_index): @@ -30,17 +32,27 @@ class DVTransform(DVVectorLike): def __init__(self, vec_in, elem_cls, op): self.m_vec_in = vec_in self.m_op = op - self.m_cptr = native.n_dvtransform_create(vec_in.m_cptr, elem_cls, op.m_cptr) + self.m_cptr = native.n_dvtransform_create(vec_in.m_cptr, elem_cls.encode('utf-8'), op.m_cptr) class DVZipped(DVVectorLike): def __init__(self, vecs, elem_names): self.m_vecs = vecs - self.m_cptr = native.n_dvzipped_create([item.m_cptr for item in vecs], elem_names) + o_vecs = ObjArray(vecs) + o_elem_names = StrArray(elem_names) + self.m_cptr = native.n_dvzipped_create(o_vecs.m_cptr, o_elem_names.m_cptr) class DVCustomVector(DVVectorLike): def __init__(self, arg_map, name_idx, code_body, elem_cls, size = -1, read_only = True): self.m_arg_map = arg_map - self.m_cptr = native.n_dvcustomvector_create( - [ (param_name, arg.m_cptr) for param_name, arg in arg_map.items()], - name_idx, code_body, elem_cls, size, read_only) + param_names = [param_name for param_name, elem in arg_map.items()] + o_param_names = StrArray(param_names) + elems = [elem for param_name, elem in arg_map.items()] + o_elems = ObjArray(elems) + + self.m_cptr = native.n_dvcustomvector_create( + o_elems.m_cptr, o_param_names.m_cptr, + name_idx.encode('utf-8'), code_body.encode('utf-8'), + elem_cls.encode('utf-8'), ctypes.c_ulonglong(size).value, read_only) + + diff --git a/python/ThrustRTC/Functor.py b/python/ThrustRTC/Functor.py index 42848e5..1db48c7 100644 --- a/python/ThrustRTC/Functor.py +++ b/python/ThrustRTC/Functor.py @@ -1,70 +1,78 @@ -import PyThrustRTC as native +from .Native import ffi, native from .DeviceViewable import DeviceViewable +from .utils import * class Functor(DeviceViewable): def __init__(self, arg_map, functor_params, code_body): self.m_arg_map = arg_map + param_names = [param_name for param_name, elem in arg_map.items()] + o_param_names = StrArray(param_names) + elems = [elem for param_name, elem in arg_map.items()] + o_elems = ObjArray(elems) + o_functor_params = StrArray(functor_params) + self.m_cptr = native.n_functor_create( - [ (param_name, arg.m_cptr) for param_name, arg in arg_map.items()], - functor_params, - code_body) + o_elems.m_cptr, o_param_names.m_cptr, + o_functor_params.m_cptr, + code_body.encode('utf-8')) class Identity(DeviceViewable): def __init__(self): - self.m_cptr = native.n_built_in_functor_create("Identity") + self.m_cptr = native.n_built_in_functor_create(b"Identity") class Maximum(DeviceViewable): def __init__(self): - self.m_cptr = native.n_built_in_functor_create("Maximum") + self.m_cptr = native.n_built_in_functor_create(b"Maximum") class Minimum(DeviceViewable): def __init__(self): - self.m_cptr = native.n_built_in_functor_create("Minimum") + self.m_cptr = native.n_built_in_functor_create(b"Minimum") + class EqualTo(DeviceViewable): def __init__(self): - self.m_cptr = native.n_built_in_functor_create("EqualTo") + self.m_cptr = native.n_built_in_functor_create(b"EqualTo") class NotEqualTo(DeviceViewable): def __init__(self): - self.m_cptr = native.n_built_in_functor_create("NotEqualTo") + self.m_cptr = native.n_built_in_functor_create(b"NotEqualTo") class Greater(DeviceViewable): def __init__(self): - self.m_cptr = native.n_built_in_functor_create("Greater") + self.m_cptr = native.n_built_in_functor_create(b"Greater") class Less(DeviceViewable): def __init__(self): - self.m_cptr = native.n_built_in_functor_create("Less") + self.m_cptr = native.n_built_in_functor_create(b"Less") class GreaterEqual(DeviceViewable): def __init__(self): - self.m_cptr = native.n_built_in_functor_create("GreaterEqual") + self.m_cptr = native.n_built_in_functor_create(b"GreaterEqual") class LessEqual(DeviceViewable): def __init__(self): - self.m_cptr = native.n_built_in_functor_create("LessEqual") + self.m_cptr = native.n_built_in_functor_create(b"LessEqual") class Plus(DeviceViewable): def __init__(self): - self.m_cptr = native.n_built_in_functor_create("Plus") + self.m_cptr = native.n_built_in_functor_create(b"Plus") class Minus(DeviceViewable): def __init__(self): - self.m_cptr = native.n_built_in_functor_create("Minus") + self.m_cptr = native.n_built_in_functor_create(b"Minus") class Multiplies(DeviceViewable): def __init__(self): - self.m_cptr = native.n_built_in_functor_create("Multiplies") + self.m_cptr = native.n_built_in_functor_create(b"Multiplies") class Divides(DeviceViewable): def __init__(self): - self.m_cptr = native.n_built_in_functor_create("Divides") + self.m_cptr = native.n_built_in_functor_create(b"Divides") class Modulus(DeviceViewable): def __init__(self): - self.m_cptr = native.n_built_in_functor_create("Modulus") + self.m_cptr = native.n_built_in_functor_create(b"Modulus") class Negate(DeviceViewable): def __init__(self): - self.m_cptr = native.n_built_in_functor_create("Negate") + self.m_cptr = native.n_built_in_functor_create(b"Negate") diff --git a/python/ThrustRTC/Merging.py b/python/ThrustRTC/Merging.py index b4dabb9..d67eda7 100644 --- a/python/ThrustRTC/Merging.py +++ b/python/ThrustRTC/Merging.py @@ -1,13 +1,13 @@ -import PyThrustRTC as native +from .Native import ffi, native def Merge(vec1, vec2, vec_out, comp = None): - cptr_comp = None + cptr_comp = ffi.NULL if comp!=None: cptr_comp = comp.m_cptr native.n_merge(vec1.m_cptr, vec2.m_cptr, vec_out.m_cptr, cptr_comp) def Merge_By_Key(keys1, keys2, value1, value2, keys_out, value_out, comp = None): - cptr_comp = None + cptr_comp = ffi.NULL if comp!=None: cptr_comp = comp.m_cptr native.n_merge_by_key(keys1.m_cptr, keys2.m_cptr, value1.m_cptr, value2.m_cptr, keys_out.m_cptr, value_out.m_cptr, cptr_comp) diff --git a/python/ThrustRTC/Native.py b/python/ThrustRTC/Native.py new file mode 100644 index 0000000..6ad137d --- /dev/null +++ b/python/ThrustRTC/Native.py @@ -0,0 +1,18 @@ +import os +import sys +import site +from .cffi import ffi + +if os.name == 'nt': + fn_thrustrtc = 'PyThrustRTC.dll' +elif os.name == "posix": + fn_thrustrtc = 'libPyThrustRTC.so' + +path_thrustrtc = sys.prefix+"/Fei/"+fn_thrustrtc +if not os.path.isfile(path_thrustrtc): + path_thrustrtc = site.USER_BASE+"/Fei/"+fn_thrustrtc + if not os.path.isfile(path_thrustrtc): + path_thrustrtc = os.path.dirname(__file__)+"/../"+fn_thrustrtc + +native = ffi.dlopen(path_thrustrtc) + diff --git a/python/ThrustRTC/PrefixSums.py b/python/ThrustRTC/PrefixSums.py index bcaf550..30e8f2b 100644 --- a/python/ThrustRTC/PrefixSums.py +++ b/python/ThrustRTC/PrefixSums.py @@ -1,16 +1,16 @@ -import PyThrustRTC as native +from .Native import ffi, native def Inclusive_Scan(vec_in, vec_out, binary_op = None): - cptr_binary_op = None + cptr_binary_op = ffi.NULL if binary_op!=None: cptr_binary_op = binary_op.m_cptr native.n_inclusive_scan(vec_in.m_cptr, vec_out.m_cptr, cptr_binary_op) def Exclusive_Scan(vec_in, vec_out, value_init = None, binary_op = None): - cptr_init = None + cptr_init = ffi.NULL if value_init!=None: cptr_init = value_init.m_cptr - cptr_binary_op = None + cptr_binary_op = ffi.NULL if binary_op!=None: cptr_binary_op = binary_op.m_cptr native.n_exclusive_scan(vec_in.m_cptr, vec_out.m_cptr, cptr_init, cptr_binary_op) @@ -22,22 +22,22 @@ def Transform_Exclusive_Scan(vec_in, vec_out, unary_op, value_init, binary_op): native.n_transform_exclusive_scan(vec_in.m_cptr, vec_out.m_cptr, unary_op.m_cptr, value_init.m_cptr, binary_op.m_cptr) def Inclusive_Scan_By_Key(vec_key, vec_value, vec_out, binary_pred = None, binary_op = None): - cptr_binary_pred = None + cptr_binary_pred = ffi.NULL if binary_pred!=None: cptr_binary_pred = binary_pred.m_cptr - cptr_binary_op = None + cptr_binary_op = ffi.NULL if binary_op!=None: cptr_binary_op = binary_op.m_cptr native.n_inclusive_scan_by_key(vec_key.m_cptr, vec_value.m_cptr, vec_out.m_cptr, cptr_binary_pred, cptr_binary_op) def Exclusive_Scan_By_Key(vec_key, vec_value, vec_out, value_init = None, binary_pred = None, binary_op = None): - cptr_init = None + cptr_init = ffi.NULL if value_init!=None: cptr_init = value_init.m_cptr - cptr_binary_pred = None + cptr_binary_pred = ffi.NULL if binary_pred!=None: cptr_binary_pred = binary_pred.m_cptr - cptr_binary_op = None + cptr_binary_op = ffi.NULL if binary_op!=None: cptr_binary_op = binary_op.m_cptr native.n_exclusive_scan_by_key(vec_key.m_cptr, vec_value.m_cptr, vec_out.m_cptr, cptr_init, cptr_binary_pred, cptr_binary_op) diff --git a/python/ThrustRTC/Reductions.py b/python/ThrustRTC/Reductions.py index fed5ea4..d5a83e8 100644 --- a/python/ThrustRTC/Reductions.py +++ b/python/ThrustRTC/Reductions.py @@ -1,4 +1,6 @@ -import PyThrustRTC as native +import numpy as np +from .Native import ffi, native +from .DeviceViewable import ValueFromDVPtr def Count(vec, value): return native.n_count(vec.m_cptr, value.m_cptr) @@ -7,89 +9,120 @@ def Count_If(vec, pred): return native.n_count_if(vec.m_cptr, pred.m_cptr) def Reduce(vec, value_init=None, binary_op=None): - cptr_init = None + cptr_init = ffi.NULL if value_init!=None: cptr_init = value_init.m_cptr - cptr_binary_op = None + cptr_binary_op = ffi.NULL if binary_op!=None: cptr_binary_op = binary_op.m_cptr - return native.n_reduce(vec.m_cptr, cptr_init, cptr_binary_op) + dvptr = native.n_reduce(vec.m_cptr, cptr_init, cptr_binary_op) + ret = ValueFromDVPtr(dvptr) + native.n_dv_destroy(dvptr) + return ret def Equal(vec1, vec2, binary_pred=None): - cptr_binary_pred = None + cptr_binary_pred = ffi.NULL if binary_pred!=None: cptr_binary_pred = binary_pred.m_cptr - return native.n_equal(vec1.m_cptr, vec2.m_cptr, cptr_binary_pred) + res = native.n_equal(vec1.m_cptr, vec2.m_cptr, cptr_binary_pred) + if res<0: + return None + return res!=0 def Min_Element(vec, comp = None): - cptr_comp = None + cptr_comp = ffi.NULL if comp!=None: cptr_comp = comp.m_cptr return native.n_min_element(vec.m_cptr, cptr_comp) def Max_Element(vec, comp = None): - cptr_comp = None + cptr_comp = ffi.NULL if comp!=None: cptr_comp = comp.m_cptr return native.n_max_element(vec.m_cptr, cptr_comp) def MinMax_Element(vec, comp = None): - cptr_comp = None + cptr_comp = ffi.NULL if comp!=None: cptr_comp = comp.m_cptr - return native.n_minmax_element(vec.m_cptr, cptr_comp) + arr_min_max = np.empty(2, dtype=np.uint64) + native.n_minmax_element(vec.m_cptr, cptr_comp, ffi.cast("unsigned long long *", arr_min_max.__array_interface__['data'][0])) + return (arr_min_max[0], arr_min_max[1]) def Inner_Product(vec1, vec2, value_init, binary_op1 = None, binary_op2 = None): - cptr_binary_op1 = None + cptr_binary_op1 = ffi.NULL if binary_op1!=None: cptr_binary_op1 = binary_op1.m_cptr - cptr_binary_op2 = None + cptr_binary_op2 = ffi.NULL if binary_op2!=None: cptr_binary_op2 = binary_op2.m_cptr - return native.n_inner_product(vec1.m_cptr, vec2.m_cptr, value_init.m_cptr, cptr_binary_op1, cptr_binary_op2,) + dvptr = native.n_inner_product(vec1.m_cptr, vec2.m_cptr, value_init.m_cptr, cptr_binary_op1, cptr_binary_op2,) + ret = ValueFromDVPtr(dvptr) + native.n_dv_destroy(dvptr) + return ret def Transform_Reduce(vec, unary_op, value_init, binary_op): - cptr_binary_op = None + cptr_binary_op = ffi.NULL if binary_op!=None: cptr_binary_op = binary_op.m_cptr - cptr_unary_op = None + cptr_unary_op = ffi.NULL if unary_op!=None: cptr_unary_op = unary_op.m_cptr - return native.n_transform_reduce(vec.m_cptr, cptr_unary_op, value_init.m_cptr, cptr_binary_op) + dvptr = native.n_transform_reduce(vec.m_cptr, cptr_unary_op, value_init.m_cptr, cptr_binary_op) + ret = ValueFromDVPtr(dvptr) + native.n_dv_destroy(dvptr) + return ret def All_Of(vec, pred): - cptr_pred = None + cptr_pred = ffi.NULL if pred!=None: cptr_pred = pred.m_cptr - return native.n_all_of(vec.m_cptr, cptr_pred) + res = native.n_all_of(vec.m_cptr, cptr_pred) + if res<0: + return None + return res!=0 def Any_Of(vec, pred): - cptr_pred = None + cptr_pred = ffi.NULL if pred!=None: cptr_pred = pred.m_cptr - return native.n_any_of(vec.m_cptr, cptr_pred) + res = native.n_any_of(vec.m_cptr, cptr_pred) + if res<0: + return None + return res!=0 def None_Of(vec, pred): - cptr_pred = None + cptr_pred = ffi.NULL if pred!=None: cptr_pred = pred.m_cptr - return native.n_none_of(vec.m_cptr, cptr_pred) + res = native.n_none_of(vec.m_cptr, cptr_pred) + if res<0: + return None + return res!=0 + def Reduce_By_Key(key_in, value_in, key_out, value_out, binary_pred = None, binary_op = None): - cptr_binary_pred = None + cptr_binary_pred = ffi.NULL if binary_pred!=None: cptr_binary_pred = binary_pred.m_cptr - cptr_binary_op = None + cptr_binary_op = ffi.NULL if binary_op!=None: cptr_binary_op = binary_op.m_cptr return native.n_reduce_by_key(key_in.m_cptr, value_in.m_cptr, key_out.m_cptr, value_out.m_cptr, cptr_binary_pred, cptr_binary_op) def Is_Partitioned(vec, pred): - return native.n_is_partitioned(vec.m_cptr, pred.m_cptr) + res = native.n_is_partitioned(vec.m_cptr, pred.m_cptr) + if res<0: + return None + return res!=0 def Is_Sorted(vec, comp = None): - cptr_comp = None + cptr_comp = ffi.NULL if comp!=None: cptr_comp = comp.m_cptr - return native.n_is_sorted(vec.m_cptr, cptr_comp) + res = native.n_is_sorted(vec.m_cptr, cptr_comp) + if res<0: + return None + return res!=0 + diff --git a/python/ThrustRTC/Reordering.py b/python/ThrustRTC/Reordering.py index a0a8fad..5e616ec 100644 --- a/python/ThrustRTC/Reordering.py +++ b/python/ThrustRTC/Reordering.py @@ -1,4 +1,4 @@ -import PyThrustRTC as native +from .Native import ffi, native def Copy_If(vec_in, vec_out, pred): return native.n_copy_if(vec_in.m_cptr, vec_out.m_cptr, pred.m_cptr) @@ -25,25 +25,25 @@ def Remove_Copy_If_Stencil(vec_in, stencil, vec_out, pred): return native.n_remove_copy_if_stencil(vec_in.m_cptr, stencil.m_cptr, vec_out.m_cptr, pred.m_cptr) def Unique(vec, binary_pred = None): - cptr_binary_pred = None + cptr_binary_pred = ffi.NULL if binary_pred!=None: cptr_binary_pred = binary_pred.m_cptr return native.n_unique(vec.m_cptr, cptr_binary_pred) def Unique_Copy(vec_in, vec_out, binary_pred = None): - cptr_binary_pred = None + cptr_binary_pred = ffi.NULL if binary_pred!=None: cptr_binary_pred = binary_pred.m_cptr return native.n_unique_copy(vec_in.m_cptr, vec_out.m_cptr, cptr_binary_pred) def Unique_By_Key(keys, values, binary_pred = None): - cptr_binary_pred = None + cptr_binary_pred = ffi.NULL if binary_pred!=None: cptr_binary_pred = binary_pred.m_cptr return native.n_unique_by_key(keys.m_cptr, values.m_cptr, cptr_binary_pred) def Unique_By_Key_Copy(keys_in, values_in, keys_out, values_out, binary_pred = None): - cptr_binary_pred = None + cptr_binary_pred = ffi.NULL if binary_pred!=None: cptr_binary_pred = binary_pred.m_cptr return native.n_unique_by_key_copy(keys_in.m_cptr, values_in.m_cptr, keys_out.m_cptr, values_out.m_cptr, cptr_binary_pred) diff --git a/python/ThrustRTC/Searching.py b/python/ThrustRTC/Searching.py index 650668c..a923940 100644 --- a/python/ThrustRTC/Searching.py +++ b/python/ThrustRTC/Searching.py @@ -1,62 +1,86 @@ -import PyThrustRTC as native +from .Native import ffi, native def Find(vec, value): - return native.n_find(vec.m_cptr, value.m_cptr) + res = native.n_find(vec.m_cptr, value.m_cptr) + if res==-1: + return None + return res def Find_If(vec, pred): - return native.n_find_if(vec.m_cptr, pred.m_cptr) + res = native.n_find_if(vec.m_cptr, pred.m_cptr) + if res==-1: + return None + return res def Find_If_Not(vec, pred): - return native.n_find_if_not(vec.m_cptr, pred.m_cptr) + res = native.n_find_if_not(vec.m_cptr, pred.m_cptr) + if res==-1: + return None + return res def Mismatch(vec1, vec2, pred=None): - cptr_pred = None + cptr_pred = ffi.NULL if pred!=None: cptr_pred = pred.m_cptr - return native.n_mismatch(vec1.m_cptr, vec2.m_cptr, cptr_pred) + res = native.n_mismatch(vec1.m_cptr, vec2.m_cptr, cptr_pred) + if res==-1: + return None + return res def Lower_Bound(vec, value, comp = None): - cptr_comp = None + cptr_comp = ffi.NULL if comp!=None: cptr_comp = comp.m_cptr return native.n_lower_bound(vec.m_cptr, value.m_cptr, cptr_comp) def Upper_Bound(vec, value, comp = None): - cptr_comp = None + cptr_comp = ffi.NULL if comp!=None: cptr_comp = comp.m_cptr return native.n_upper_bound(vec.m_cptr, value.m_cptr, cptr_comp) def Binary_Search(vec, value, comp = None): - cptr_comp = None + cptr_comp = ffi.NULL if comp!=None: cptr_comp = comp.m_cptr - return native.n_binary_search(vec.m_cptr, value.m_cptr, cptr_comp) + res = native.n_binary_search(vec.m_cptr, value.m_cptr, cptr_comp) + if res <0: + return None + return res!=0 def Partition_Point(vec, pred): return native.n_partition_point(vec.m_cptr, pred.m_cptr) def Is_Sorted_Until(vec, comp = None): - cptr_comp = None + cptr_comp = ffi.NULL if comp!=None: cptr_comp = comp.m_cptr return native.n_is_sorted_until(vec.m_cptr, cptr_comp) def Lower_Bound_V(vec, values, result, comp = None): - cptr_comp = None + cptr_comp = ffi.NULL if comp!=None: cptr_comp = comp.m_cptr - return native.n_lower_bound_v(vec.m_cptr, values.m_cptr, result.m_cptr, cptr_comp) + res = native.n_lower_bound_v(vec.m_cptr, values.m_cptr, result.m_cptr, cptr_comp) + if res == -1: + return None + return res def Upper_Bound_V(vec, values, result, comp = None): - cptr_comp = None + cptr_comp = ffi.NULL if comp!=None: cptr_comp = comp.m_cptr - return native.n_upper_bound_v(vec.m_cptr, values.m_cptr, result.m_cptr, cptr_comp) + res = native.n_upper_bound_v(vec.m_cptr, values.m_cptr, result.m_cptr, cptr_comp) + if res == -1: + return None + return res def Binary_Search_V(vec, values, result, comp = None): - cptr_comp = None + cptr_comp = ffi.NULL if comp!=None: cptr_comp = comp.m_cptr - return native.n_binary_search_v(vec.m_cptr, values.m_cptr, result.m_cptr, cptr_comp) + res = native.n_binary_search_v(vec.m_cptr, values.m_cptr, result.m_cptr, cptr_comp) + if res == -1: + return None + return res diff --git a/python/ThrustRTC/Sorting.py b/python/ThrustRTC/Sorting.py index b5f3e00..80c48f4 100644 --- a/python/ThrustRTC/Sorting.py +++ b/python/ThrustRTC/Sorting.py @@ -1,13 +1,13 @@ -import PyThrustRTC as native +from .Native import ffi, native -def Sort(vec, comp = None,): - cptr_comp = None +def Sort(vec, comp = None): + cptr_comp = ffi.NULL if comp!=None: cptr_comp = comp.m_cptr native.n_sort(vec.m_cptr, cptr_comp) def Sort_By_Key(keys, values, comp = None): - cptr_comp = None + cptr_comp = ffi.NULL if comp!=None: cptr_comp = comp.m_cptr native.n_sort_by_key(keys.m_cptr, values.m_cptr, cptr_comp) diff --git a/python/ThrustRTC/Transformations.py b/python/ThrustRTC/Transformations.py index b1752a5..5036872 100644 --- a/python/ThrustRTC/Transformations.py +++ b/python/ThrustRTC/Transformations.py @@ -1,11 +1,8 @@ -import PyThrustRTC as native +from .Native import ffi, native def Fill(vec, value): native.n_fill(vec.m_cptr, value.m_cptr) -def For_Each(vec, f): - native.n_for_each(vec.m_cptr, f.m_cptr) - def Replace(vec, old_value, new_value): native.n_replace(vec.m_cptr, old_value.m_cptr, new_value.m_cptr) @@ -18,17 +15,20 @@ def Replace_Copy(vec_in, vec_out, old_value, new_value): def Replace_Copy_If(vec_in, vec_out, pred, new_value): native.n_replace_copy_if(vec_in.m_cptr, vec_out.m_cptr, pred.m_cptr, new_value.m_cptr) +def For_Each(vec, f): + native.n_for_each(vec.m_cptr, f.m_cptr) + def Adjacent_Difference(vec_in, vec_out, binary_op=None): - cptr_binary_op = None + cptr_binary_op = ffi.NULL if binary_op!=None: cptr_binary_op = binary_op.m_cptr - native.n_adjacent_difference(vec_in.m_cptr, vec_out.m_cptr, cptr_binary_op,) + native.n_adjacent_difference(vec_in.m_cptr, vec_out.m_cptr, cptr_binary_op) def Sequence(vec, value_init=None, value_step=None): - cptr_init = None + cptr_init = ffi.NULL if value_init!=None: cptr_init = value_init.m_cptr - cptr_step = None + cptr_step = ffi.NULL if value_step!=None: cptr_step = value_step.m_cptr native.n_sequence(vec.m_cptr, cptr_init, cptr_step) @@ -50,3 +50,4 @@ def Transform_If_Stencil(vec_in, vec_stencil, vec_out, op, pred): def Transform_Binary_If_Stencil(vec_in1, vec_in2, vec_stencil, vec_out, op, pred): native.n_transform_binary_if_stencil(vec_in1.m_cptr, vec_in2.m_cptr, vec_stencil.m_cptr, vec_out.m_cptr, op.m_cptr, pred.m_cptr) + diff --git a/python/ThrustRTC/__init__.py b/python/ThrustRTC/__init__.py index c927ea8..372782e 100644 --- a/python/ThrustRTC/__init__.py +++ b/python/ThrustRTC/__init__.py @@ -1,9 +1,3 @@ -import PyThrustRTC as native -import os - -def set_libnvrtc_path(path): - native.n_set_libnvrtc_path(path) - from .Context import * from .DeviceViewable import * from .DVVector import device_vector, device_vector_from_numpy, device_vector_from_list, device_vector_from_dvs, DVNumbaVector diff --git a/python/ThrustRTC/cffi.py b/python/ThrustRTC/cffi.py new file mode 100644 index 0000000..919f89f --- /dev/null +++ b/python/ThrustRTC/cffi.py @@ -0,0 +1,8 @@ +# auto-generated file +import _cffi_backend + +ffi = _cffi_backend.FFI('ThrustRTC.cffi', + _version = 0x2601, + _types = b'\x00\x00\x8F\x0D\x00\x01\x13\x03\x00\x00\x00\x0F\x00\x00\x9B\x0D\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x9E\x0D\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x0E\x0D\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x0E\x0D\x00\x00\x01\x11\x00\x00\x07\x01\x00\x00\x07\x01\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x0E\x0D\x00\x00\x01\x11\x00\x00\x07\x01\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x0E\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x0E\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x07\x01\x00\x00\x08\x01\x00\x00\x00\x0F\x00\x00\x0E\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x08\x01\x00\x00\x00\x0F\x00\x00\x0E\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x93\x03\x00\x00\x00\x0F\x00\x00\x0E\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x0E\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x0E\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x07\x01\x00\x00\x00\x0F\x00\x00\x0E\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x0E\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x0E\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\xA4\x0D\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\xA4\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\xA4\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x1F\x0D\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x1F\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x1F\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x1F\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x1F\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x1F\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x93\x0D\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x93\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x93\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x01\x11\x03\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x8F\x11\x00\x00\x0C\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x8F\x11\x00\x00\x0C\x01\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x0E\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x0D\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x07\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x0B\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x08\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x08\x01\x00\x00\x08\x01\x00\x00\x08\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x0C\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x0C\x01\x00\x00\x8F\x03\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x0C\x01\x00\x01\x12\x03\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x01\x11\x00\x00\x8F\x11\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x01\x11\x00\x00\x8F\x11\x00\x00\x8F\x11\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x01\x11\x00\x00\x8F\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x01\x11\x00\x00\x0C\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x01\x11\x00\x00\x0C\x01\x00\x00\x0C\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x8F\x11\x00\x00\x8F\x11\x00\x00\x8F\x11\x00\x00\x0C\x01\x00\x00\x08\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x8F\x11\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x01\x13\x0D\x00\x00\x8F\x11\x00\x00\x00\x0F\x00\x01\x13\x0D\x00\x00\x8F\x11\x00\x00\x8F\x11\x00\x00\x00\x0F\x00\x01\x13\x0D\x00\x00\x8F\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x01\x13\x0D\x00\x00\x08\x01\x00\x00\x00\x0F\x00\x01\x13\x0D\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x01\x13\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x0C\x01\x00\x00\x0C\x01\x00\x00\x00\x0F\x00\x01\x13\x0D\x00\x00\x00\x0F\x00\x00\x02\x01\x00\x01\x13\x03\x00\x00\x00\x01', + _globals = (b'\x00\x00\xFB\x23n_add_built_in_header',0,b'\x00\x00\xF8\x23n_add_code_block',0,b'\x00\x00\xFF\x23n_add_constant_object',0,b'\x00\x00\xF8\x23n_add_include_dir',0,b'\x00\x00\xF8\x23n_add_inlcude_filename',0,b'\x00\x00\x2B\x23n_adjacent_difference',0,b'\x00\x00\x17\x23n_all_of',0,b'\x00\x00\x17\x23n_any_of',0,b'\x00\x00\x2B\x23n_binary_search',0,b'\x00\x00\x30\x23n_binary_search_v',0,b'\x00\x00\x8E\x23n_built_in_functor_create',0,b'\x00\x00\x17\x23n_copy',0,b'\x00\x00\x68\x23n_copy_if',0,b'\x00\x00\x6D\x23n_copy_if_stencil',0,b'\x00\x00\x85\x23n_count',0,b'\x00\x00\x85\x23n_count_if',0,b'\x00\x00\xA9\x23n_dim3_create',0,b'\x00\x01\x06\x23n_dim3_destroy',0,b'\x00\x01\x06\x23n_dv_destroy',0,b'\x00\x00\x00\x23n_dv_name_view_cls',0,b'\x00\x00\xA0\x23n_dvbool_create',0,b'\x00\x00\x09\x23n_dvbool_value',0,b'\x00\x00\xCA\x23n_dvconstant_create',0,b'\x00\x00\xCA\x23n_dvcounter_create',0,b'\x00\x00\xD7\x23n_dvcustomvector_create',0,b'\x00\x00\x91\x23n_dvdiscard_create',0,b'\x00\x00\x9A\x23n_dvdouble_create',0,b'\x00\x00\x03\x23n_dvdouble_value',0,b'\x00\x00\x9D\x23n_dvfloat_create',0,b'\x00\x00\x06\x23n_dvfloat_value',0,b'\x00\x00\xA0\x23n_dvint16_create',0,b'\x00\x00\x09\x23n_dvint16_value',0,b'\x00\x00\xA0\x23n_dvint32_create',0,b'\x00\x00\x09\x23n_dvint32_value',0,b'\x00\x00\xA3\x23n_dvint64_create',0,b'\x00\x00\x55\x23n_dvint64_value',0,b'\x00\x00\xA0\x23n_dvint8_create',0,b'\x00\x00\x09\x23n_dvint8_value',0,b'\x00\x00\xD3\x23n_dvpermutation_create',0,b'\x00\x00\xCE\x23n_dvrange_create',0,b'\x00\x00\xB9\x23n_dvreverse_create',0,b'\x00\x00\xC5\x23n_dvtransform_create',0,b'\x00\x00\xD3\x23n_dvtuple_create',0,b'\x00\x00\xA6\x23n_dvuint16_create',0,b'\x00\x00\x61\x23n_dvuint16_value',0,b'\x00\x00\xA6\x23n_dvuint32_create',0,b'\x00\x00\x61\x23n_dvuint32_value',0,b'\x00\x00\xAE\x23n_dvuint64_create',0,b'\x00\x00\x82\x23n_dvuint64_value',0,b'\x00\x00\xA6\x23n_dvuint8_create',0,b'\x00\x00\x61\x23n_dvuint8_value',0,b'\x00\x00\x95\x23n_dvvector_create',0,b'\x00\x00\xB9\x23n_dvvector_from_dvs',0,b'\x00\x01\x09\x23n_dvvector_to_host',0,b'\x00\x00\x95\x23n_dvvectoradaptor_create',0,b'\x00\x00\x00\x23n_dvvectorlike_name_elem_cls',0,b'\x00\x00\x82\x23n_dvvectorlike_size',0,b'\x00\x00\xD3\x23n_dvzipped_create',0,b'\x00\x00\x2B\x23n_equal',0,b'\x00\x00\x30\x23n_exclusive_scan',0,b'\x00\x00\x44\x23n_exclusive_scan_by_key',0,b'\x00\x00\x17\x23n_fill',0,b'\x00\x00\x58\x23n_find',0,b'\x00\x00\x58\x23n_find_if',0,b'\x00\x00\x58\x23n_find_if_not',0,b'\x00\x00\xC0\x23n_for_create',0,b'\x00\x01\x06\x23n_for_destroy',0,b'\x00\x00\x17\x23n_for_each',0,b'\x00\x00\x0C\x23n_for_launch',0,b'\x00\x00\x12\x23n_for_launch_n',0,b'\x00\x00\x09\x23n_for_num_params',0,b'\x00\x00\xE5\x23n_functor_create',0,b'\x00\x00\x2B\x23n_gather',0,b'\x00\x00\x3D\x23n_gather_if',0,b'\x00\x00\x2B\x23n_inclusive_scan',0,b'\x00\x00\x3D\x23n_inclusive_scan_by_key',0,b'\x00\x00\xF1\x23n_inner_product',0,b'\x00\x00\x17\x23n_is_partitioned',0,b'\x00\x00\x17\x23n_is_sorted',0,b'\x00\x00\x85\x23n_is_sorted_until',0,b'\x00\x00\x1B\x23n_kernel_calc_number_blocks',0,b'\x00\x00\x21\x23n_kernel_calc_optimal_block_size',0,b'\x00\x00\xBC\x23n_kernel_create',0,b'\x00\x01\x06\x23n_kernel_destroy',0,b'\x00\x00\x36\x23n_kernel_launch',0,b'\x00\x00\x09\x23n_kernel_num_params',0,b'\x00\x00\x89\x23n_lower_bound',0,b'\x00\x00\x30\x23n_lower_bound_v',0,b'\x00\x00\x85\x23n_max_element',0,b'\x00\x00\x30\x23n_merge',0,b'\x00\x00\x4C\x23n_merge_by_key',0,b'\x00\x00\x85\x23n_min_element',0,b'\x00\x00\x26\x23n_minmax_element',0,b'\x00\x00\x5C\x23n_mismatch',0,b'\x00\x00\x17\x23n_none_of',0,b'\x00\x00\x64\x23n_partition',0,b'\x00\x00\x6D\x23n_partition_copy',0,b'\x00\x00\x73\x23n_partition_copy_stencil',0,b'\x00\x00\x85\x23n_partition_point',0,b'\x00\x00\x68\x23n_partition_stencil',0,b'\x00\x00\xB5\x23n_pointer_array_create',0,b'\x00\x01\x06\x23n_pointer_array_destroy',0,b'\x00\x00\xE0\x23n_reduce',0,b'\x00\x00\x7A\x23n_reduce_by_key',0,b'\x00\x00\x64\x23n_remove',0,b'\x00\x00\x68\x23n_remove_copy',0,b'\x00\x00\x68\x23n_remove_copy_if',0,b'\x00\x00\x6D\x23n_remove_copy_if_stencil',0,b'\x00\x00\x64\x23n_remove_if',0,b'\x00\x00\x68\x23n_remove_if_stencil',0,b'\x00\x00\x2B\x23n_replace',0,b'\x00\x00\x30\x23n_replace_copy',0,b'\x00\x00\x30\x23n_replace_copy_if',0,b'\x00\x00\x2B\x23n_replace_if',0,b'\x00\x00\x2B\x23n_scatter',0,b'\x00\x00\x3D\x23n_scatter_if',0,b'\x00\x00\x2B\x23n_sequence',0,b'\x00\x00\xF8\x23n_set_libnvrtc_path',0,b'\x00\x01\x03\x23n_set_verbose',0,b'\x00\x00\x17\x23n_sort',0,b'\x00\x00\x2B\x23n_sort_by_key',0,b'\x00\x00\xB1\x23n_string_array_create',0,b'\x00\x01\x06\x23n_string_array_destroy',0,b'\x00\x00\x17\x23n_swap',0,b'\x00\x00\x17\x23n_tabulate',0,b'\x00\x00\x2B\x23n_transform',0,b'\x00\x00\x30\x23n_transform_binary',0,b'\x00\x00\x44\x23n_transform_binary_if_stencil',0,b'\x00\x00\x3D\x23n_transform_exclusive_scan',0,b'\x00\x00\x30\x23n_transform_if',0,b'\x00\x00\x3D\x23n_transform_if_stencil',0,b'\x00\x00\x30\x23n_transform_inclusive_scan',0,b'\x00\x00\xEB\x23n_transform_reduce',0,b'\x00\x00\x64\x23n_unique',0,b'\x00\x00\x68\x23n_unique_by_key',0,b'\x00\x00\x73\x23n_unique_by_key_copy',0,b'\x00\x00\x68\x23n_unique_copy',0,b'\x00\x00\x89\x23n_upper_bound',0,b'\x00\x00\x30\x23n_upper_bound_v',0,b'\x00\x01\x0F\x23n_wait',0), +) diff --git a/python/ThrustRTC/cffi_build.py b/python/ThrustRTC/cffi_build.py new file mode 100644 index 0000000..9d5fe9a --- /dev/null +++ b/python/ThrustRTC/cffi_build.py @@ -0,0 +1,181 @@ +import cffi + +ffibuilder = cffi.FFI() +ffibuilder.set_source("ThrustRTC.cffi", None) + +ffibuilder.cdef(""" +// utils +void* n_string_array_create(unsigned long long size, const char* const* strs); +void n_string_array_destroy(void* ptr_arr); +void* n_pointer_array_create(unsigned long long size, const void* const* ptrs); +void n_pointer_array_destroy(void* ptr_arr); +void* n_dim3_create(unsigned x, unsigned y, unsigned z); +void n_dim3_destroy(void* cptr); + +// Context +void n_set_libnvrtc_path(const char* path); +void n_set_verbose(unsigned verbose); +void n_add_include_dir(const char* dir); +void n_add_built_in_header(const char* filename, const char* filecontent); +void n_add_inlcude_filename(const char* fn); +void n_add_code_block(const char* line); +void n_add_constant_object(const char* name, void* cptr); +void n_wait(); + +void* n_kernel_create(void* ptr_param_list, const char* body); +void n_kernel_destroy(void* cptr); +int n_kernel_num_params(void* cptr); +int n_kernel_calc_optimal_block_size(void* ptr_kernel, void* ptr_arg_list, unsigned sharedMemBytes); +int n_kernel_calc_number_blocks(void* ptr_kernel, void* ptr_arg_list, int sizeBlock, unsigned sharedMemBytes); +int n_kernel_launch(void* ptr_kernel, void* ptr_gridDim, void* ptr_blockDim, void* ptr_arg_list, int sharedMemBytes); + +void* n_for_create(void* ptr_param_list, const char* name_iter, const char* body); +void n_for_destroy(void* cptr); +int n_for_num_params(void* cptr); +int n_for_launch(void* cptr, int begin, int end, void* ptr_arg_list); +int n_for_launch_n(void* cptr, int n, void* ptr_arg_list); + +// DeviceViewable +const char* n_dv_name_view_cls(void* cptr); +void n_dv_destroy(void* cptr); +void* n_dvint8_create(int v); +int n_dvint8_value(void* cptr); +void* n_dvuint8_create(unsigned v); +unsigned n_dvuint8_value(void* cptr); +void* n_dvint16_create(int v); +int n_dvint16_value(void* cptr); +void* n_dvuint16_create(unsigned v); +unsigned n_dvuint16_value(void* cptr); +void* n_dvint32_create(int v); +int n_dvint32_value(void* cptr); +void* n_dvuint32_create(unsigned v); +unsigned n_dvuint32_value(void* cptr); +void* n_dvint64_create(long long v); +long long n_dvint64_value(void* cptr); +void* n_dvuint64_create(unsigned long long v); +unsigned long long n_dvuint64_value(void* cptr); +void* n_dvfloat_create(float v); +float n_dvfloat_value(void* cptr); +void* n_dvdouble_create(double v); +double n_dvdouble_value(void* cptr); +void* n_dvbool_create(int v); +int n_dvbool_value(void* cptr); + +// DVVector +const char* n_dvvectorlike_name_elem_cls(void* cptr); +unsigned long long n_dvvectorlike_size(void* cptr); +void* n_dvvector_create(const char* elem_cls, unsigned long long size, void* hdata); +void n_dvvector_to_host(void* cptr, void* hdata, unsigned long long begin, unsigned long long end); +void* n_dvvector_from_dvs(void* ptr_dvs); +void* n_dvvectoradaptor_create(const char* elem_cls, unsigned long long size, void* data); +void* n_dvrange_create(void* ptr_in, unsigned long long begin, unsigned long long end); + +// Tuple +void* n_dvtuple_create(void* ptr_dvs, void* ptr_names); + +// Fake-Vectors +void* n_dvconstant_create(void* cptr, unsigned long long size); +void* n_dvcounter_create(void* cptr, unsigned long long size); +void* n_dvdiscard_create(const char* elem_cls, unsigned long long size); +void* n_dvpermutation_create(void* ptr_value, void* ptr_index); +void* n_dvreverse_create(void* ptr_value); +void* n_dvtransform_create(void* ptr_in, const char* elem_cls, void* ptr_op); +void* n_dvzipped_create(void* ptr_vecs, void* ptr_elem_names); +void* n_dvcustomvector_create(void* ptr_dvs, void* ptr_names, const char* name_idx, const char* body, const char* elem_cls, unsigned long long size, unsigned read_only); + +// Functor +void* n_functor_create(void* ptr_dvs, void* ptr_names, void* ptr_functor_params, const char* code_body); +void* n_built_in_functor_create(const char* name_built_in_view_cls); + +// Transformations +int n_fill(void* ptr_vec, void* ptr_value); +int n_replace(void* ptr_vec, void* ptr_old_value, void* ptr_new_value); +int n_replace_if(void* ptr_vec, void* p_pred, void* ptr_new_value); +int n_replace_copy(void* ptr_in, void* ptr_out, void* ptr_old_value, void* ptr_new_value); +int n_replace_copy_if(void* ptr_in, void* ptr_out, void* p_pred, void* ptr_new_value); +int n_for_each(void* ptr_vec, void* ptr_f); +int n_adjacent_difference(void* ptr_in, void* ptr_out, void* ptr_binary_op); +int n_sequence(void* ptr_vec, void* ptr_value_init, void* ptr_value_step); +int n_tabulate(void* ptr_vec, void* ptr_op); +int n_transform(void* ptr_in, void* ptr_out, void* ptr_op); +int n_transform_binary(void* ptr_in1, void* ptr_in2, void* ptr_out, void* ptr_op); +int n_transform_if(void* ptr_in, void* ptr_out, void* ptr_op, void* ptr_pred); +int n_transform_if_stencil(void* ptr_in, void* ptr_stencil, void* ptr_out, void* ptr_op, void* ptr_pred); +int n_transform_binary_if_stencil(void* ptr_in1, void* ptr_in2, void* ptr_stencil, void* ptr_out, void* ptr_op, void* ptr_pred); + +// Copying +int n_gather(void* ptr_map, void* ptr_in, void* ptr_out); +int n_gather_if(void* ptr_map, void* ptr_stencil, void* ptr_in, void* ptr_out, void* ptr_pred); +int n_scatter(void* ptr_in, void* ptr_map, void* ptr_out); +int n_scatter_if(void* ptr_in, void* ptr_map, void* ptr_stencil, void* ptr_out, void* ptr_pred); +int n_copy(void* ptr_in, void* ptr_out); +int n_swap(void* ptr_vec1, void* ptr_vec2); + +// Redutions +unsigned long long n_count(void* ptr_vec, void* ptr_value); +unsigned long long n_count_if(void* ptr_vec, void* ptr_pred); +void* n_reduce(void* ptr_vec, void* ptr_init, void* ptr_bin_op); +unsigned n_reduce_by_key(void* ptr_key_in, void* ptr_value_in, void* ptr_key_out, void* ptr_value_out, void* ptr_binary_pred, void* ptr_binary_op); +int n_equal(void* ptr_vec1, void* ptr_vec2, void* ptr_binary_pred); +unsigned long long n_min_element(void* ptr_vec, void* ptr_comp); +unsigned long long n_max_element(void* ptr_vec, void* ptr_comp); +int n_minmax_element(void* ptr_vec, void* ptr_comp, unsigned long long* ret); +void* n_inner_product(void* ptr_vec1, void* ptr_vec2, void* ptr_init, void* ptr_binary_op1, void* ptr_binary_op2); +void* n_transform_reduce(void* ptr_vec, void* ptr_unary_op, void* ptr_init, void* ptr_binary_op); +int n_all_of(void* ptr_vec, void* ptr_pred); +int n_any_of(void* ptr_vec, void* ptr_pred); +int n_none_of(void* ptr_vec, void* ptr_pred); +int n_is_partitioned(void* ptr_vec, void* ptr_pred); +int n_is_sorted(void* ptr_vec, void* ptr_comp); + +// PrefixSums +int n_inclusive_scan(void* ptr_vec_in, void* ptr_vec_out, void* ptr_binary_op); +int n_exclusive_scan(void* ptr_vec_in, void* ptr_vec_out, void* ptr_init, void* ptr_binary_op); +int n_inclusive_scan_by_key(void* ptr_vec_key, void* ptr_vec_value, void* ptr_vec_out, void* ptr_binary_pred, void* ptr_binary_op); +int n_exclusive_scan_by_key(void* ptr_vec_key, void* ptr_vec_value, void* ptr_vec_out, void* ptr_init, void* ptr_binary_pred, void* ptr_binary_op); +int n_transform_inclusive_scan(void* ptr_vec_in, void* ptr_vec_out, void* ptr_unary_op, void* ptr_binary_op); +int n_transform_exclusive_scan(void* ptr_vec_in, void* ptr_vec_out, void* ptr_unary_op, void* ptr_init, void* ptr_binary_op); + +// Reordering +unsigned n_copy_if(void* ptr_in, void* ptr_out, void* ptr_pred); +unsigned n_copy_if_stencil(void* ptr_in, void* ptr_stencil, void* ptr_out, void* ptr_pred); +unsigned n_remove(void* ptr_vec, void* ptr_value); +unsigned n_remove_copy(void* ptr_in, void* ptr_out, void* ptr_value); +unsigned n_remove_if(void* ptr_vec, void* ptr_pred); +unsigned n_remove_copy_if(void* ptr_in, void* ptr_out, void* ptr_pred); +unsigned n_remove_if_stencil(void* ptr_vec, void* ptr_stencil, void* ptr_pred); +unsigned n_remove_copy_if_stencil(void* ptr_in, void* ptr_stencil, void* ptr_out, void* ptr_pred); +unsigned n_unique(void* ptr_vec, void* ptr_binary_pred); +unsigned n_unique_copy(void* ptr_vec_in, void* ptr_vec_out, void* ptr_binary_pred); +unsigned n_unique_by_key(void* ptr_keys, void* ptr_values, void* ptr_binary_pred); +unsigned n_unique_by_key_copy(void* ptr_keys_in, void* ptr_values_in, void* ptr_key_out, void* ptr_values_out, void* ptr_binary_pred); +unsigned n_partition(void* ptr_vec, void* ptr_pred); +unsigned n_partition_stencil(void* ptr_vec, void* ptr_stencil, void* ptr_pred); +unsigned n_partition_copy(void* ptr_vec_in, void* ptr_vec_true, void* ptr_vec_false, void* ptr_pred); +unsigned n_partition_copy_stencil(void* ptr_vec_in, void* ptr_stencil, void* ptr_vec_true, void* ptr_vec_false, void* ptr_pred); + +// Searching +long long n_find(void* ptr_vec, void* ptr_value); +long long n_find_if(void* ptr_vec, void* ptr_pred); +long long n_find_if_not(void* ptr_vec, void* ptr_pred); +long long n_mismatch(void* ptr_vec1, void* ptr_vec2, void* ptr_pred); +unsigned long long n_lower_bound(void* ptr_vec, void* ptr_value, void* ptr_comp); +unsigned long long n_upper_bound(void* ptr_vec, void* ptr_value, void* ptr_comp); +int n_binary_search(void* ptr_vec, void* ptr_value, void* ptr_comp); +int n_lower_bound_v(void* ptr_vec, void* ptr_values, void* ptr_result, void* ptr_comp); +int n_upper_bound_v(void* ptr_vec, void* ptr_values, void* ptr_result, void* ptr_comp); +int n_binary_search_v(void* ptr_vec, void* ptr_values, void* ptr_result, void* ptr_comp); +unsigned long long n_partition_point(void* ptr_vec, void* ptr_pred); +unsigned long long n_is_sorted_until(void* ptr_vec, void* ptr_comp); + +// Merging +int n_merge(void* ptr_vec1, void* ptr_vec2, void* ptr_vec_out, void* ptr_comp); +int n_merge_by_key(void* ptr_keys1, void* ptr_keys2, void* ptr_value1, void* ptr_value2, void* ptr_keys_out, void* ptr_value_out, void* ptr_comp); + +// Sorting +int n_sort(void* ptr_vec, void* ptr_comp); +int n_sort_by_key(void* ptr_keys, void* ptr_values, void* ptr_comp); +""") + + +ffibuilder.compile() diff --git a/python/ThrustRTC/utils.py b/python/ThrustRTC/utils.py new file mode 100644 index 0000000..0ecc39a --- /dev/null +++ b/python/ThrustRTC/utils.py @@ -0,0 +1,34 @@ +from .Native import ffi,native + +class StrArray: + def __init__(self, arr): + c_strs = [ffi.from_buffer('char[]', s.encode('utf-8')) for s in arr] + self.m_cptr = native.n_string_array_create(len(c_strs), c_strs) + + def __del__(self): + native.n_string_array_destroy(self.m_cptr) + +class ObjArray: + def __init__(self, arr): + c_ptrs = [obj.m_cptr for obj in arr] + self.m_cptr = native.n_pointer_array_create(len(c_ptrs), c_ptrs) + + def __del__(self): + native.n_pointer_array_destroy(self.m_cptr) + +class Dim3: + def __init__(self, t): + tp = [1,1,1] + if type(t) is tuple: + tp[0:len(t)] = t[:] + else: + tp[0]=t + + self.m_cptr = native.n_dim3_create(tp[0],tp[1],tp[2]) + + def __del__(self): + native.n_dim3_destroy(self.m_cptr) + + + + diff --git a/python/adjacent_difference.hpp b/python/adjacent_difference.hpp deleted file mode 100644 index 3d672ab..0000000 --- a/python/adjacent_difference.hpp +++ /dev/null @@ -1,30 +0,0 @@ -#include -#include "TRTCContext.h" -#include "adjacent_difference.h" - -static PyObject* n_adjacent_difference(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - PyObject* py_binary_op = PyTuple_GetItem(args, 2); - Functor* binary_op = nullptr; - if (py_binary_op != Py_None) - { - binary_op = (Functor*)PyLong_AsVoidPtr(py_binary_op); - } - - if (binary_op == nullptr) - { - if (TRTC_Adjacent_Difference(*vec_in, *vec_out)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; - } - else - { - if (TRTC_Adjacent_Difference(*vec_in, *vec_out, *binary_op)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; - } -} diff --git a/python/api.h b/python/api.h new file mode 100644 index 0000000..f1608af --- /dev/null +++ b/python/api.h @@ -0,0 +1,184 @@ +#pragma once + +#if defined(WIN32) || defined(_WIN32) || defined(__WIN32__) || defined(__NT__) +#define PY_THRUSTRTC_API __declspec(dllexport) +#else +#define PY_THRUSTRTC_API +#endif + +extern "C" +{ + // utils + PY_THRUSTRTC_API void* n_string_array_create(unsigned long long size, const char* const* strs); + PY_THRUSTRTC_API void n_string_array_destroy(void* ptr_arr); + PY_THRUSTRTC_API void* n_pointer_array_create(unsigned long long size, const void* const* ptrs); + PY_THRUSTRTC_API void n_pointer_array_destroy(void* ptr_arr); + PY_THRUSTRTC_API void* n_dim3_create(unsigned x, unsigned y, unsigned z); + PY_THRUSTRTC_API void n_dim3_destroy(void* cptr); + + // Context + PY_THRUSTRTC_API void n_set_libnvrtc_path(const char* path); + PY_THRUSTRTC_API void n_set_verbose(unsigned verbose); + PY_THRUSTRTC_API void n_add_include_dir(const char* dir); + PY_THRUSTRTC_API void n_add_built_in_header(const char* filename, const char* filecontent); + PY_THRUSTRTC_API void n_add_inlcude_filename(const char* fn); + PY_THRUSTRTC_API void n_add_code_block(const char* line); + PY_THRUSTRTC_API void n_add_constant_object(const char* name, void* cptr); + PY_THRUSTRTC_API void n_wait(); + + PY_THRUSTRTC_API void* n_kernel_create(void* ptr_param_list, const char* body); + PY_THRUSTRTC_API void n_kernel_destroy(void* cptr); + PY_THRUSTRTC_API int n_kernel_num_params(void* cptr); + PY_THRUSTRTC_API int n_kernel_calc_optimal_block_size(void* ptr_kernel, void* ptr_arg_list, unsigned sharedMemBytes); + PY_THRUSTRTC_API int n_kernel_calc_number_blocks(void* ptr_kernel, void* ptr_arg_list, int sizeBlock, unsigned sharedMemBytes); + PY_THRUSTRTC_API int n_kernel_launch(void* ptr_kernel, void* ptr_gridDim, void* ptr_blockDim, void* ptr_arg_list, int sharedMemBytes); + + PY_THRUSTRTC_API void* n_for_create(void* ptr_param_list, const char* name_iter, const char* body); + PY_THRUSTRTC_API void n_for_destroy(void* cptr); + PY_THRUSTRTC_API int n_for_num_params(void* cptr); + PY_THRUSTRTC_API int n_for_launch(void* cptr, int begin, int end, void* ptr_arg_list); + PY_THRUSTRTC_API int n_for_launch_n(void* cptr, int n, void* ptr_arg_list); + + // DeviceViewable + PY_THRUSTRTC_API const char* n_dv_name_view_cls(void* cptr); + PY_THRUSTRTC_API void n_dv_destroy(void* cptr); + PY_THRUSTRTC_API void* n_dvint8_create(int v); + PY_THRUSTRTC_API int n_dvint8_value(void* cptr); + PY_THRUSTRTC_API void* n_dvuint8_create(unsigned v); + PY_THRUSTRTC_API unsigned n_dvuint8_value(void* cptr); + PY_THRUSTRTC_API void* n_dvint16_create(int v); + PY_THRUSTRTC_API int n_dvint16_value(void* cptr); + PY_THRUSTRTC_API void* n_dvuint16_create(unsigned v); + PY_THRUSTRTC_API unsigned n_dvuint16_value(void* cptr); + PY_THRUSTRTC_API void* n_dvint32_create(int v); + PY_THRUSTRTC_API int n_dvint32_value(void* cptr); + PY_THRUSTRTC_API void* n_dvuint32_create(unsigned v); + PY_THRUSTRTC_API unsigned n_dvuint32_value(void* cptr); + PY_THRUSTRTC_API void* n_dvint64_create(long long v); + PY_THRUSTRTC_API long long n_dvint64_value(void* cptr); + PY_THRUSTRTC_API void* n_dvuint64_create(unsigned long long v); + PY_THRUSTRTC_API unsigned long long n_dvuint64_value(void* cptr); + PY_THRUSTRTC_API void* n_dvfloat_create(float v); + PY_THRUSTRTC_API float n_dvfloat_value(void* cptr); + PY_THRUSTRTC_API void* n_dvdouble_create(double v); + PY_THRUSTRTC_API double n_dvdouble_value(void* cptr); + PY_THRUSTRTC_API void* n_dvbool_create(int v); + PY_THRUSTRTC_API int n_dvbool_value(void* cptr); + + // DVVector + PY_THRUSTRTC_API const char* n_dvvectorlike_name_elem_cls(void* cptr); + PY_THRUSTRTC_API unsigned long long n_dvvectorlike_size(void* cptr); + PY_THRUSTRTC_API void* n_dvvector_create(const char* elem_cls, unsigned long long size, void* hdata); + PY_THRUSTRTC_API void n_dvvector_to_host(void* cptr, void* hdata, unsigned long long begin, unsigned long long end); + PY_THRUSTRTC_API void* n_dvvector_from_dvs(void* ptr_dvs); + PY_THRUSTRTC_API void* n_dvvectoradaptor_create(const char* elem_cls, unsigned long long size, void* data); + PY_THRUSTRTC_API void* n_dvrange_create(void* ptr_in, unsigned long long begin, unsigned long long end); + + // Tuple + PY_THRUSTRTC_API void* n_dvtuple_create(void* ptr_dvs, void* ptr_names); + + // Fake-Vectors + PY_THRUSTRTC_API void* n_dvconstant_create(void* cptr, unsigned long long size); + PY_THRUSTRTC_API void* n_dvcounter_create(void* cptr, unsigned long long size); + PY_THRUSTRTC_API void* n_dvdiscard_create(const char* elem_cls, unsigned long long size); + PY_THRUSTRTC_API void* n_dvpermutation_create(void* ptr_value, void* ptr_index); + PY_THRUSTRTC_API void* n_dvreverse_create(void* ptr_value); + PY_THRUSTRTC_API void* n_dvtransform_create(void* ptr_in, const char* elem_cls, void* ptr_op); + PY_THRUSTRTC_API void* n_dvzipped_create(void* ptr_vecs, void* ptr_elem_names); + PY_THRUSTRTC_API void* n_dvcustomvector_create(void* ptr_dvs, void* ptr_names, const char* name_idx, const char* body, const char* elem_cls, unsigned long long size, unsigned read_only); + + // Functor + PY_THRUSTRTC_API void* n_functor_create(void* ptr_dvs, void* ptr_names, void* ptr_functor_params, const char* code_body); + PY_THRUSTRTC_API void* n_built_in_functor_create(const char* name_built_in_view_cls); + + // Transformations + PY_THRUSTRTC_API int n_fill(void* ptr_vec, void* ptr_value); + PY_THRUSTRTC_API int n_replace(void* ptr_vec, void* ptr_old_value, void* ptr_new_value); + PY_THRUSTRTC_API int n_replace_if(void* ptr_vec, void* p_pred, void* ptr_new_value); + PY_THRUSTRTC_API int n_replace_copy(void* ptr_in, void* ptr_out, void* ptr_old_value, void* ptr_new_value); + PY_THRUSTRTC_API int n_replace_copy_if(void* ptr_in, void* ptr_out, void* p_pred, void* ptr_new_value); + PY_THRUSTRTC_API int n_for_each(void* ptr_vec, void* ptr_f); + PY_THRUSTRTC_API int n_adjacent_difference(void* ptr_in, void* ptr_out, void* ptr_binary_op); + PY_THRUSTRTC_API int n_sequence(void* ptr_vec, void* ptr_value_init, void* ptr_value_step); + PY_THRUSTRTC_API int n_tabulate(void* ptr_vec, void* ptr_op); + PY_THRUSTRTC_API int n_transform(void* ptr_in, void* ptr_out, void* ptr_op); + PY_THRUSTRTC_API int n_transform_binary(void* ptr_in1, void* ptr_in2, void* ptr_out, void* ptr_op); + PY_THRUSTRTC_API int n_transform_if(void* ptr_in, void* ptr_out, void* ptr_op, void* ptr_pred); + PY_THRUSTRTC_API int n_transform_if_stencil(void* ptr_in, void* ptr_stencil, void* ptr_out, void* ptr_op, void* ptr_pred); + PY_THRUSTRTC_API int n_transform_binary_if_stencil(void* ptr_in1, void* ptr_in2, void* ptr_stencil, void* ptr_out, void* ptr_op, void* ptr_pred); + + // Copying + PY_THRUSTRTC_API int n_gather(void* ptr_map, void* ptr_in, void* ptr_out); + PY_THRUSTRTC_API int n_gather_if(void* ptr_map, void* ptr_stencil, void* ptr_in, void* ptr_out, void* ptr_pred); + PY_THRUSTRTC_API int n_scatter(void* ptr_in, void* ptr_map, void* ptr_out); + PY_THRUSTRTC_API int n_scatter_if(void* ptr_in, void* ptr_map, void* ptr_stencil, void* ptr_out, void* ptr_pred); + PY_THRUSTRTC_API int n_copy(void* ptr_in, void* ptr_out); + PY_THRUSTRTC_API int n_swap(void* ptr_vec1, void* ptr_vec2); + + // Redutions + PY_THRUSTRTC_API unsigned long long n_count(void* ptr_vec, void* ptr_value); + PY_THRUSTRTC_API unsigned long long n_count_if(void* ptr_vec, void* ptr_pred); + PY_THRUSTRTC_API void* n_reduce(void* ptr_vec, void* ptr_init, void* ptr_bin_op); + PY_THRUSTRTC_API unsigned n_reduce_by_key(void* ptr_key_in, void* ptr_value_in, void* ptr_key_out, void* ptr_value_out, void* ptr_binary_pred, void* ptr_binary_op); + PY_THRUSTRTC_API int n_equal(void* ptr_vec1, void* ptr_vec2, void* ptr_binary_pred); + PY_THRUSTRTC_API unsigned long long n_min_element(void* ptr_vec, void* ptr_comp); + PY_THRUSTRTC_API unsigned long long n_max_element(void* ptr_vec, void* ptr_comp); + PY_THRUSTRTC_API int n_minmax_element(void* ptr_vec, void* ptr_comp, unsigned long long* ret); + PY_THRUSTRTC_API void* n_inner_product(void* ptr_vec1, void* ptr_vec2, void* ptr_init, void* ptr_binary_op1, void* ptr_binary_op2); + PY_THRUSTRTC_API void* n_transform_reduce(void* ptr_vec, void* ptr_unary_op, void* ptr_init, void* ptr_binary_op); + PY_THRUSTRTC_API int n_all_of(void* ptr_vec, void* ptr_pred); + PY_THRUSTRTC_API int n_any_of(void* ptr_vec, void* ptr_pred); + PY_THRUSTRTC_API int n_none_of(void* ptr_vec, void* ptr_pred); + PY_THRUSTRTC_API int n_is_partitioned(void* ptr_vec, void* ptr_pred); + PY_THRUSTRTC_API int n_is_sorted(void* ptr_vec, void* ptr_comp); + + // PrefixSums + PY_THRUSTRTC_API int n_inclusive_scan(void* ptr_vec_in, void* ptr_vec_out, void* ptr_binary_op); + PY_THRUSTRTC_API int n_exclusive_scan(void* ptr_vec_in, void* ptr_vec_out, void* ptr_init, void* ptr_binary_op); + PY_THRUSTRTC_API int n_inclusive_scan_by_key(void* ptr_vec_key, void* ptr_vec_value, void* ptr_vec_out, void* ptr_binary_pred, void* ptr_binary_op); + PY_THRUSTRTC_API int n_exclusive_scan_by_key(void* ptr_vec_key, void* ptr_vec_value, void* ptr_vec_out, void* ptr_init, void* ptr_binary_pred, void* ptr_binary_op); + PY_THRUSTRTC_API int n_transform_inclusive_scan(void* ptr_vec_in, void* ptr_vec_out, void* ptr_unary_op, void* ptr_binary_op); + PY_THRUSTRTC_API int n_transform_exclusive_scan(void* ptr_vec_in, void* ptr_vec_out, void* ptr_unary_op, void* ptr_init, void* ptr_binary_op); + + // Reordering + PY_THRUSTRTC_API unsigned n_copy_if(void* ptr_in, void* ptr_out, void* ptr_pred); + PY_THRUSTRTC_API unsigned n_copy_if_stencil(void* ptr_in, void* ptr_stencil, void* ptr_out, void* ptr_pred); + PY_THRUSTRTC_API unsigned n_remove(void* ptr_vec, void* ptr_value); + PY_THRUSTRTC_API unsigned n_remove_copy(void* ptr_in, void* ptr_out, void* ptr_value); + PY_THRUSTRTC_API unsigned n_remove_if(void* ptr_vec, void* ptr_pred); + PY_THRUSTRTC_API unsigned n_remove_copy_if(void* ptr_in, void* ptr_out, void* ptr_pred); + PY_THRUSTRTC_API unsigned n_remove_if_stencil(void* ptr_vec, void* ptr_stencil, void* ptr_pred); + PY_THRUSTRTC_API unsigned n_remove_copy_if_stencil(void* ptr_in, void* ptr_stencil, void* ptr_out, void* ptr_pred); + PY_THRUSTRTC_API unsigned n_unique(void* ptr_vec, void* ptr_binary_pred); + PY_THRUSTRTC_API unsigned n_unique_copy(void* ptr_vec_in, void* ptr_vec_out, void* ptr_binary_pred); + PY_THRUSTRTC_API unsigned n_unique_by_key(void* ptr_keys, void* ptr_values, void* ptr_binary_pred); + PY_THRUSTRTC_API unsigned n_unique_by_key_copy(void* ptr_keys_in, void* ptr_values_in, void* ptr_key_out, void* ptr_values_out, void* ptr_binary_pred); + PY_THRUSTRTC_API unsigned n_partition(void* ptr_vec, void* ptr_pred); + PY_THRUSTRTC_API unsigned n_partition_stencil(void* ptr_vec, void* ptr_stencil, void* ptr_pred); + PY_THRUSTRTC_API unsigned n_partition_copy(void* ptr_vec_in, void* ptr_vec_true, void* ptr_vec_false, void* ptr_pred); + PY_THRUSTRTC_API unsigned n_partition_copy_stencil(void* ptr_vec_in, void* ptr_stencil, void* ptr_vec_true, void* ptr_vec_false, void* ptr_pred); + + // Searching + PY_THRUSTRTC_API long long n_find(void* ptr_vec, void* ptr_value); + PY_THRUSTRTC_API long long n_find_if(void* ptr_vec, void* ptr_pred); + PY_THRUSTRTC_API long long n_find_if_not(void* ptr_vec, void* ptr_pred); + PY_THRUSTRTC_API long long n_mismatch(void* ptr_vec1, void* ptr_vec2, void* ptr_pred); + PY_THRUSTRTC_API unsigned long long n_lower_bound(void* ptr_vec, void* ptr_value, void* ptr_comp); + PY_THRUSTRTC_API unsigned long long n_upper_bound(void* ptr_vec, void* ptr_value, void* ptr_comp); + PY_THRUSTRTC_API int n_binary_search(void* ptr_vec, void* ptr_value, void* ptr_comp); + PY_THRUSTRTC_API int n_lower_bound_v(void* ptr_vec, void* ptr_values, void* ptr_result, void* ptr_comp); + PY_THRUSTRTC_API int n_upper_bound_v(void* ptr_vec, void* ptr_values, void* ptr_result, void* ptr_comp); + PY_THRUSTRTC_API int n_binary_search_v(void* ptr_vec, void* ptr_values, void* ptr_result, void* ptr_comp); + PY_THRUSTRTC_API unsigned long long n_partition_point(void* ptr_vec, void* ptr_pred); + PY_THRUSTRTC_API unsigned long long n_is_sorted_until(void* ptr_vec, void* ptr_comp); + + // Merging + PY_THRUSTRTC_API int n_merge(void* ptr_vec1, void* ptr_vec2, void* ptr_vec_out, void* ptr_comp); + PY_THRUSTRTC_API int n_merge_by_key(void* ptr_keys1, void* ptr_keys2, void* ptr_value1, void* ptr_value2, void* ptr_keys_out, void* ptr_value_out, void* ptr_comp); + + // Sorting + PY_THRUSTRTC_API int n_sort(void* ptr_vec, void* ptr_comp); + PY_THRUSTRTC_API int n_sort_by_key(void* ptr_keys, void* ptr_values, void* ptr_comp); + +} + diff --git a/python/api_Context.cpp b/python/api_Context.cpp new file mode 100644 index 0000000..c1b5904 --- /dev/null +++ b/python/api_Context.cpp @@ -0,0 +1,200 @@ +#include "api.h" +#include "TRTCContext.h" +#include +#include + +typedef std::vector StrArray; +typedef std::vector PtrArray; + +void n_set_libnvrtc_path(const char* path) +{ + set_libnvrtc_path(path); +} + +void n_set_verbose(unsigned verbose) +{ + TRTC_Set_Verbose(verbose!=0); +} + +void n_add_include_dir(const char* dir) +{ + TRTC_Add_Include_Dir(dir); +} + +void n_add_built_in_header(const char* filename, const char* filecontent) +{ + TRTC_Add_Built_In_Header(filename, filecontent); +} + +void n_add_inlcude_filename(const char* fn) +{ + TRTC_Add_Inlcude_Filename(fn); +} + +void n_add_code_block(const char* line) +{ + TRTC_Add_Code_Block(line); +} + +void n_add_constant_object(const char* name, void* cptr) +{ + DeviceViewable* dv = (DeviceViewable*)cptr; + TRTC_Add_Constant_Object(name, *dv); +} + +void n_wait() +{ + TRTC_Wait(); +} + +void* n_kernel_create(void* ptr_param_list, const char* body) +{ + StrArray* param_list = (StrArray*)ptr_param_list; + size_t num_params = param_list->size(); + std::vector params(num_params); + for (size_t i = 0; i < num_params; i++) + params[i] = (*param_list)[i].c_str(); + TRTC_Kernel* cptr = new TRTC_Kernel(params, body); + return cptr; +} + +void n_kernel_destroy(void* cptr) +{ + TRTC_Kernel* kernel = (TRTC_Kernel*)cptr; + delete kernel; +} + +int n_kernel_num_params(void* cptr) +{ + TRTC_Kernel* kernel = (TRTC_Kernel*)cptr; + return (int)kernel->num_params(); +} + + +int n_kernel_calc_optimal_block_size(void* ptr_kernel, void* ptr_arg_list, unsigned sharedMemBytes) +{ + TRTC_Kernel* kernel = (TRTC_Kernel*)ptr_kernel; + size_t num_params = kernel->num_params(); + PtrArray* arg_list = (PtrArray*)ptr_arg_list; + + size_t size = arg_list->size(); + if (num_params != size) + { + printf("Wrong number of arguments received. %d required, %d received.", (int)num_params, (int)size); + return -1; + } + + int sizeBlock; + if (kernel->calc_optimal_block_size(arg_list->data(), sizeBlock, sharedMemBytes)) + return sizeBlock; + else + return -1; +} + + +int n_kernel_calc_number_blocks(void* ptr_kernel, void* ptr_arg_list, int sizeBlock, unsigned sharedMemBytes) +{ + TRTC_Kernel* kernel = (TRTC_Kernel*)ptr_kernel; + size_t num_params = kernel->num_params(); + PtrArray* arg_list = (PtrArray*)ptr_arg_list; + + size_t size = arg_list->size(); + if (num_params != size) + { + printf("Wrong number of arguments received. %d required, %d received.", (int)num_params, (int)size); + return -1; + } + + int numBlocks; + if (kernel->calc_number_blocks(arg_list->data(), sizeBlock, numBlocks, sharedMemBytes)) + return numBlocks; + else + return -1; +} + +int n_kernel_launch(void* ptr_kernel, void* ptr_gridDim, void* ptr_blockDim, void* ptr_arg_list, int sharedMemBytes) +{ + TRTC_Kernel* kernel = (TRTC_Kernel*)ptr_kernel; + size_t num_params = kernel->num_params(); + + dim_type* gridDim = (dim_type*)ptr_gridDim; + dim_type* blockDim = (dim_type*)ptr_blockDim; + + PtrArray* arg_list = (PtrArray*)ptr_arg_list; + + size_t size = arg_list->size(); + if (num_params != size) + { + printf("Wrong number of arguments received. %d required, %d received.", (int)num_params, (int)size); + return -1; + } + + if (kernel->launch(*gridDim, *blockDim, arg_list->data(), sharedMemBytes)) + return 0; + else + return -1; +} + +void* n_for_create(void* ptr_param_list, const char* name_iter, const char* body) +{ + StrArray* param_list = (StrArray*)ptr_param_list; + size_t num_params = param_list->size(); + std::vector params(num_params); + for (size_t i = 0; i < num_params; i++) + params[i] = (*param_list)[i].c_str(); + TRTC_For* cptr = new TRTC_For(params, name_iter, body); + return cptr; +} + + +void n_for_destroy(void* cptr) +{ + TRTC_For* p_for = (TRTC_For*)cptr; + delete p_for; +} + +int n_for_num_params(void* cptr) +{ + TRTC_For* p_for = (TRTC_For*)cptr; + return (int)p_for->num_params(); +} + +int n_for_launch(void* cptr, int begin, int end, void* ptr_arg_list) +{ + TRTC_For* p_for = (TRTC_For*)cptr; + size_t num_params = p_for->num_params(); + + PtrArray* arg_list = (PtrArray*)ptr_arg_list; + + size_t size = arg_list->size(); + if (num_params != size) + { + printf("Wrong number of arguments received. %d required, %d received.", (int)num_params, (int)size); + return -1; + } + + if (p_for->launch(begin, end, arg_list->data())) + return 0; + else + return -1; +} + +int n_for_launch_n(void* cptr, int n, void* ptr_arg_list) +{ + TRTC_For* p_for = (TRTC_For*)cptr; + size_t num_params = p_for->num_params(); + + PtrArray* arg_list = (PtrArray*)ptr_arg_list; + + size_t size = arg_list->size(); + if (num_params != size) + { + printf("Wrong number of arguments received. %d required, %d received.", (int)num_params, (int)size); + return -1; + } + + if (p_for->launch_n(n, arg_list->data())) + return 0; + else + return -1; +} diff --git a/python/api_Copying.cpp b/python/api_Copying.cpp new file mode 100644 index 0000000..774f813 --- /dev/null +++ b/python/api_Copying.cpp @@ -0,0 +1,96 @@ +#include "api.h" +#include "TRTCContext.h" +#include "gather.h" +#include "scatter.h" +#include "copy.h" +#include "swap.h" + +int n_gather(void* ptr_map, void* ptr_in, void* ptr_out) +{ + DVVectorLike* vec_map = (DVVectorLike*)ptr_map; + DVVectorLike* vec_in = (DVVectorLike*)ptr_in; + DVVectorLike* vec_out = (DVVectorLike*)ptr_out; + if (TRTC_Gather(*vec_map, *vec_in, *vec_out)) + return 0; + else + return -1; +} + +int n_gather_if(void* ptr_map, void* ptr_stencil, void* ptr_in, void* ptr_out, void* ptr_pred) +{ + DVVectorLike* vec_map = (DVVectorLike*)ptr_map; + DVVectorLike* vec_stencil = (DVVectorLike*)ptr_stencil; + DVVectorLike* vec_in = (DVVectorLike*)ptr_in; + DVVectorLike* vec_out = (DVVectorLike*)ptr_out; + Functor* pred = (Functor*)ptr_pred; + + if (pred == nullptr) + { + if (TRTC_Gather_If(*vec_map, *vec_stencil, *vec_in, *vec_out)) + return 0; + else + return -1; + } + else + { + if (TRTC_Gather_If(*vec_map, *vec_stencil, *vec_in, *vec_out, *pred)) + return 0; + else + return -1; + } +} + +int n_scatter(void* ptr_in, void* ptr_map, void* ptr_out) +{ + DVVectorLike* vec_in = (DVVectorLike*)ptr_in; + DVVectorLike* vec_map = (DVVectorLike*)ptr_map; + DVVectorLike* vec_out = (DVVectorLike*)ptr_out; + if (TRTC_Scatter(*vec_in, *vec_map, *vec_out)) + return 0; + else + return -1; +} + +int n_scatter_if(void* ptr_in, void* ptr_map, void* ptr_stencil, void* ptr_out, void* ptr_pred) +{ + DVVectorLike* vec_in = (DVVectorLike*)ptr_in; + DVVectorLike* vec_map = (DVVectorLike*)ptr_map; + DVVectorLike* vec_stencil = (DVVectorLike*)ptr_stencil; + DVVectorLike* vec_out = (DVVectorLike*)ptr_out; + Functor* pred = (Functor*)ptr_pred; + + if (pred == nullptr) + { + if (TRTC_Scatter_If(*vec_in, *vec_map, *vec_stencil, *vec_out)) + return 0; + else + return -1; + } + else + { + if (TRTC_Scatter_If(*vec_in, *vec_map, *vec_stencil, *vec_out, *pred)) + return 0; + else + return -1; + } +} + +int n_copy(void* ptr_in, void* ptr_out) +{ + DVVectorLike* vec_in = (DVVectorLike*)ptr_in; + DVVectorLike* vec_out = (DVVectorLike*)ptr_out; + if (TRTC_Copy(*vec_in, *vec_out)) + return 0; + else + return 1; +} + +int n_swap(void* ptr_vec1, void* ptr_vec2) +{ + DVVectorLike* vec1 = (DVVectorLike*)ptr_vec1; + DVVectorLike* vec2 = (DVVectorLike*)ptr_vec2; + if (TRTC_Swap(*vec1, *vec2)) + return 0; + else + return -1; +} diff --git a/python/api_DVTuple.cpp b/python/api_DVTuple.cpp new file mode 100644 index 0000000..2b00bc3 --- /dev/null +++ b/python/api_DVTuple.cpp @@ -0,0 +1,21 @@ +#include "api.h" +#include "TRTCContext.h" +#include "DVTuple.h" + +typedef std::vector StrArray; +typedef std::vector PtrArray; + +void* n_dvtuple_create(void* ptr_dvs, void* ptr_names) +{ + PtrArray* dvs = (PtrArray*)ptr_dvs; + StrArray* names = (StrArray*)ptr_names; + size_t num_params = dvs->size(); + std::vector arg_map(num_params); + for (size_t i = 0; i < num_params; i++) + { + arg_map[i].obj_name = (*names)[i].c_str(); + arg_map[i].obj = (*dvs)[i]; + } + return new DVTuple(arg_map); +} + diff --git a/python/api_DVVector.cpp b/python/api_DVVector.cpp new file mode 100644 index 0000000..afff613 --- /dev/null +++ b/python/api_DVVector.cpp @@ -0,0 +1,71 @@ +#include "api.h" +#include "DVVector.h" +#include "fake_vectors/DVRange.h" + +typedef std::vector PtrArray; + +const char* n_dvvectorlike_name_elem_cls(void* cptr) +{ + DVVectorLike* dvvec = (DVVectorLike*)cptr; + return dvvec->name_elem_cls().c_str(); +} + +unsigned long long n_dvvectorlike_size(void* cptr) +{ + DVVectorLike* dvvec = (DVVectorLike*)cptr; + return dvvec->size(); +} + +void* n_dvvector_create(const char* elem_cls, unsigned long long size, void* hdata) +{ + return new DVVector(elem_cls, size, hdata); +} + +void n_dvvector_to_host(void* cptr, void* hdata, unsigned long long begin, unsigned long long end) +{ + DVVector* dvvec = (DVVector*)cptr; + dvvec->to_host(hdata, begin, end); +} + +void* n_dvvector_from_dvs(void* ptr_dvs) +{ + PtrArray* dvs = (PtrArray*)ptr_dvs; + size_t num_items = dvs->size(); + if (num_items < 1) return nullptr; + std::string elem_cls = (*dvs)[0]->name_view_cls(); + for (size_t i = 1; i < num_items; i++) + { + if ((*dvs)[i]->name_view_cls() != elem_cls) + return nullptr; + } + size_t elem_size = TRTC_Size_Of(elem_cls.c_str()); + std::vector buf(elem_size*num_items); + for (size_t i = 0; i < num_items; i++) + { + memcpy(buf.data() + elem_size * i, (*dvs)[i]->view().data(), elem_size); + } + return new DVVector(elem_cls.data(), num_items, buf.data()); +} + +void* n_dvvectoradaptor_create(const char* elem_cls, unsigned long long size, void* data) +{ + return new DVVectorAdaptor(elem_cls, size, data); +} + +void* n_dvrange_create(void* ptr_in, unsigned long long begin, unsigned long long end) +{ + DVVectorLike* vec_value = (DVVectorLike*)ptr_in; + DVVector* p_vec = dynamic_cast(vec_value); + if (p_vec) + { + return new DVVectorAdaptor(*p_vec, begin, end); + } + + DVVectorAdaptor* p_vec_adpt = dynamic_cast(vec_value); + if (p_vec_adpt) + { + return new DVVectorAdaptor(*p_vec_adpt, begin, end); + } + + return new DVRange(*vec_value, begin, end); +} diff --git a/python/api_DeviceViewable.cpp b/python/api_DeviceViewable.cpp new file mode 100644 index 0000000..1e3a338 --- /dev/null +++ b/python/api_DeviceViewable.cpp @@ -0,0 +1,135 @@ +#include "api.h" +#include "DeviceViewable.h" + +const char* n_dv_name_view_cls(void* cptr) +{ + DeviceViewable* dv = (DeviceViewable*)cptr; + return dv->name_view_cls().c_str(); +} + +void n_dv_destroy(void* cptr) +{ + DeviceViewable* dv = (DeviceViewable*)cptr; + delete dv; +} + +void* n_dvint8_create(int v) +{ + return new DVInt8((int8_t)v); +} + +int n_dvint8_value(void* cptr) +{ + DeviceViewable* dv = (DeviceViewable*)cptr; + return (int)*(int8_t*)dv->view().data(); +} + +void* n_dvuint8_create(unsigned v) +{ + return new DVUInt8((uint8_t)v); +} + +unsigned n_dvuint8_value(void* cptr) +{ + DeviceViewable* dv = (DeviceViewable*)cptr; + return (unsigned)*(uint8_t*)dv->view().data(); +} + +void* n_dvint16_create(int v) +{ + return new DVInt16((int16_t)v); +} + +int n_dvint16_value(void* cptr) +{ + DeviceViewable* dv = (DeviceViewable*)cptr; + return (int)*(int16_t*)dv->view().data(); +} + +void* n_dvuint16_create(unsigned v) +{ + return new DVUInt16((uint16_t)v); +} + +unsigned n_dvuint16_value(void* cptr) +{ + DeviceViewable* dv = (DeviceViewable*)cptr; + return (unsigned)*(uint16_t*)dv->view().data(); +} + +void* n_dvint32_create(int v) +{ + return new DVInt32(v); +} + +int n_dvint32_value(void* cptr) +{ + DeviceViewable* dv = (DeviceViewable*)cptr; + return *(int32_t*)dv->view().data(); +} + +void* n_dvuint32_create(unsigned v) +{ + return new DVUInt32(v); +} + +unsigned n_dvuint32_value(void* cptr) +{ + DeviceViewable* dv = (DeviceViewable*)cptr; + return *(uint32_t*)dv->view().data(); +} + +void* n_dvint64_create(long long v) +{ + return new DVInt64(v); +} + +long long n_dvint64_value(void* cptr) +{ + DeviceViewable* dv = (DeviceViewable*)cptr; + return *(int64_t*)dv->view().data(); +} + +void* n_dvuint64_create(unsigned long long v) +{ + return new DVUInt64(v); +} + +unsigned long long n_dvuint64_value(void* cptr) +{ + DeviceViewable* dv = (DeviceViewable*)cptr; + return *(uint64_t*)dv->view().data(); +} + +void* n_dvfloat_create(float v) +{ + return new DVFloat(v); +} + +float n_dvfloat_value(void* cptr) +{ + DeviceViewable* dv = (DeviceViewable*)cptr; + return *(float*)dv->view().data(); +} + +void* n_dvdouble_create(double v) +{ + return new DVDouble(v); +} + +double n_dvdouble_value(void* cptr) +{ + DeviceViewable* dv = (DeviceViewable*)cptr; + return *(double*)dv->view().data(); +} + +void* n_dvbool_create(int v) +{ + return new DVBool(v!=0); +} + +int n_dvbool_value(void* cptr) +{ + DeviceViewable* dv = (DeviceViewable*)cptr; + return *(bool*)dv->view().data()?1:0; +} diff --git a/python/api_FakeVectors.cpp b/python/api_FakeVectors.cpp new file mode 100644 index 0000000..58d8ac7 --- /dev/null +++ b/python/api_FakeVectors.cpp @@ -0,0 +1,89 @@ +#include "api.h" +#include "TRTCContext.h" +#include "fake_vectors/DVConstant.h" +#include "fake_vectors/DVCounter.h" +#include "fake_vectors/DVDiscard.h" +#include "fake_vectors/DVPermutation.h" +#include "fake_vectors/DVReverse.h" +#include "fake_vectors/DVTransform.h" +#include "fake_vectors/DVZipped.h" +#include "fake_vectors/DVCustomVector.h" + +typedef std::vector PtrArray; +typedef std::vector StrArray; + +void* n_dvconstant_create(void* cptr, unsigned long long size) +{ + DeviceViewable* dvobj = (DeviceViewable*)cptr; + return new DVConstant(*dvobj, size); +} + +void* n_dvcounter_create(void* cptr, unsigned long long size) +{ + DeviceViewable* dvobj_init = (DeviceViewable*)cptr; + return new DVCounter(*dvobj_init, size); +} + +void* n_dvdiscard_create(const char* elem_cls, unsigned long long size) +{ + return new DVDiscard(elem_cls, size); +} + +void* n_dvpermutation_create(void* ptr_value, void* ptr_index) +{ + DVVectorLike* vec_value = (DVVectorLike*)ptr_value; + DVVectorLike* vec_index = (DVVectorLike*)ptr_index; + return new DVPermutation(*vec_value, *vec_index); +} + +void* n_dvreverse_create(void* ptr_value) +{ + DVVectorLike* vec_value = (DVVectorLike*)ptr_value; + return new DVReverse(*vec_value); +} + +void* n_dvtransform_create(void* ptr_in, const char* elem_cls, void* ptr_op) +{ + DVVectorLike* vec_in = (DVVectorLike*)ptr_in; + Functor* op = (Functor*)ptr_op; + return new DVTransform(*vec_in, elem_cls, *op); +} + +void* n_dvzipped_create(void* ptr_vecs, void* ptr_elem_names) +{ + std::vector* vecs = (std::vector*)ptr_vecs; + size_t num_vecs = vecs->size(); + StrArray* str_elem_names = (StrArray*)ptr_elem_names; + size_t num_elems = str_elem_names->size(); + if (num_elems != num_vecs) + { + printf("Number of vectors %d mismatch with number of element names %d.", (int)num_vecs, (int)num_elems); + return nullptr; + } + std::vector elem_names(num_elems); + for (size_t i = 0; i < num_elems; i++) + elem_names[i] = (*str_elem_names)[i].c_str(); + return new DVZipped(*vecs, elem_names); +} + +void* n_dvcustomvector_create(void* ptr_dvs, void* ptr_names, const char* name_idx, const char* body, const char* elem_cls, unsigned long long size, unsigned read_only) +{ + PtrArray* dvs = (PtrArray*)ptr_dvs; + StrArray* names = (StrArray*)ptr_names; + size_t num_params = dvs->size(); + std::vector arg_map(num_params); + for (size_t i = 0; i < num_params; i++) + { + arg_map[i].obj_name = (*names)[i].c_str(); + arg_map[i].obj = (*dvs)[i]; + } + return new DVCustomVector(arg_map, name_idx, body, elem_cls, size, read_only!=0); +} + + + + + + + + diff --git a/python/api_Functor.cpp b/python/api_Functor.cpp new file mode 100644 index 0000000..baff0b6 --- /dev/null +++ b/python/api_Functor.cpp @@ -0,0 +1,34 @@ +#include "api.h" +#include "TRTCContext.h" +#include "functor.h" + +typedef std::vector StrArray; +typedef std::vector PtrArray; + + +void* n_functor_create(void* ptr_dvs, void* ptr_names, void* ptr_functor_params, const char* code_body) +{ + PtrArray* dvs = (PtrArray*)ptr_dvs; + StrArray* names = (StrArray*)ptr_names; + size_t num_params = dvs->size(); + std::vector arg_map(num_params); + for (size_t i = 0; i < num_params; i++) + { + arg_map[i].obj_name = (*names)[i].c_str(); + arg_map[i].obj = (*dvs)[i]; + } + + StrArray* str_functor_params = (StrArray*)ptr_functor_params; + size_t num_functor_params = str_functor_params->size(); + std::vector functor_params(num_functor_params); + for (size_t i = 0; i < num_functor_params; i++) + functor_params[i] = (*str_functor_params)[i].c_str(); + + return new Functor(arg_map, functor_params, code_body); +} + +void* n_built_in_functor_create(const char* name_built_in_view_cls) +{ + return new Functor(name_built_in_view_cls); +} + diff --git a/python/api_Merging.cpp b/python/api_Merging.cpp new file mode 100644 index 0000000..37f1937 --- /dev/null +++ b/python/api_Merging.cpp @@ -0,0 +1,54 @@ +#include "api.h" +#include "TRTCContext.h" +#include "merge.h" + +int n_merge(void* ptr_vec1, void* ptr_vec2, void* ptr_vec_out, void* ptr_comp) +{ + DVVectorLike* vec1 = (DVVectorLike*)ptr_vec1; + DVVectorLike* vec2 = (DVVectorLike*)ptr_vec2; + DVVectorLike* vec_out = (DVVectorLike*)ptr_vec_out; + Functor* comp = (Functor*)ptr_comp; + + if (comp == nullptr) + { + if (TRTC_Merge(*vec1, *vec2, *vec_out)) + return 0; + else + return -1; + } + else + { + if (TRTC_Merge(*vec1, *vec2, *vec_out, *comp)) + return 0; + else + return -1; + } +} + +int n_merge_by_key(void* ptr_keys1, void* ptr_keys2, void* ptr_value1, void* ptr_value2, void* ptr_keys_out, void* ptr_value_out, void* ptr_comp) +{ + + DVVectorLike* keys1 = (DVVectorLike*)ptr_keys1; + DVVectorLike* keys2 = (DVVectorLike*)ptr_keys2; + DVVectorLike* value1 = (DVVectorLike*)ptr_value1; + DVVectorLike* value2 = (DVVectorLike*)ptr_value2; + DVVectorLike* keys_out = (DVVectorLike*)ptr_keys_out; + DVVectorLike* value_out = (DVVectorLike*)ptr_value_out; + Functor* comp = (Functor*)ptr_comp; + + if (comp == nullptr) + { + if (TRTC_Merge_By_Key(*keys1, *keys2, *value1, *value2, *keys_out, *value_out)) + return 0; + else + return -1; + } + else + { + if (TRTC_Merge_By_Key(*keys1, *keys2, *value1, *value2, *keys_out, *value_out, *comp)) + return 0; + else + return -1; + } + +} diff --git a/python/api_PrefixSums.cpp b/python/api_PrefixSums.cpp new file mode 100644 index 0000000..a212c82 --- /dev/null +++ b/python/api_PrefixSums.cpp @@ -0,0 +1,113 @@ +#include "api.h" +#include "TRTCContext.h" +#include "scan.h" +#include "transform_scan.h" + +int n_inclusive_scan(void* ptr_vec_in, void* ptr_vec_out, void* ptr_binary_op) +{ + DVVectorLike* vec_in = (DVVectorLike*)ptr_vec_in; + DVVectorLike* vec_out = (DVVectorLike*)ptr_vec_out; + Functor* binary_op = (Functor*)ptr_binary_op; + + if (binary_op == nullptr) + { + if (!TRTC_Inclusive_Scan(*vec_in, *vec_out)) return -1; + } + else + { + if (!TRTC_Inclusive_Scan(*vec_in, *vec_out, *binary_op)) return -1; + } + return 0; +} + +int n_exclusive_scan(void* ptr_vec_in, void* ptr_vec_out, void* ptr_init, void* ptr_binary_op) +{ + DVVectorLike* vec_in = (DVVectorLike*)ptr_vec_in; + DVVectorLike* vec_out = (DVVectorLike*)ptr_vec_out; + DeviceViewable* init = (DeviceViewable*)ptr_init; + Functor* binary_op = (Functor*)ptr_binary_op; + + if (init == nullptr) + { + if (!TRTC_Exclusive_Scan(*vec_in, *vec_out)) return -1; + } + else if (binary_op == nullptr) + { + if (!TRTC_Exclusive_Scan(*vec_in, *vec_out, *init)) return -1; + } + else + { + if (!TRTC_Exclusive_Scan(*vec_in, *vec_out, *init, *binary_op)) return -1; + } + return 0; +} + +int n_inclusive_scan_by_key(void* ptr_vec_key, void* ptr_vec_value, void* ptr_vec_out, void* ptr_binary_pred, void* ptr_binary_op) +{ + DVVectorLike* vec_key = (DVVectorLike*)ptr_vec_key; + DVVectorLike* vec_value = (DVVectorLike*)ptr_vec_value; + DVVectorLike* vec_out = (DVVectorLike*)ptr_vec_out; + Functor* binary_pred = (Functor*)ptr_binary_pred; + Functor* binary_op = (Functor*)ptr_binary_op; + + if (binary_pred == nullptr) + { + if (!TRTC_Inclusive_Scan_By_Key(*vec_key, *vec_value, *vec_out)) return -1; + } + else if (binary_op == nullptr) + { + if (!TRTC_Inclusive_Scan_By_Key(*vec_key, *vec_value, *vec_out, *binary_pred)) return -1; + } + else + { + if (!TRTC_Inclusive_Scan_By_Key(*vec_key, *vec_value, *vec_out, *binary_pred, *binary_op)) return -1; + } + return 0; +} + +int n_exclusive_scan_by_key(void* ptr_vec_key, void* ptr_vec_value, void* ptr_vec_out, void* ptr_init, void* ptr_binary_pred, void* ptr_binary_op) +{ + DVVectorLike* vec_key = (DVVectorLike*)ptr_vec_key; + DVVectorLike* vec_value = (DVVectorLike*)ptr_vec_value; + DVVectorLike* vec_out = (DVVectorLike*)ptr_vec_out; + DeviceViewable* init = (DeviceViewable*)ptr_init; + Functor* binary_pred = (Functor*)ptr_binary_pred; + Functor* binary_op = (Functor*)ptr_binary_op; + + if (init == nullptr) + { + if (!TRTC_Exclusive_Scan_By_Key(*vec_key, *vec_value, *vec_out)) return -1; + } + else if (binary_pred == nullptr) + { + if (!TRTC_Exclusive_Scan_By_Key(*vec_key, *vec_value, *vec_out, *init)) return -1; + } + else if (binary_op == nullptr) + { + if (!TRTC_Exclusive_Scan_By_Key(*vec_key, *vec_value, *vec_out, *init, *binary_pred)) return -1; + } + else + { + if (!TRTC_Exclusive_Scan_By_Key(*vec_key, *vec_value, *vec_out, *init, *binary_pred, *binary_op)) return -1; + } + return 0; +} + +int n_transform_inclusive_scan(void* ptr_vec_in, void* ptr_vec_out, void* ptr_unary_op, void* ptr_binary_op) +{ + DVVectorLike* vec_in = (DVVectorLike*)ptr_vec_in; + DVVectorLike* vec_out = (DVVectorLike*)ptr_vec_out; + Functor* unary_op = (Functor*)ptr_unary_op; + Functor* binary_op = (Functor*)ptr_binary_op; + return TRTC_Transform_Inclusive_Scan(*vec_in, *vec_out, *unary_op, *binary_op) ? 0 : -1; +} + +int n_transform_exclusive_scan(void* ptr_vec_in, void* ptr_vec_out, void* ptr_unary_op, void* ptr_init, void* ptr_binary_op) +{ + DVVectorLike* vec_in = (DVVectorLike*)ptr_vec_in; + DVVectorLike* vec_out = (DVVectorLike*)ptr_vec_out; + Functor* unary_op = (Functor*)ptr_unary_op; + DeviceViewable* init = (DeviceViewable*)ptr_init; + Functor* binary_op = (Functor*)ptr_binary_op; + return TRTC_Transform_Exclusive_Scan(*vec_in, *vec_out, *unary_op, *init, *binary_op) ? 0 : 1; +} diff --git a/python/api_Reductions.cpp b/python/api_Reductions.cpp new file mode 100644 index 0000000..cc48fda --- /dev/null +++ b/python/api_Reductions.cpp @@ -0,0 +1,256 @@ +#include "api.h" +#include "TRTCContext.h" +#include "count.h" +#include "reduce.h" +#include "equal.h" +#include "extrema.h" +#include "inner_product.h" +#include "transform_reduce.h" +#include "logical.h" +#include "partition.h" +#include "sort.h" + +unsigned long long n_count(void* ptr_vec, void* ptr_value) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + DeviceViewable* value = (DeviceViewable*)ptr_value; + size_t res; + if (TRTC_Count(*vec, *value, res)) + return (unsigned long long)res; + else + return (unsigned long long)(-1); + +} + +unsigned long long n_count_if(void* ptr_vec, void* ptr_pred) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + Functor* pred = (Functor*)ptr_pred; + size_t res; + if (TRTC_Count_If(*vec, *pred, res)) + return (unsigned long long)res; + else + return (unsigned long long)(-1); +} + +void* n_reduce(void* ptr_vec, void* ptr_init, void* ptr_bin_op) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + DeviceViewable* init = (DeviceViewable*)ptr_init; + Functor* binary_op = (Functor*)ptr_bin_op; + + ViewBuf ret; + if (init == nullptr) + { + if (!TRTC_Reduce(*vec, ret)) + return nullptr; + } + else if (binary_op == nullptr) + { + if (!TRTC_Reduce(*vec, *init, ret)) + return nullptr; + } + else + { + if (!TRTC_Reduce(*vec, *init, *binary_op, ret)) + return nullptr; + } + return dv_from_viewbuf(ret, vec->name_elem_cls().c_str()); +} + +unsigned n_reduce_by_key(void* ptr_key_in, void* ptr_value_in, void* ptr_key_out, void* ptr_value_out, void* ptr_binary_pred, void* ptr_binary_op) +{ + DVVectorLike* key_in = (DVVectorLike*)ptr_key_in; + DVVectorLike* value_in = (DVVectorLike*)ptr_value_in; + DVVectorLike* key_out = (DVVectorLike*)ptr_key_out; + DVVectorLike* value_out = (DVVectorLike*)ptr_value_out; + Functor* binary_pred = (Functor*)ptr_binary_pred;; + Functor* binary_op = (Functor*)ptr_binary_op; + + if (binary_pred == nullptr) + { + return TRTC_Reduce_By_Key(*key_in, *value_in, *key_out, *value_out); + } + else if (binary_op == nullptr) + { + return TRTC_Reduce_By_Key(*key_in, *value_in, *key_out, *value_out, *binary_pred); + } + else + { + return TRTC_Reduce_By_Key(*key_in, *value_in, *key_out, *value_out, *binary_pred, *binary_op); + } +} + +int n_equal(void* ptr_vec1, void* ptr_vec2, void* ptr_binary_pred) +{ + DVVectorLike* vec1 = (DVVectorLike*)ptr_vec1; + DVVectorLike* vec2 = (DVVectorLike*)ptr_vec2; + Functor* binary_pred = (Functor*)ptr_binary_pred; + + bool res; + if (binary_pred == nullptr) + if (TRTC_Equal(*vec1, *vec2, res)) + return res ? 1 : 0; + else + return -1; + else + if (TRTC_Equal(*vec1, *vec2, *binary_pred, res)) + return res ? 1 : 0; + else + return -1; +} + +unsigned long long n_min_element(void* ptr_vec, void* ptr_comp) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + Functor* comp = (Functor*)ptr_comp; + + size_t id_min; + if (comp == nullptr) + if (TRTC_Min_Element(*vec, id_min)) + return (unsigned long long)id_min; + else + return (unsigned long long)(-1); + else + if (TRTC_Min_Element(*vec, *comp, id_min)) + return (unsigned long long)id_min; + else + return (unsigned long long)(-1); +} + + +unsigned long long n_max_element(void* ptr_vec, void* ptr_comp) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + Functor* comp = (Functor*)ptr_comp; + + size_t id_max; + if (comp == nullptr) + if (TRTC_Max_Element(*vec, id_max)) + return (unsigned long long)id_max; + else + return (unsigned long long)(-1); + else + if (TRTC_Max_Element(*vec, *comp, id_max)) + return (unsigned long long)id_max; + else + return (unsigned long long)(-1); +} + + +int n_minmax_element(void* ptr_vec, void* ptr_comp, unsigned long long* ret) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + Functor* comp = (Functor*)ptr_comp; + + size_t id_min, id_max; + if (comp == nullptr) + { + if (!TRTC_MinMax_Element(*vec, id_min, id_max)) return -1; + } + else + { + if (!TRTC_MinMax_Element(*vec, *comp, id_min, id_max)) return -1; + } + + ret[0] = id_min; + ret[1] = id_max; + return 0; +} + +void* n_inner_product(void* ptr_vec1, void* ptr_vec2, void* ptr_init, void* ptr_binary_op1, void* ptr_binary_op2) +{ + DVVectorLike* vec1 = (DVVectorLike*)ptr_vec1; + DVVectorLike* vec2 = (DVVectorLike*)ptr_vec2; + DeviceViewable* init = (DeviceViewable*)ptr_init; + + Functor* binary_op1 = (Functor*)ptr_binary_op1; + Functor* binary_op2 = (Functor*)ptr_binary_op2; + + ViewBuf ret; + if (binary_op1==nullptr || binary_op2 == nullptr) + { + if (!TRTC_Inner_Product(*vec1, *vec2, *init, ret)) return nullptr; + } + else + { + if (!TRTC_Inner_Product(*vec1, *vec2, *init, ret, *binary_op1, *binary_op2)) return nullptr; + } + return dv_from_viewbuf(ret, init->name_view_cls().c_str()); +} + +void* n_transform_reduce(void* ptr_vec, void* ptr_unary_op, void* ptr_init, void* ptr_binary_op) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + Functor* unary_op = (Functor*)ptr_unary_op; + DeviceViewable* init = (DeviceViewable*)ptr_init; + Functor* binary_op = (Functor*)ptr_binary_op; + ViewBuf ret; + if (!TRTC_Transform_Reduce(*vec, *unary_op, *init, *binary_op, ret)) return nullptr; + return dv_from_viewbuf(ret, init->name_view_cls().c_str()); +} + +int n_all_of(void* ptr_vec, void* ptr_pred) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + Functor* pred = (Functor*)ptr_pred; + bool res; + if (TRTC_All_Of(*vec, *pred, res)) + return res ? 1 : 0; + else + return -1; +} + +int n_any_of(void* ptr_vec, void* ptr_pred) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + Functor* pred = (Functor*)ptr_pred; + bool res; + if (TRTC_Any_Of(*vec, *pred, res)) + return res ? 1 : 0; + else + return -1; +} + +int n_none_of(void* ptr_vec, void* ptr_pred) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + Functor* pred = (Functor*)ptr_pred; + bool res; + if (TRTC_None_Of(*vec, *pred, res)) + return res ? 1 : 0; + else + return -1; +} + +int n_is_partitioned(void* ptr_vec, void* ptr_pred) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + Functor* pred = (Functor*)ptr_pred; + bool res; + if (TRTC_Is_Partitioned(*vec, *pred, res)) + return res ? 1 : 0; + else + return -1; +} + +int n_is_sorted(void* ptr_vec, void* ptr_comp) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + Functor* comp = (Functor*)ptr_comp; + bool res; + if (comp == nullptr) + { + if (TRTC_Is_Sorted(*vec, res)) + return res ? 1 : 0; + else + return -1; + } + else + { + if (TRTC_Is_Sorted(*vec, *comp, res)) + return res ? 1 : 0; + else + return -1; + } +} diff --git a/python/api_Reordering.cpp b/python/api_Reordering.cpp new file mode 100644 index 0000000..8736357 --- /dev/null +++ b/python/api_Reordering.cpp @@ -0,0 +1,165 @@ +#include "api.h" +#include "TRTCContext.h" +#include "copy.h" +#include "remove.h" +#include "unique.h" +#include "partition.h" + +unsigned n_copy_if(void* ptr_in, void* ptr_out, void* ptr_pred) +{ + DVVectorLike* vec_in = (DVVectorLike*)ptr_in; + DVVectorLike* vec_out = (DVVectorLike*)ptr_out; + Functor* pred = (Functor*)ptr_pred; + return TRTC_Copy_If(*vec_in, *vec_out, *pred); +} + +unsigned n_copy_if_stencil(void* ptr_in, void* ptr_stencil, void* ptr_out, void* ptr_pred) +{ + DVVectorLike* vec_in = (DVVectorLike*)ptr_in; + DVVectorLike* vec_stencil = (DVVectorLike*)ptr_stencil; + DVVectorLike* vec_out = (DVVectorLike*)ptr_out; + Functor* pred = (Functor*)ptr_pred; + return TRTC_Copy_If_Stencil(*vec_in, *vec_stencil, *vec_out, *pred); +} + +unsigned n_remove(void* ptr_vec, void* ptr_value) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + DeviceViewable* value = (DeviceViewable*)ptr_value; + return TRTC_Remove(*vec, *value); +} + +unsigned n_remove_copy(void* ptr_in, void* ptr_out, void* ptr_value) +{ + DVVectorLike* vec_in = (DVVectorLike*)ptr_in; + DVVectorLike* vec_out = (DVVectorLike*)ptr_out; + DeviceViewable* value = (DeviceViewable*)ptr_value; + return TRTC_Remove_Copy(*vec_in, *vec_out, *value); +} + +unsigned n_remove_if(void* ptr_vec, void* ptr_pred) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + Functor* pred = (Functor*)ptr_pred; + return TRTC_Remove_If(*vec, *pred); +} + +unsigned n_remove_copy_if(void* ptr_in, void* ptr_out, void* ptr_pred) +{ + DVVectorLike* vec_in = (DVVectorLike*)ptr_in; + DVVectorLike* vec_out = (DVVectorLike*)ptr_out; + Functor* pred = (Functor*)ptr_pred; + return TRTC_Remove_Copy_If(*vec_in, *vec_out, *pred); +} + +unsigned n_remove_if_stencil(void* ptr_vec, void* ptr_stencil, void* ptr_pred) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + DVVectorLike* stencil = (DVVectorLike*)ptr_stencil; + Functor* pred = (Functor*)ptr_pred; + return TRTC_Remove_If_Stencil(*vec, *stencil, *pred); +} + +unsigned n_remove_copy_if_stencil(void* ptr_in, void* ptr_stencil, void* ptr_out, void* ptr_pred) +{ + DVVectorLike* vec_in = (DVVectorLike*)ptr_in; + DVVectorLike* stencil = (DVVectorLike*)ptr_stencil; + DVVectorLike* vec_out = (DVVectorLike*)ptr_out; + Functor* pred = (Functor*)ptr_pred; + return TRTC_Remove_Copy_If_Stencil(*vec_in, *stencil, *vec_out, *pred); +} + +unsigned n_unique(void* ptr_vec, void* ptr_binary_pred) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + Functor* binary_pred = (Functor*)ptr_binary_pred; + if (binary_pred == nullptr) + { + return TRTC_Unique(*vec); + } + else + { + return TRTC_Unique(*vec, *binary_pred); + } +} + +unsigned n_unique_copy(void* ptr_vec_in, void* ptr_vec_out, void* ptr_binary_pred) +{ + DVVectorLike* vec_in = (DVVectorLike*)ptr_vec_in; + DVVectorLike* vec_out = (DVVectorLike*)ptr_vec_out; + Functor* binary_pred = (Functor*)ptr_binary_pred; + if (binary_pred == nullptr) + { + return TRTC_Unique_Copy(*vec_in, *vec_out); + } + else + { + return TRTC_Unique_Copy(*vec_in, *vec_out, *binary_pred); + } +} + +unsigned n_unique_by_key(void* ptr_keys, void* ptr_values, void* ptr_binary_pred) +{ + DVVectorLike* keys = (DVVectorLike*)ptr_keys; + DVVectorLike* values = (DVVectorLike*)ptr_values; + Functor* binary_pred = (Functor*)ptr_binary_pred; + if (binary_pred == nullptr) + { + return TRTC_Unique_By_Key(*keys, *values); + } + else + { + return TRTC_Unique_By_Key(*keys, *values, *binary_pred); + } +} + +unsigned n_unique_by_key_copy(void* ptr_keys_in, void* ptr_values_in, void* ptr_key_out, void* ptr_values_out, void* ptr_binary_pred) +{ + DVVectorLike* keys_in = (DVVectorLike*)ptr_keys_in; + DVVectorLike* values_in = (DVVectorLike*)ptr_values_in; + DVVectorLike* keys_out = (DVVectorLike*)ptr_values_in; + DVVectorLike* values_out = (DVVectorLike*)ptr_values_out; + Functor* binary_pred = (Functor*)ptr_binary_pred; + if (binary_pred == nullptr) + { + return TRTC_Unique_By_Key_Copy(*keys_in, *values_in, *keys_out, *values_out); + } + else + { + return TRTC_Unique_By_Key_Copy(*keys_in, *values_in, *keys_out, *values_out, *binary_pred); + } +} + +unsigned n_partition(void* ptr_vec, void* ptr_pred) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + Functor* pred = (Functor*)ptr_pred; + return TRTC_Partition(*vec, *pred); +} + +unsigned n_partition_stencil(void* ptr_vec, void* ptr_stencil, void* ptr_pred) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + DVVectorLike* stencil = (DVVectorLike*)ptr_stencil; + Functor* pred = (Functor*)ptr_pred; + return TRTC_Partition_Stencil(*vec, *stencil, *pred); +} + +unsigned n_partition_copy(void* ptr_vec_in, void* ptr_vec_true, void* ptr_vec_false, void* ptr_pred) +{ + DVVectorLike* vec_in = (DVVectorLike*)ptr_vec_in; + DVVectorLike* vec_true = (DVVectorLike*)ptr_vec_true; + DVVectorLike* vec_false = (DVVectorLike*)ptr_vec_false; + Functor* pred = (Functor*)ptr_pred; + return TRTC_Partition_Copy(*vec_in, *vec_true, *vec_false, *pred); +} + +unsigned n_partition_copy_stencil(void* ptr_vec_in, void* ptr_stencil, void* ptr_vec_true, void* ptr_vec_false, void* ptr_pred) +{ + DVVectorLike* vec_in = (DVVectorLike*)ptr_vec_in; + DVVectorLike* stencil = (DVVectorLike*)ptr_stencil; + DVVectorLike* vec_true = (DVVectorLike*)ptr_vec_true; + DVVectorLike* vec_false = (DVVectorLike*)ptr_vec_false; + Functor* pred = (Functor*)ptr_pred; + return TRTC_Partition_Copy_Stencil(*vec_in, *stencil, *vec_true, *vec_false, *pred); +} diff --git a/python/api_Searching.cpp b/python/api_Searching.cpp new file mode 100644 index 0000000..c76f7a2 --- /dev/null +++ b/python/api_Searching.cpp @@ -0,0 +1,232 @@ +#include "api.h" +#include "TRTCContext.h" +#include "find.h" +#include "mismatch.h" +#include "binary_search.h" +#include "partition.h" +#include "sort.h" + +long long n_find(void* ptr_vec, void* ptr_value) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + DeviceViewable* value = (DeviceViewable*)ptr_value; + size_t res; + if (TRTC_Find(*vec, *value, res)) + return (long long)res; + else + return -1; +} + +long long n_find_if(void* ptr_vec, void* ptr_pred) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + Functor* pred = (Functor*)ptr_pred; + size_t res; + if (TRTC_Find_If(*vec, *pred, res)) + return (long long)res; + else + return -1; +} + +long long n_find_if_not(void* ptr_vec, void* ptr_pred) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + Functor* pred = (Functor*)ptr_pred; + size_t res; + if (TRTC_Find_If_Not(*vec, *pred, res)) + return (long long)res; + else + return -1; +} + +long long n_mismatch(void* ptr_vec1, void* ptr_vec2, void* ptr_pred) +{ + DVVectorLike* vec1 = (DVVectorLike*)ptr_vec1; + DVVectorLike* vec2 = (DVVectorLike*)ptr_vec2; + Functor* pred = (Functor*)(DVVectorLike*)ptr_pred; + size_t res; + if (pred == nullptr) + { + if (TRTC_Mismatch(*vec1, *vec2, res)) + return (long long)res; + else + return -1; + } + else + { + if (TRTC_Mismatch(*vec1, *vec2, *pred, res)) + return (long long)res; + else + return -1; + } +} + +unsigned long long n_lower_bound(void* ptr_vec, void* ptr_value, void* ptr_comp) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + DeviceViewable* value = (DeviceViewable*)ptr_value; + Functor* comp = (Functor*)ptr_comp; + if (comp == nullptr) + { + size_t res; + if (TRTC_Lower_Bound(*vec, *value, res)) + return (unsigned long long)res; + else + return (unsigned long long)(-1); + } + else + { + size_t res; + if (TRTC_Lower_Bound(*vec, *value, *comp, res)) + return (unsigned long long)res; + else + return (unsigned long long)(-1); + } +} + +unsigned long long n_upper_bound(void* ptr_vec, void* ptr_value, void* ptr_comp) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + DeviceViewable* value = (DeviceViewable*)ptr_value; + Functor* comp = (Functor*)ptr_comp; + if (comp == nullptr) + { + size_t res; + if (TRTC_Upper_Bound(*vec, *value, res)) + return (unsigned long long)res; + else + return (unsigned long long)(-1); + } + else + { + size_t res; + if (TRTC_Upper_Bound(*vec, *value, *comp, res)) + return (unsigned long long)res; + else + return (unsigned long long)(-1); + } +} + +int n_binary_search(void* ptr_vec, void* ptr_value, void* ptr_comp) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + DeviceViewable* value = (DeviceViewable*)ptr_value; + Functor* comp = (Functor*)ptr_comp; + if (comp == nullptr) + { + bool res; + if (TRTC_Binary_Search(*vec, *value, res)) + return res ? 1 : 0; + else + return -1; + } + else + { + bool res; + if (TRTC_Binary_Search(*vec, *value, *comp, res)) + return res ? 1 : 0; + else + return -1; + } +} + +int n_lower_bound_v(void* ptr_vec, void* ptr_values, void* ptr_result, void* ptr_comp) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + DVVectorLike* values = (DVVectorLike*)ptr_values; + DVVectorLike* result = (DVVectorLike*)ptr_result; + Functor* comp = (Functor*)ptr_comp; + + if (comp == nullptr) + { + if (TRTC_Lower_Bound_V(*vec, *values, *result)) + return 0; + else + return -1; + } + else + { + if (TRTC_Lower_Bound_V(*vec, *values, *result, *comp)) + return 0; + else + return -1; + } +} + +int n_upper_bound_v(void* ptr_vec, void* ptr_values, void* ptr_result, void* ptr_comp) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + DVVectorLike* values = (DVVectorLike*)ptr_values; + DVVectorLike* result = (DVVectorLike*)ptr_result; + Functor* comp = (Functor*)ptr_comp; + + if (comp == nullptr) + { + if (TRTC_Upper_Bound_V(*vec, *values, *result)) + return 0; + else + return -1; + } + else + { + if (TRTC_Upper_Bound_V(*vec, *values, *result, *comp)) + return 0; + else + return -1; + } +} + +int n_binary_search_v(void* ptr_vec, void* ptr_values, void* ptr_result, void* ptr_comp) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + DVVectorLike* values = (DVVectorLike*)ptr_values; + DVVectorLike* result = (DVVectorLike*)ptr_result; + Functor* comp = (Functor*)ptr_comp; + + if (comp == nullptr) + { + if (TRTC_Binary_Search_V(*vec, *values, *result)) + return 0; + else + return -1; + } + else + { + if (TRTC_Binary_Search_V(*vec, *values, *result, *comp)) + return 0; + else + return -1; + } +} + +unsigned long long n_partition_point(void* ptr_vec, void* ptr_pred) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + Functor* pred = (Functor*)ptr_pred; + size_t pp; + if (TRTC_Partition_Point(*vec, *pred, pp)) + return (unsigned long long)pp; + else + return (unsigned long long)(-1); +} + +unsigned long long n_is_sorted_until(void* ptr_vec, void* ptr_comp) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + Functor* comp = (Functor*)ptr_comp; + size_t res; + if (comp == nullptr) + { + if (TRTC_Is_Sorted_Until(*vec, res)) + return (unsigned long long)res; + else + return (unsigned long long)(-1); + } + else + { + if (TRTC_Is_Sorted_Until(*vec, *comp, res)) + return (unsigned long long)res; + else + return (unsigned long long)(-1); + } +} diff --git a/python/api_Sorting.cpp b/python/api_Sorting.cpp new file mode 100644 index 0000000..e0fb336 --- /dev/null +++ b/python/api_Sorting.cpp @@ -0,0 +1,47 @@ +#include "api.h" +#include "TRTCContext.h" +#include "sort.h" + +int n_sort(void* ptr_vec, void* ptr_comp) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + Functor* comp = (Functor*)ptr_comp; + + if (comp == nullptr) + { + if (TRTC_Sort(*vec)) + return 0; + else + return -1; + } + else + { + if (TRTC_Sort(*vec, *comp)) + return 0; + else + return -1; + } +} + +int n_sort_by_key(void* ptr_keys, void* ptr_values, void* ptr_comp) +{ + DVVectorLike* keys = (DVVectorLike*)ptr_keys; + DVVectorLike* values = (DVVectorLike*)ptr_values; + Functor* comp = (Functor*)ptr_comp; + + if (comp == nullptr) + { + if (TRTC_Sort_By_Key(*keys, *values)) + return 0; + else + return -1; + } + else + { + if (TRTC_Sort_By_Key(*keys, *values, *comp)) + return 0; + else + return -1; + } +} + diff --git a/python/api_Transformations.cpp b/python/api_Transformations.cpp new file mode 100644 index 0000000..37fb4a8 --- /dev/null +++ b/python/api_Transformations.cpp @@ -0,0 +1,200 @@ +#include "api.h" +#include "TRTCContext.h" +#include "fill.h" +#include "replace.h" +#include "for_each.h" +#include "adjacent_difference.h" +#include "sequence.h" +#include "tabulate.h" +#include "transform.h" + +int n_fill(void* ptr_vec, void* ptr_value) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + DeviceViewable* value = (DeviceViewable*)ptr_value; + if (TRTC_Fill(*vec, *value)) + return 0; + else + return -1; +} + +int n_replace(void* ptr_vec, void* ptr_old_value, void* ptr_new_value) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + DeviceViewable* old_value = (DeviceViewable*)ptr_old_value; + DeviceViewable* new_value = (DeviceViewable*)ptr_new_value; + if (TRTC_Replace(*vec, *old_value, *new_value)) + return 0; + else + return -1; +} + +int n_replace_if(void* ptr_vec, void* p_pred, void* ptr_new_value) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + Functor* pred = (Functor*)p_pred; + DeviceViewable* new_value = (DeviceViewable*)ptr_new_value; + if (TRTC_Replace_If(*vec, *pred, *new_value)) + return 0; + else + return -1; +} + +int n_replace_copy(void* ptr_in, void* ptr_out, void* ptr_old_value, void* ptr_new_value) +{ + DVVectorLike* vec_in = (DVVectorLike*)ptr_in; + DVVectorLike* vec_out = (DVVectorLike*)ptr_out; + DeviceViewable* old_value = (DeviceViewable*)ptr_old_value; + DeviceViewable* new_value = (DeviceViewable*)ptr_new_value; + if (TRTC_Replace_Copy(*vec_in, *vec_out, *old_value, *new_value)) + return 0; + else + return -1; +} + +int n_replace_copy_if(void* ptr_in, void* ptr_out, void* p_pred, void* ptr_new_value) +{ + DVVectorLike* vec_in = (DVVectorLike*)ptr_in; + DVVectorLike* vec_out = (DVVectorLike*)ptr_out; + Functor* pred = (Functor*)p_pred; + DeviceViewable* new_value = (DeviceViewable*)ptr_new_value; + if (TRTC_Replace_Copy_If(*vec_in, *vec_out, *pred, *new_value)) + return 0; + else + return -1; +} + +int n_for_each(void* ptr_vec, void* ptr_f) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + Functor* f = (Functor*)ptr_f; + if (TRTC_For_Each(*vec, *f)) + return 0; + else + return -1; +} + +int n_adjacent_difference(void* ptr_in, void* ptr_out, void* ptr_binary_op) +{ + DVVectorLike* vec_in = (DVVectorLike*)ptr_in; + DVVectorLike* vec_out = (DVVectorLike*)ptr_out; + Functor* binary_op = (Functor*)ptr_binary_op; + + if (binary_op == nullptr) + { + if (TRTC_Adjacent_Difference(*vec_in, *vec_out)) + return 0; + else + return -1; + } + else + { + if (TRTC_Adjacent_Difference(*vec_in, *vec_out, *binary_op)) + return 0; + else + return -1; + } +} + +int n_sequence(void* ptr_vec, void* ptr_value_init, void* ptr_value_step) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + DeviceViewable* value_init = (DeviceViewable*)ptr_value_init; + DeviceViewable* value_step = (DeviceViewable*)ptr_value_step; + if (value_init == nullptr) + { + if (TRTC_Sequence(*vec)) + return 0; + else + return -1; + } + else if (value_step == nullptr) + { + if (TRTC_Sequence(*vec, *value_init)) + return 0; + else + return -1; + } + else + { + if (TRTC_Sequence(*vec, *value_init, *value_step)) + return 0; + else + return -1; + } +} + +int n_tabulate(void* ptr_vec, void* ptr_op) +{ + DVVectorLike* vec = (DVVectorLike*)ptr_vec; + Functor* op = (Functor*)ptr_op; + if (TRTC_Tabulate(*vec, *op)) + return 0; + else + return -1; +} + + +int n_transform(void* ptr_in, void* ptr_out, void* ptr_op) +{ + DVVectorLike* vec_in = (DVVectorLike*)ptr_in; + DVVectorLike* vec_out = (DVVectorLike*)ptr_out; + Functor* op = (Functor*)ptr_op; + if (TRTC_Transform(*vec_in, *vec_out, *op)) + return 0; + else + return -1; +} + +int n_transform_binary(void* ptr_in1, void* ptr_in2, void* ptr_out, void* ptr_op) +{ + DVVectorLike* vec_in1 = (DVVectorLike*)ptr_in1; + DVVectorLike* vec_in2 = (DVVectorLike*)ptr_in2; + DVVectorLike* vec_out = (DVVectorLike*)ptr_out; + Functor* op = (Functor*)ptr_op; + if (TRTC_Transform_Binary(*vec_in1, *vec_in2, *vec_out, *op)) + return 0; + else + return -1; +} + +int n_transform_if(void* ptr_in, void* ptr_out, void* ptr_op, void* ptr_pred) +{ + DVVectorLike* vec_in = (DVVectorLike*)ptr_in; + DVVectorLike* vec_out = (DVVectorLike*)ptr_out; + Functor* op = (Functor*)ptr_op; + Functor* pred = (Functor*)ptr_pred; + TRTC_Transform_If(*vec_in, *vec_out, *op, *pred); + return 0; +} + +int n_transform_if_stencil(void* ptr_in, void* ptr_stencil, void* ptr_out, void* ptr_op, void* ptr_pred) +{ + DVVectorLike* vec_in = (DVVectorLike*)ptr_in; + DVVectorLike* vec_stencil = (DVVectorLike*)ptr_stencil; + DVVectorLike* vec_out = (DVVectorLike*)ptr_out; + Functor* op = (Functor*)ptr_op; + Functor* pred = (Functor*)ptr_pred; + if (TRTC_Transform_If_Stencil(*vec_in, *vec_stencil, *vec_out, *op, *pred)) + return 0; + else + return -1; +} + +int n_transform_binary_if_stencil(void* ptr_in1, void* ptr_in2, void* ptr_stencil, void* ptr_out, void* ptr_op, void* ptr_pred) +{ + DVVectorLike* vec_in1 = (DVVectorLike*)ptr_in1; + DVVectorLike* vec_in2 = (DVVectorLike*)ptr_in2; + DVVectorLike* vec_stencil = (DVVectorLike*)ptr_stencil; + DVVectorLike* vec_out = (DVVectorLike*)ptr_out; + Functor* op = (Functor*)ptr_op; + Functor* pred = (Functor*)ptr_pred; + if (TRTC_Transform_Binary_If_Stencil(*vec_in1, *vec_in2, *vec_stencil, *vec_out, *op, *pred)) + return 0; + else + return -1; +} + + + + diff --git a/python/api_utils.cpp b/python/api_utils.cpp new file mode 100644 index 0000000..a7fc6b8 --- /dev/null +++ b/python/api_utils.cpp @@ -0,0 +1,50 @@ +#include "api.h" +#include "TRTCContext.h" +#include "DeviceViewable.h" +#include +#include +#include + +typedef std::vector StrArray; +typedef std::vector PtrArray; + +void* n_string_array_create(unsigned long long size, const char* const* strs) +{ + StrArray* ret = new StrArray(size); + for (size_t i = 0; i < size; i++) + (*ret)[i] = strs[i]; + + return ret; +} + +void n_string_array_destroy(void* ptr_arr) +{ + StrArray* arr = (StrArray*)ptr_arr; + delete arr; +} + +void* n_pointer_array_create(unsigned long long size, const void* const* ptrs) +{ + PtrArray* ret = new PtrArray(size); + memcpy(ret->data(), ptrs, sizeof(void*)*size); + return ret; +} + +void n_pointer_array_destroy(void* ptr_arr) +{ + PtrArray* arr = (PtrArray*)ptr_arr; + delete arr; +} + +void* n_dim3_create(unsigned x, unsigned y, unsigned z) +{ + dim_type* ret = new dim_type({ x,y,z }); + return ret; +} + +void n_dim3_destroy(void* cptr) +{ + dim_type* v = (dim_type*)cptr; + delete v; +} + diff --git a/python/binary_search.hpp b/python/binary_search.hpp deleted file mode 100644 index 66f8af6..0000000 --- a/python/binary_search.hpp +++ /dev/null @@ -1,160 +0,0 @@ -#include -#include "TRTCContext.h" -#include "binary_search.h" - -static PyObject* n_lower_bound(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DeviceViewable* value = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - Functor* comp = nullptr; - PyObject* py_comp = PyTuple_GetItem(args, 2); - if (py_comp != Py_None) - comp = (Functor*)PyLong_AsVoidPtr(py_comp); - if (comp == nullptr) - { - size_t res; - if (TRTC_Lower_Bound(*vec, *value, res)) - return PyLong_FromUnsignedLongLong((unsigned long long)res); - else - Py_RETURN_NONE; - } - else - { - size_t res; - if (TRTC_Lower_Bound(*vec, *value, *comp, res)) - return PyLong_FromUnsignedLongLong((unsigned long long)res); - else - Py_RETURN_NONE; - } -} - -static PyObject* n_upper_bound(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DeviceViewable* value = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - Functor* comp = nullptr; - PyObject* py_comp = PyTuple_GetItem(args, 2); - if (py_comp != Py_None) - comp = (Functor*)PyLong_AsVoidPtr(py_comp); - if (comp == nullptr) - { - size_t res; - if (TRTC_Upper_Bound(*vec, *value, res)) - return PyLong_FromUnsignedLongLong((unsigned long long)res); - else - Py_RETURN_NONE; - } - else - { - size_t res; - if (TRTC_Upper_Bound(*vec, *value, *comp, res)) - return PyLong_FromUnsignedLongLong((unsigned long long)res); - else - Py_RETURN_NONE; - } -} - -static PyObject* n_binary_search(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DeviceViewable* value = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - Functor* comp = nullptr; - PyObject* py_comp = PyTuple_GetItem(args, 2); - if (py_comp != Py_None) - comp = (Functor*)PyLong_AsVoidPtr(py_comp); - if (comp == nullptr) - { - bool res; - if (TRTC_Binary_Search(*vec, *value, res)) - return PyBool_FromLong(res?(long)1:(long)0); - else - Py_RETURN_NONE; - } - else - { - bool res; - if (TRTC_Binary_Search(*vec, *value, *comp, res)) - return PyBool_FromLong(res ? (long)1 : (long)0); - else - Py_RETURN_NONE; - } -} - -static PyObject* n_lower_bound_v(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* values = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DVVectorLike* result = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - Functor* comp = nullptr; - PyObject* py_comp = PyTuple_GetItem(args, 3); - if (py_comp != Py_None) - comp = (Functor*)PyLong_AsVoidPtr(py_comp); - - if (comp == nullptr) - { - if (TRTC_Lower_Bound_V(*vec, *values, *result)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; - } - else - { - if (TRTC_Lower_Bound_V(*vec, *values, *result, *comp)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; - } -} - -static PyObject* n_upper_bound_v(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* values = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DVVectorLike* result = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - Functor* comp = nullptr; - PyObject* py_comp = PyTuple_GetItem(args, 3); - if (py_comp != Py_None) - comp = (Functor*)PyLong_AsVoidPtr(py_comp); - - if (comp == nullptr) - { - if (TRTC_Upper_Bound_V(*vec, *values, *result)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; - } - else - { - if (TRTC_Upper_Bound_V(*vec, *values, *result, *comp)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; - } -} - -static PyObject* n_binary_search_v(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* values = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DVVectorLike* result = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - Functor* comp = nullptr; - PyObject* py_comp = PyTuple_GetItem(args, 3); - if (py_comp != Py_None) - comp = (Functor*)PyLong_AsVoidPtr(py_comp); - - if (comp == nullptr) - { - if (TRTC_Binary_Search_V(*vec, *values, *result)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; - } - else - { - if (TRTC_Binary_Search_V(*vec, *values, *result, *comp)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; - } -} - diff --git a/python/copy.hpp b/python/copy.hpp deleted file mode 100644 index ec31f6e..0000000 --- a/python/copy.hpp +++ /dev/null @@ -1,38 +0,0 @@ -#include -#include "TRTCContext.h" -#include "copy.h" - -static PyObject* n_copy(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - if (TRTC_Copy(*vec_in, *vec_out)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; -} - -static PyObject* n_copy_if(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - Functor* pred = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - uint32_t res = TRTC_Copy_If(*vec_in, *vec_out, *pred); - if (res != uint32_t(-1)) - return PyLong_FromUnsignedLong((unsigned long)res); - else - Py_RETURN_NONE; -} - -static PyObject* n_copy_if_stencil(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_stencil = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - Functor* pred = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 3)); - uint32_t res = TRTC_Copy_If_Stencil(*vec_in, *vec_stencil, *vec_out, *pred); - if (res != uint32_t(-1)) - return PyLong_FromUnsignedLong((unsigned long)res); - else - Py_RETURN_NONE; -} diff --git a/python/count.hpp b/python/count.hpp deleted file mode 100644 index 7cc658c..0000000 --- a/python/count.hpp +++ /dev/null @@ -1,26 +0,0 @@ -#include -#include "TRTCContext.h" -#include "count.h" - -static PyObject* n_count(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DeviceViewable* value = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - size_t res; - if (TRTC_Count(*vec, *value, res)) - return PyLong_FromUnsignedLongLong((unsigned long long)res); - else - Py_RETURN_NONE; - -} - -static PyObject* n_count_if(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - Functor* pred = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - size_t res; - if (TRTC_Count_If(*vec, *pred, res)) - return PyLong_FromUnsignedLongLong((unsigned long long)res); - else - Py_RETURN_NONE; -} diff --git a/python/equal.hpp b/python/equal.hpp deleted file mode 100644 index 15be115..0000000 --- a/python/equal.hpp +++ /dev/null @@ -1,25 +0,0 @@ -#include -#include "TRTCContext.h" -#include "equal.h" - -static PyObject* n_equal(PyObject* self, PyObject* args) -{ - DVVectorLike* vec1 = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec2 = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - PyObject* py_binary_pred = PyTuple_GetItem(args, 2); - Functor* binary_pred = nullptr; - if (py_binary_pred != Py_None) - binary_pred = (Functor*)PyLong_AsVoidPtr(py_binary_pred); - - bool res; - if (binary_pred == nullptr) - if (TRTC_Equal(*vec1, *vec2, res)) - return PyBool_FromLong(res ? 1 : 0); - else - Py_RETURN_NONE; - else - if (TRTC_Equal(*vec1, *vec2, *binary_pred, res)) - return PyBool_FromLong(res ? 1 : 0); - else - Py_RETURN_NONE; -} diff --git a/python/extrema.hpp b/python/extrema.hpp deleted file mode 100644 index 155289c..0000000 --- a/python/extrema.hpp +++ /dev/null @@ -1,73 +0,0 @@ -#include -#include "TRTCContext.h" -#include "extrema.h" - -static PyObject* n_min_element(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - PyObject* py_comp = PyTuple_GetItem(args, 1); - Functor* comp = nullptr; - if (py_comp != Py_None) - comp = (Functor*)PyLong_AsVoidPtr(py_comp); - - size_t id_min; - if (comp==nullptr) - if (TRTC_Min_Element(*vec, id_min)) - return PyLong_FromUnsignedLongLong((unsigned long long)id_min); - else - Py_RETURN_NONE; - else - if (TRTC_Min_Element(*vec, *comp, id_min)) - return PyLong_FromUnsignedLongLong((unsigned long long)id_min); - else - Py_RETURN_NONE; -} - - -static PyObject* n_max_element(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - PyObject* py_comp = PyTuple_GetItem(args, 1); - Functor* comp = nullptr; - if (py_comp != Py_None) - comp = (Functor*)PyLong_AsVoidPtr(py_comp); - - size_t id_max; - if (comp == nullptr) - if (TRTC_Max_Element(*vec, id_max)) - return PyLong_FromUnsignedLongLong((unsigned long long)id_max); - else - Py_RETURN_NONE; - else - if (TRTC_Max_Element(*vec, *comp, id_max)) - return PyLong_FromUnsignedLongLong((unsigned long long)id_max); - else - Py_RETURN_NONE; -} - - -static PyObject* n_minmax_element(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - PyObject* py_comp = PyTuple_GetItem(args, 1); - Functor* comp = nullptr; - if (py_comp != Py_None) - comp = (Functor*)PyLong_AsVoidPtr(py_comp); - - size_t id_min, id_max; - if (comp == nullptr) - { - if (!TRTC_MinMax_Element(*vec, id_min, id_max)) - Py_RETURN_NONE; - } - else - { - if (!TRTC_MinMax_Element(*vec, *comp, id_min, id_max)) - Py_RETURN_NONE; - } - - PyObject* ret = PyTuple_New(2); - PyTuple_SetItem(ret, 0, PyLong_FromUnsignedLongLong((unsigned long long)id_min)); - PyTuple_SetItem(ret, 1, PyLong_FromUnsignedLongLong((unsigned long long)id_max)); - return ret; -} diff --git a/python/fill.hpp b/python/fill.hpp deleted file mode 100644 index 18fe30e..0000000 --- a/python/fill.hpp +++ /dev/null @@ -1,13 +0,0 @@ -#include -#include "TRTCContext.h" -#include "fill.h" - -static PyObject* n_fill(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DeviceViewable* value = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - if(TRTC_Fill(*vec, *value)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; -} diff --git a/python/find.hpp b/python/find.hpp deleted file mode 100644 index 98a16ea..0000000 --- a/python/find.hpp +++ /dev/null @@ -1,38 +0,0 @@ -#include -#include "TRTCContext.h" -#include "find.h" - -static PyObject* n_find(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DeviceViewable* value = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - size_t res; - if (TRTC_Find(*vec, *value, res)) - return PyLong_FromLongLong((long long)res); - else - Py_RETURN_NONE; -} - -static PyObject* n_find_if(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - Functor* pred = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - size_t res; - if (TRTC_Find_If(*vec, *pred, res)) - return PyLong_FromLongLong((long long)res); - else - Py_RETURN_NONE; -} - - -static PyObject* n_find_if_not(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - Functor* pred = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - size_t res; - if (TRTC_Find_If_Not(*vec, *pred, res)) - return PyLong_FromLongLong((long long)res); - else - Py_RETURN_NONE; -} - diff --git a/python/for_each.hpp b/python/for_each.hpp deleted file mode 100644 index c7b3835..0000000 --- a/python/for_each.hpp +++ /dev/null @@ -1,13 +0,0 @@ -#include -#include "TRTCContext.h" -#include "for_each.h" - -static PyObject* n_for_each(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - Functor* f = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - if (TRTC_For_Each(*vec, *f)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; -} diff --git a/python/functor.hpp b/python/functor.hpp deleted file mode 100644 index 06869d2..0000000 --- a/python/functor.hpp +++ /dev/null @@ -1,30 +0,0 @@ -#include -#include "functor.h" - -static PyObject* n_functor_create(PyObject* self, PyObject* args) -{ - PyObject* pyArgMap = PyTuple_GetItem(args, 0); - ssize_t num_params = PyList_Size(pyArgMap); - std::vector arg_map(num_params); - for (ssize_t i = 0; i < num_params; i++) - { - PyObject* pyCapturedDeviceViewable = PyList_GetItem(pyArgMap, i); - arg_map[i].obj_name = PyUnicode_AsUTF8(PyTuple_GetItem(pyCapturedDeviceViewable, 0)); - arg_map[i].obj = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(pyCapturedDeviceViewable, 1)); - } - PyObject* py_functor_params = PyTuple_GetItem(args, 1); - ssize_t num_functor_params = PyList_Size(py_functor_params); - std::vector functor_params(num_functor_params); - for (ssize_t i = 0; i < num_functor_params; i++) - functor_params[i] = PyUnicode_AsUTF8(PyList_GetItem(py_functor_params, i)); - - PyObject* py_code_body = PyTuple_GetItem(args, 2); - const char* code_body = PyUnicode_AsUTF8(py_code_body); - return PyLong_FromVoidPtr(new Functor(arg_map, functor_params, code_body)); -} - -static PyObject* n_built_in_functor_create(PyObject* self, PyObject* args) -{ - const char* name_built_in_view_cls = PyUnicode_AsUTF8(PyTuple_GetItem(args, 0)); - return PyLong_FromVoidPtr(new Functor(name_built_in_view_cls)); -} diff --git a/python/gather.hpp b/python/gather.hpp deleted file mode 100644 index 33711bf..0000000 --- a/python/gather.hpp +++ /dev/null @@ -1,42 +0,0 @@ -#include -#include "TRTCContext.h" -#include "gather.h" - -static PyObject* n_gather(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_map = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - if (TRTC_Gather(*vec_map, *vec_in, *vec_out)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; -} - -static PyObject* n_gather_if(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_map = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_stencil = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DVVectorLike* vec_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 3)); - PyObject *py_pred = PyTuple_GetItem(args, 4); - Functor* pred = nullptr; - if (py_pred != Py_None) - pred = (Functor*)PyLong_AsVoidPtr(py_pred); - - if (pred == nullptr) - { - if (TRTC_Gather_If(*vec_map, *vec_stencil, *vec_in, *vec_out)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; - } - else - { - if (TRTC_Gather_If(*vec_map, *vec_stencil, *vec_in, *vec_out, *pred)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; - } -} - diff --git a/python/inner_product.hpp b/python/inner_product.hpp deleted file mode 100644 index b099491..0000000 --- a/python/inner_product.hpp +++ /dev/null @@ -1,36 +0,0 @@ -#include -#include "TRTCContext.h" -#include "inner_product.h" -#include "viewbuf_to_python.hpp" - -static PyObject* n_inner_product(PyObject* self, PyObject* args) -{ - DVVectorLike* vec1 = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec2 = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DeviceViewable* init = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - - PyObject* py_binary_op1 = PyTuple_GetItem(args, 3); - Functor* binary_op1 = nullptr; - if (py_binary_op1 != Py_None) - binary_op1 = (Functor*)PyLong_AsVoidPtr(py_binary_op1); - - PyObject* py_binary_op2 = PyTuple_GetItem(args, 4); - Functor* binary_op2 = nullptr; - if (py_binary_op2 != Py_None) - binary_op2 = (Functor*)PyLong_AsVoidPtr(py_binary_op2); - - ViewBuf ret; - if (binary_op1==nullptr || binary_op2 == nullptr) - { - if (TRTC_Inner_Product(*vec1, *vec2, *init, ret)) - return PyValue_FromViewBuf(ret, init->name_view_cls().c_str()); - Py_RETURN_NONE; - } - else - { - if (TRTC_Inner_Product(*vec1, *vec2, *init, ret, *binary_op1, *binary_op2)) - return PyValue_FromViewBuf(ret, init->name_view_cls().c_str()); - Py_RETURN_NONE; - } -} - diff --git a/python/logical.hpp b/python/logical.hpp deleted file mode 100644 index 3cfe53b..0000000 --- a/python/logical.hpp +++ /dev/null @@ -1,36 +0,0 @@ -#include -#include "TRTCContext.h" -#include "logical.h" - -static PyObject* n_all_of(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - Functor* pred = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - bool res; - if (TRTC_All_Of(*vec, *pred, res)) - return PyBool_FromLong(res ? 1 : 0); - else - Py_RETURN_NONE; -} - -static PyObject* n_any_of(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - Functor* pred = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - bool res; - if (TRTC_Any_Of(*vec, *pred, res)) - return PyBool_FromLong(res ? 1 : 0); - else - Py_RETURN_NONE; -} - -static PyObject* n_none_of(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - Functor* pred = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - bool res; - if (TRTC_None_Of(*vec, *pred, res)) - return PyBool_FromLong(res ? 1 : 0); - else - Py_RETURN_NONE; -} diff --git a/python/merge.hpp b/python/merge.hpp deleted file mode 100644 index f121312..0000000 --- a/python/merge.hpp +++ /dev/null @@ -1,60 +0,0 @@ -#include -#include "TRTCContext.h" -#include "merge.h" - -static PyObject* n_merge(PyObject* self, PyObject* args) -{ - DVVectorLike* vec1 = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec2 = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - PyObject* py_comp = PyTuple_GetItem(args, 3); - Functor* comp = nullptr; - if (py_comp != Py_None) - comp = (Functor*)PyLong_AsVoidPtr(py_comp); - - if (comp == nullptr) - { - if (TRTC_Merge(*vec1, *vec2, *vec_out)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; - } - else - { - if (TRTC_Merge(*vec1, *vec2, *vec_out, *comp)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; - } -} - -static PyObject* n_merge_by_key(PyObject* self, PyObject* args) -{ - - DVVectorLike* keys1 = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* keys2 = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DVVectorLike* value1 = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - DVVectorLike* value2 = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 3)); - DVVectorLike* keys_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 4)); - DVVectorLike* value_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 5)); - PyObject* py_comp = PyTuple_GetItem(args, 6); - Functor* comp = nullptr; - if (py_comp != Py_None) - comp = (Functor*)PyLong_AsVoidPtr(py_comp); - - if (comp == nullptr) - { - if (TRTC_Merge_By_Key(*keys1, *keys2, *value1, *value2, *keys_out, *value_out)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; - } - else - { - if (TRTC_Merge_By_Key(*keys1, *keys2, *value1, *value2, *keys_out, *value_out, *comp)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; - } - -} diff --git a/python/mismatch.hpp b/python/mismatch.hpp deleted file mode 100644 index 5cfdbec..0000000 --- a/python/mismatch.hpp +++ /dev/null @@ -1,28 +0,0 @@ -#include -#include "TRTCContext.h" -#include "mismatch.h" - -static PyObject* n_mismatch(PyObject* self, PyObject* args) -{ - DVVectorLike* vec1 = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec2 = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - Functor* pred = nullptr; - PyObject* PyPred = PyTuple_GetItem(args, 2); - if (PyPred != Py_None) - pred = (Functor*)(DVVectorLike*)PyLong_AsVoidPtr(PyPred); - size_t res; - if (pred == nullptr) - { - if (TRTC_Mismatch(*vec1, *vec2, res)) - return PyLong_FromLongLong((long long)res); - else - Py_RETURN_NONE; - } - else - { - if (TRTC_Mismatch(*vec1, *vec2, *pred, res)) - return PyLong_FromLongLong((long long)res); - else - Py_RETURN_NONE; - } -} diff --git a/python/partition.hpp b/python/partition.hpp deleted file mode 100644 index 8585b81..0000000 --- a/python/partition.hpp +++ /dev/null @@ -1,77 +0,0 @@ -#include -#include "TRTCContext.h" -#include "partition.h" - -static PyObject* n_partition(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - Functor* pred = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - uint32_t res = TRTC_Partition(*vec, *pred); - if (res != uint32_t(-1)) - return PyLong_FromUnsignedLong((unsigned long)res); - else - Py_RETURN_NONE; -} - -static PyObject* n_partition_stencil(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* stencil = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - Functor* pred = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - uint32_t res = TRTC_Partition_Stencil(*vec, *stencil, *pred); - if (res != uint32_t(-1)) - return PyLong_FromUnsignedLong((unsigned long)res); - else - Py_RETURN_NONE; -} - -static PyObject* n_partition_copy(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_true = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DVVectorLike* vec_false = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - Functor* pred = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 3)); - uint32_t res = TRTC_Partition_Copy(*vec_in, *vec_true, *vec_false, *pred); - if (res != uint32_t(-1)) - return PyLong_FromUnsignedLong((unsigned long)res); - else - Py_RETURN_NONE; -} - -static PyObject* n_partition_copy_stencil(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* stencil = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DVVectorLike* vec_true = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - DVVectorLike* vec_false = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 3)); - Functor* pred = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 4)); - uint32_t res = TRTC_Partition_Copy_Stencil(*vec_in, *stencil, *vec_true, *vec_false, *pred); - if (res != uint32_t(-1)) - return PyLong_FromUnsignedLong((unsigned long)res); - else - Py_RETURN_NONE; -} - - -static PyObject* n_partition_point(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - Functor* pred = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - size_t pp; - if (TRTC_Partition_Point(*vec, *pred, pp)) - return PyLong_FromUnsignedLongLong((unsigned long long)pp); - else - Py_RETURN_NONE; -} - -static PyObject* n_is_partitioned(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - Functor* pred = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - bool res; - if (TRTC_Is_Partitioned(*vec, *pred, res)) - return PyBool_FromLong(res ? 1 : 0); - else - Py_RETURN_NONE; - -} diff --git a/python/reduce.hpp b/python/reduce.hpp deleted file mode 100644 index 933119d..0000000 --- a/python/reduce.hpp +++ /dev/null @@ -1,79 +0,0 @@ -#include -#include "TRTCContext.h" -#include "reduce.h" -#include "viewbuf_to_python.hpp" - -static PyObject* n_reduce(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - PyObject* pyinit = PyTuple_GetItem(args, 1); - DeviceViewable* init = nullptr; - if (pyinit != Py_None) - init = (DeviceViewable*)PyLong_AsVoidPtr(pyinit); - PyObject* py_binary_op = PyTuple_GetItem(args, 2); - Functor* binary_op = nullptr; - if (py_binary_op != Py_None) - binary_op = (Functor*)PyLong_AsVoidPtr(py_binary_op); - - ViewBuf ret; - if (init == nullptr) - { - if (TRTC_Reduce(*vec, ret)) - return PyValue_FromViewBuf(ret, vec->name_elem_cls().c_str()); - Py_RETURN_NONE; - } - else if (binary_op == nullptr) - { - if (TRTC_Reduce(*vec, *init, ret)) - return PyValue_FromViewBuf(ret, vec->name_elem_cls().c_str()); - Py_RETURN_NONE; - } - else - { - if (TRTC_Reduce(*vec, *init, *binary_op, ret)) - return PyValue_FromViewBuf(ret, vec->name_elem_cls().c_str()); - Py_RETURN_NONE; - } -} - -static PyObject* n_reduce_by_key(PyObject* self, PyObject* args) -{ - DVVectorLike* key_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* value_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DVVectorLike* key_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - DVVectorLike* value_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 3)); - PyObject* py_binary_pred = PyTuple_GetItem(args, 4); - Functor* binary_pred = nullptr; - if (py_binary_pred != Py_None) - binary_pred = (Functor*)PyLong_AsVoidPtr(py_binary_pred); - PyObject* py_binary_op = PyTuple_GetItem(args, 5); - Functor* binary_op = nullptr; - if (py_binary_op != Py_None) - binary_op = (Functor*)PyLong_AsVoidPtr(py_binary_op); - - if (binary_pred == nullptr) - { - uint32_t res = TRTC_Reduce_By_Key(*key_in, *value_in, *key_out, *value_out); - if (res != uint32_t(-1)) - return PyLong_FromUnsignedLong((unsigned long)res); - else - Py_RETURN_NONE; - } - else if (binary_op == nullptr) - { - uint32_t res = TRTC_Reduce_By_Key(*key_in, *value_in, *key_out, *value_out, *binary_pred); - if (res != uint32_t(-1)) - return PyLong_FromUnsignedLong((unsigned long)res); - else - Py_RETURN_NONE; - } - else - { - uint32_t res = TRTC_Reduce_By_Key(*key_in, *value_in, *key_out, *value_out, *binary_pred, *binary_op); - if (res != uint32_t(-1)) - return PyLong_FromUnsignedLong((unsigned long)res); - else - Py_RETURN_NONE; - } -} - diff --git a/python/remove.hpp b/python/remove.hpp deleted file mode 100644 index e65a897..0000000 --- a/python/remove.hpp +++ /dev/null @@ -1,76 +0,0 @@ -#include -#include "TRTCContext.h" -#include "remove.h" - -static PyObject* n_remove(PyObject* self, PyObject* args) -{ - DVVectorLike* vec= (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DeviceViewable* value = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - uint32_t res = TRTC_Remove(*vec, *value); - if (res != uint32_t(-1)) - return PyLong_FromUnsignedLong((unsigned long)res); - else - Py_RETURN_NONE; -} - -static PyObject* n_remove_copy(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DeviceViewable* value = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - uint32_t res = TRTC_Remove_Copy(*vec_in, *vec_out, *value); - if (res != uint32_t(-1)) - return PyLong_FromUnsignedLong((unsigned long)res); - else - Py_RETURN_NONE; -} - -static PyObject* n_remove_if(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - Functor* pred = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - uint32_t res = TRTC_Remove_If(*vec, *pred); - if (res != uint32_t(-1)) - return PyLong_FromUnsignedLong((unsigned long)res); - else - Py_RETURN_NONE; -} - -static PyObject* n_remove_copy_if(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - Functor* pred = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - uint32_t res = TRTC_Remove_Copy_If(*vec_in, *vec_out, *pred); - if (res != uint32_t(-1)) - return PyLong_FromUnsignedLong((unsigned long)res); - else - Py_RETURN_NONE; -} - -static PyObject* n_remove_if_stencil(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* stencil = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - Functor* pred = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - uint32_t res = TRTC_Remove_If_Stencil(*vec, *stencil, *pred); - if (res != uint32_t(-1)) - return PyLong_FromUnsignedLong((unsigned long)res); - else - Py_RETURN_NONE; -} - -static PyObject* n_remove_copy_if_stencil(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* stencil = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - Functor* pred = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 3)); - uint32_t res = TRTC_Remove_Copy_If_Stencil(*vec_in, *stencil, *vec_out, *pred); - if (res != uint32_t(-1)) - return PyLong_FromUnsignedLong((unsigned long)res); - else - Py_RETURN_NONE; -} - - diff --git a/python/replace.hpp b/python/replace.hpp deleted file mode 100644 index 46019d3..0000000 --- a/python/replace.hpp +++ /dev/null @@ -1,49 +0,0 @@ -#include -#include "TRTCContext.h" -#include "replace.h" - -static PyObject* n_replace(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DeviceViewable* old_value = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DeviceViewable* new_value = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - if (TRTC_Replace(*vec, *old_value, *new_value)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; -} - -static PyObject* n_replace_if(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - Functor* pred = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DeviceViewable* new_value = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - if (TRTC_Replace_If(*vec, *pred, *new_value)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; -} - -static PyObject* n_replace_copy(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DeviceViewable* old_value = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - DeviceViewable* new_value = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 3)); - if (TRTC_Replace_Copy(*vec_in, *vec_out, *old_value, *new_value)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; -} - -static PyObject* n_replace_copy_if(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - Functor* pred = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - DeviceViewable* new_value = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 3)); - if(TRTC_Replace_Copy_If(*vec_in, *vec_out, *pred, *new_value)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; -} diff --git a/python/scan.hpp b/python/scan.hpp deleted file mode 100644 index 477d66c..0000000 --- a/python/scan.hpp +++ /dev/null @@ -1,138 +0,0 @@ -#include -#include "TRTCContext.h" -#include "scan.h" - -static PyObject* n_inclusive_scan(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - PyObject* py_binary_op = PyTuple_GetItem(args, 2); - Functor* binary_op = nullptr; - if (py_binary_op != Py_None) - binary_op = (Functor*)PyLong_AsVoidPtr(py_binary_op); - - if (binary_op == nullptr) - { - if (TRTC_Inclusive_Scan(*vec_in, *vec_out)) - return PyLong_FromLong(0); - Py_RETURN_NONE; - } - else - { - if (TRTC_Inclusive_Scan(*vec_in, *vec_out, *binary_op)) - return PyLong_FromLong(0); - Py_RETURN_NONE; - } -} - -static PyObject* n_exclusive_scan(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - PyObject* pyinit = PyTuple_GetItem(args, 2); - DeviceViewable* init = nullptr; - if (pyinit != Py_None) - init = (DeviceViewable*)PyLong_AsVoidPtr(pyinit); - PyObject* py_binary_op = PyTuple_GetItem(args, 3); - Functor* binary_op = nullptr; - if (py_binary_op != Py_None) - binary_op = (Functor*)PyLong_AsVoidPtr(py_binary_op); - - if (init == nullptr) - { - if (TRTC_Exclusive_Scan(*vec_in, *vec_out)) - return PyLong_FromLong(0); - Py_RETURN_NONE; - } - else if (binary_op == nullptr) - { - if (TRTC_Exclusive_Scan(*vec_in, *vec_out, *init)) - return PyLong_FromLong(0); - Py_RETURN_NONE; - } - else - { - if (TRTC_Exclusive_Scan(*vec_in, *vec_out, *init, *binary_op)) - return PyLong_FromLong(0); - Py_RETURN_NONE; - } -} - -static PyObject* n_inclusive_scan_by_key(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_key = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_value = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - PyObject* py_binary_pred = PyTuple_GetItem(args, 3); - Functor* binary_pred = nullptr; - if (py_binary_pred != Py_None) - binary_pred = (Functor*)PyLong_AsVoidPtr(py_binary_pred); - PyObject* py_binary_op = PyTuple_GetItem(args, 4); - Functor* binary_op = nullptr; - if (py_binary_op != Py_None) - binary_op = (Functor*)PyLong_AsVoidPtr(py_binary_op); - - if (binary_pred == nullptr) - { - if (TRTC_Inclusive_Scan_By_Key(*vec_key, *vec_value, *vec_out)) - return PyLong_FromLong(0); - Py_RETURN_NONE; - } - else if (binary_op == nullptr) - { - if (TRTC_Inclusive_Scan_By_Key(*vec_key, *vec_value, *vec_out, *binary_pred)) - return PyLong_FromLong(0); - Py_RETURN_NONE; - } - else - { - if (TRTC_Inclusive_Scan_By_Key(*vec_key, *vec_value, *vec_out, *binary_pred, *binary_op)) - return PyLong_FromLong(0); - Py_RETURN_NONE; - } -} - -static PyObject* n_exclusive_scan_by_key(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_key = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_value = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - PyObject* pyinit = PyTuple_GetItem(args, 3); - DeviceViewable* init = nullptr; - if (pyinit != Py_None) - init = (DeviceViewable*)PyLong_AsVoidPtr(pyinit); - PyObject* py_binary_pred = PyTuple_GetItem(args, 4); - Functor* binary_pred = nullptr; - if (py_binary_pred != Py_None) - binary_pred = (Functor*)PyLong_AsVoidPtr(py_binary_pred); - PyObject* py_binary_op = PyTuple_GetItem(args, 5); - Functor* binary_op = nullptr; - if (py_binary_op != Py_None) - binary_op = (Functor*)PyLong_AsVoidPtr(py_binary_op); - - if (init == nullptr) - { - if (TRTC_Exclusive_Scan_By_Key(*vec_key, *vec_value, *vec_out)) - return PyLong_FromLong(0); - Py_RETURN_NONE; - } - if (binary_pred == nullptr) - { - if (TRTC_Exclusive_Scan_By_Key(*vec_key, *vec_value, *vec_out, *init)) - return PyLong_FromLong(0); - Py_RETURN_NONE; - } - else if (binary_op == nullptr) - { - if (TRTC_Exclusive_Scan_By_Key(*vec_key, *vec_value, *vec_out, *init, *binary_pred)) - return PyLong_FromLong(0); - Py_RETURN_NONE; - } - else - { - if (TRTC_Exclusive_Scan_By_Key(*vec_key, *vec_value, *vec_out, *init, *binary_pred, *binary_op)) - return PyLong_FromLong(0); - Py_RETURN_NONE; - } -} - diff --git a/python/scatter.hpp b/python/scatter.hpp deleted file mode 100644 index d588d3f..0000000 --- a/python/scatter.hpp +++ /dev/null @@ -1,42 +0,0 @@ -#include -#include "TRTCContext.h" -#include "scatter.h" - -static PyObject* n_scatter(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_map = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - if (TRTC_Scatter(*vec_in, *vec_map, *vec_out)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; -} - -static PyObject* n_scatter_if(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_map = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DVVectorLike* vec_stencil = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 3)); - PyObject *py_pred = PyTuple_GetItem(args, 4); - Functor* pred = nullptr; - if (py_pred != Py_None) - pred = (Functor*)PyLong_AsVoidPtr(py_pred); - - if (pred == nullptr) - { - if (TRTC_Scatter_If(*vec_in, *vec_map, *vec_stencil, *vec_out)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; - } - else - { - if (TRTC_Scatter_If(*vec_in, *vec_map, *vec_stencil, *vec_out, *pred)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; - } -} - diff --git a/python/sequence.hpp b/python/sequence.hpp deleted file mode 100644 index 2b9c7a1..0000000 --- a/python/sequence.hpp +++ /dev/null @@ -1,38 +0,0 @@ -#include -#include "TRTCContext.h" -#include "sequence.h" - -static PyObject* n_sequence(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - PyObject* py_value_init = PyTuple_GetItem(args, 1); - DeviceViewable* value_init = nullptr; - if (py_value_init!=Py_None) - value_init = (DeviceViewable*)PyLong_AsVoidPtr(py_value_init); - PyObject* py_value_step = PyTuple_GetItem(args, 2); - DeviceViewable* value_step = nullptr; - if (py_value_step != Py_None) - value_step = (DeviceViewable*)PyLong_AsVoidPtr(py_value_step); - - if (value_init == nullptr) - { - if (TRTC_Sequence(*vec)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; - } - else if (value_step == nullptr) - { - if (TRTC_Sequence(*vec, *value_init)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; - } - else - { - if (TRTC_Sequence(*vec, *value_init, *value_step)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; - } -} diff --git a/python/setup.py b/python/setup.py index 8b2cd49..fd802e6 100644 --- a/python/setup.py +++ b/python/setup.py @@ -1,97 +1,15 @@ -#!/usr/bin/python3 - -from setuptools import setup, Extension +from setuptools import setup from codecs import open import os -import platform here = os.path.abspath(os.path.dirname(__file__)) with open(os.path.join(here, 'README.md'), encoding='utf-8') as f: long_description = f.read() -extra_compile_args=[] -define_macros = [] - -if os.name == 'nt': - define_macros += [ - ('THRUST_RTC_DLL_EXPORT', None), - ('_CRT_SECURE_NO_DEPRECATE', None), - ('_SCL_SECURE_NO_DEPRECATE', None), - ('_CRT_SECURE_NO_WARNINGS', None) - ] -else: - extra_compile_args = ['-std=c++11'] - -sources = [ -'PyThrustRTC.cpp', -'../thirdparty/unqlite/unqlite.c', -'../thirdparty/crc64/crc64.cpp', -'../internal/launch_calc.cpp', -'../internal/cuda_wrapper.cpp', -'../internal/nvtrc_wrapper.cpp', -'../internal/general_reduce.cpp', -'../internal/general_scan.cpp', -'../internal/general_scan_by_key.cpp', -'../internal/general_copy_if.cpp', -'../internal/general_find.cpp', -'../internal/merge_sort.cpp', -'../internal/radix_sort.cpp', -'../TRTCContext.cpp', -'../DVVector.cpp', -'../DVTuple.cpp', -'../fake_vectors/DVRange.cpp', -'../fake_vectors/DVConstant.cpp', -'../fake_vectors/DVCounter.cpp', -'../fake_vectors/DVDiscard.cpp', -'../fake_vectors/DVPermutation.cpp', -'../fake_vectors/DVReverse.cpp', -'../fake_vectors/DVTransform.cpp', -'../fake_vectors/DVZipped.cpp', -'../fake_vectors/DVCustomVector.cpp', -'../functor.cpp', -'../fill.cpp', -'../replace.cpp', -'../for_each.cpp', -'../adjacent_difference.cpp', -'../sequence.cpp', -'../tabulate.cpp', -'../transform.cpp', -'../gather.cpp', -'../scatter.cpp', -'../copy.cpp', -'../swap.cpp', -'../count.cpp', -'../reduce.cpp', -'../equal.cpp', -'../extrema.cpp', -'../inner_product.cpp', -'../transform_reduce.cpp', -'../logical.cpp', -'../scan.cpp', -'../transform_scan.cpp', -'../scan_by_key.cpp', -'../remove.cpp', -'../unique.cpp', -'../partition.cpp', -'../find.cpp', -'../mismatch.cpp', -'../binary_search.cpp', -'../merge.cpp', -'../sort.cpp' -] - - -module_PyThrustRTC = Extension( - 'PyThrustRTC', - sources = sources, - include_dirs = ['.', '../thirdparty/crc64', '../thirdparty/unqlite', '..', '../internal'], - define_macros = define_macros, - extra_compile_args=extra_compile_args) - setup( name = 'ThrustRTC', - version = '0.2.1', + version = '0.3.3', description = 'Thrust for Python based on NVRTC', long_description=long_description, long_description_content_type='text/markdown', @@ -101,7 +19,7 @@ author_email='hyangfeih@gmail.com', keywords='GPU CUDA Thrust', packages=['ThrustRTC'], - ext_modules=[module_PyThrustRTC], - install_requires = ['numpy'] + data_files=[("Fei", ["PyThrustRTC.dll", "libPyThrustRTC.so"])], + install_requires = ['cffi','numpy'], ) diff --git a/python/sort.hpp b/python/sort.hpp deleted file mode 100644 index 6352180..0000000 --- a/python/sort.hpp +++ /dev/null @@ -1,102 +0,0 @@ -#include -#include "TRTCContext.h" -#include "sort.h" - -static PyObject* n_sort(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - PyObject* py_comp = PyTuple_GetItem(args, 1); - Functor* comp = nullptr; - if (py_comp != Py_None) - comp = (Functor*)PyLong_AsVoidPtr(py_comp); - - if (comp == nullptr) - { - if (TRTC_Sort(*vec)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; - } - else - { - if (TRTC_Sort(*vec, *comp)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; - } -} - -static PyObject* n_sort_by_key(PyObject* self, PyObject* args) -{ - DVVectorLike* keys = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* values = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - PyObject* py_comp = PyTuple_GetItem(args, 2); - Functor* comp = nullptr; - if (py_comp != Py_None) - comp = (Functor*)PyLong_AsVoidPtr(py_comp); - - if (comp == nullptr) - { - if (TRTC_Sort_By_Key(*keys, *values)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; - } - else - { - if (TRTC_Sort_By_Key(*keys, *values, *comp)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; - } -} - - -static PyObject* n_is_sorted(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - PyObject* py_comp = PyTuple_GetItem(args, 1); - Functor* comp = nullptr; - if (py_comp != Py_None) - comp = (Functor*)PyLong_AsVoidPtr(py_comp); - bool res; - if (comp == nullptr) - { - if (TRTC_Is_Sorted(*vec, res)) - return PyBool_FromLong(res ? 1 : 0); - else - Py_RETURN_NONE; - } - else - { - if (TRTC_Is_Sorted(*vec, *comp, res)) - return PyBool_FromLong(res ? 1 : 0); - else - Py_RETURN_NONE; - } -} - -static PyObject* n_is_sorted_until(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - PyObject* py_comp = PyTuple_GetItem(args, 1); - Functor* comp = nullptr; - if (py_comp != Py_None) - comp = (Functor*)PyLong_AsVoidPtr(py_comp); - size_t res; - if (comp == nullptr) - { - if (TRTC_Is_Sorted_Until(*vec, res)) - return PyLong_FromUnsignedLongLong((unsigned long long)res); - else - Py_RETURN_NONE; - } - else - { - if (TRTC_Is_Sorted_Until(*vec, *comp, res)) - return PyLong_FromUnsignedLongLong((unsigned long long)res); - else - Py_RETURN_NONE; - } -} - diff --git a/python/swap.hpp b/python/swap.hpp deleted file mode 100644 index d4f69bc..0000000 --- a/python/swap.hpp +++ /dev/null @@ -1,13 +0,0 @@ -#include -#include "TRTCContext.h" -#include "swap.h" - -static PyObject* n_swap(PyObject* self, PyObject* args) -{ - DVVectorLike* vec1 = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec2 = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - if (TRTC_Swap(*vec1, *vec2)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; -} diff --git a/python/tabulate.hpp b/python/tabulate.hpp deleted file mode 100644 index 32bb236..0000000 --- a/python/tabulate.hpp +++ /dev/null @@ -1,13 +0,0 @@ -#include -#include "TRTCContext.h" -#include "tabulate.h" - -static PyObject* n_tabulate(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - Functor* op = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - if (TRTC_Tabulate(*vec, *op)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; -} diff --git a/python/test/CMakeLists.txt b/python/test/CMakeLists.txt index 3a7d97d..65f401f 100644 --- a/python/test/CMakeLists.txt +++ b/python/test/CMakeLists.txt @@ -1,4 +1,3 @@ - set(PYTHON_TEST test_trtc.py test_for.py @@ -29,7 +28,7 @@ set(PYTHON_TEST test_find.py test_mismatch.py test_binary_search.py - test_merge.py + test_merge.py test_sort.py test_numba.py @@ -46,6 +45,4 @@ set(PYTHON_TEST ../demo/k-means.py ) -install(FILES ${PYTHON_TEST} DESTINATION test_python ) - - +install(FILES ${PYTHON_TEST} DESTINATION test_python ) \ No newline at end of file diff --git a/python/transform.hpp b/python/transform.hpp deleted file mode 100644 index 6200793..0000000 --- a/python/transform.hpp +++ /dev/null @@ -1,64 +0,0 @@ -#include -#include "TRTCContext.h" -#include "transform.h" - -static PyObject* n_transform(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - Functor* op = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - if (TRTC_Transform(*vec_in,* vec_out, *op)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; -} - -static PyObject* n_transform_binary(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_in1 = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_in2 = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - Functor* op = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 3)); - if(TRTC_Transform_Binary(*vec_in1, *vec_in2, *vec_out, *op)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; -} - -static PyObject* n_transform_if(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - Functor* op = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - Functor* pred = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 3)); - TRTC_Transform_If(*vec_in, *vec_out, *op, *pred); - return PyLong_FromLong(0); -} - -static PyObject* n_transform_if_stencil(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_stencil = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - Functor* op = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 3)); - Functor* pred = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 4)); - if(TRTC_Transform_If_Stencil(*vec_in, *vec_stencil, *vec_out, *op, *pred)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; -} - -static PyObject* n_transform_binary_if_stencil(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_in1 = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_in2 = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DVVectorLike* vec_stencil = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 3)); - Functor* op = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 4)); - Functor* pred = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 5)); - if(TRTC_Transform_Binary_If_Stencil(*vec_in1, *vec_in2, *vec_stencil, *vec_out, *op, *pred)) - return PyLong_FromLong(0); - else - Py_RETURN_NONE; -} - diff --git a/python/transform_reduce.hpp b/python/transform_reduce.hpp deleted file mode 100644 index 1663823..0000000 --- a/python/transform_reduce.hpp +++ /dev/null @@ -1,16 +0,0 @@ -#include -#include "TRTCContext.h" -#include "transform_reduce.h" -#include "viewbuf_to_python.hpp" - -static PyObject* n_transform_reduce(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - Functor* unary_op = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DeviceViewable* init = (DeviceViewable*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - Functor* binary_op = (Functor*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 3)); - ViewBuf ret; - if (TRTC_Transform_Reduce(*vec, *unary_op, *init, *binary_op, ret)) - return PyValue_FromViewBuf(ret, init->name_view_cls().c_str()); - Py_RETURN_NONE; -} diff --git a/python/transform_scan.hpp b/python/transform_scan.hpp deleted file mode 100644 index 8fce659..0000000 --- a/python/transform_scan.hpp +++ /dev/null @@ -1,31 +0,0 @@ -#include -#include "TRTCContext.h" -#include "transform_scan.h" - -static PyObject* n_transform_inclusive_scan(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - PyObject* py_unary_op = PyTuple_GetItem(args, 2); - Functor* unary_op = (Functor*)PyLong_AsVoidPtr(py_unary_op); - PyObject* py_binary_op = PyTuple_GetItem(args, 3); - Functor* binary_op = (Functor*)PyLong_AsVoidPtr(py_binary_op); - if (TRTC_Transform_Inclusive_Scan(*vec_in, *vec_out, *unary_op, *binary_op)) - return PyLong_FromLong(0); - Py_RETURN_NONE; -} - -static PyObject* n_transform_exclusive_scan(PyObject* self, PyObject* args) -{ - DVVectorLike* vec_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - PyObject* py_unary_op = PyTuple_GetItem(args, 2); - Functor* unary_op = (Functor*)PyLong_AsVoidPtr(py_unary_op); - PyObject* pyinit = PyTuple_GetItem(args, 3); - DeviceViewable* init = (DeviceViewable*)PyLong_AsVoidPtr(pyinit); - PyObject* py_binary_op = PyTuple_GetItem(args, 4); - Functor* binary_op = (Functor*)PyLong_AsVoidPtr(py_binary_op); - if (TRTC_Transform_Exclusive_Scan(*vec_in, *vec_out, *unary_op, *init, *binary_op)) - return PyLong_FromLong(0); - Py_RETURN_NONE; -} diff --git a/python/unique.hpp b/python/unique.hpp deleted file mode 100644 index 066ddc8..0000000 --- a/python/unique.hpp +++ /dev/null @@ -1,110 +0,0 @@ -#include -#include "TRTCContext.h" -#include "unique.h" - -static PyObject* n_unique(PyObject* self, PyObject* args) -{ - DVVectorLike* vec = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - Functor* binary_pred = nullptr; - PyObject* py_binary_pred = PyTuple_GetItem(args, 1); - if (py_binary_pred!= Py_None) - binary_pred=(Functor*)PyLong_AsVoidPtr(py_binary_pred); - if (binary_pred == nullptr) - { - uint32_t res = TRTC_Unique(*vec); - if (res != uint32_t(-1)) - return PyLong_FromUnsignedLong((unsigned long)res); - else - Py_RETURN_NONE; - } - else - { - uint32_t res = TRTC_Unique(*vec, *binary_pred); - if (res != uint32_t(-1)) - return PyLong_FromUnsignedLong((unsigned long)res); - else - Py_RETURN_NONE; - } -} - -static PyObject* n_unique_copy(PyObject* self, PyObject* args) -{ - - DVVectorLike* vec_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* vec_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - Functor* binary_pred = nullptr; - PyObject* py_binary_pred = PyTuple_GetItem(args, 2); - if (py_binary_pred != Py_None) - binary_pred = (Functor*)PyLong_AsVoidPtr(py_binary_pred); - if (binary_pred == nullptr) - { - uint32_t res = TRTC_Unique_Copy(*vec_in, *vec_out); - if (res != uint32_t(-1)) - return PyLong_FromUnsignedLong((unsigned long)res); - else - Py_RETURN_NONE; - } - else - { - uint32_t res = TRTC_Unique_Copy(*vec_in, *vec_out, *binary_pred); - if (res != uint32_t(-1)) - return PyLong_FromUnsignedLong((unsigned long)res); - else - Py_RETURN_NONE; - } -} - -static PyObject* n_unique_by_key(PyObject* self, PyObject* args) -{ - DVVectorLike* keys = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* values = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - Functor* binary_pred = nullptr; - PyObject* py_binary_pred = PyTuple_GetItem(args, 2); - if (py_binary_pred != Py_None) - binary_pred = (Functor*)PyLong_AsVoidPtr(py_binary_pred); - if (binary_pred == nullptr) - { - uint32_t res = TRTC_Unique_By_Key(*keys, *values); - if (res != uint32_t(-1)) - return PyLong_FromUnsignedLong((unsigned long)res); - else - Py_RETURN_NONE; - } - else - { - uint32_t res = TRTC_Unique_By_Key(*keys, *values, *binary_pred); - if (res != uint32_t(-1)) - return PyLong_FromUnsignedLong((unsigned long)res); - else - Py_RETURN_NONE; - } -} - -static PyObject* n_unique_by_key_copy(PyObject* self, PyObject* args) -{ - DVVectorLike* keys_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 0)); - DVVectorLike* values_in = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 1)); - DVVectorLike* keys_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 2)); - DVVectorLike* values_out = (DVVectorLike*)PyLong_AsVoidPtr(PyTuple_GetItem(args, 3)); - Functor* binary_pred = nullptr; - PyObject* py_binary_pred = PyTuple_GetItem(args, 4); - if (py_binary_pred != Py_None) - binary_pred = (Functor*)PyLong_AsVoidPtr(py_binary_pred); - if (binary_pred == nullptr) - { - uint32_t res = TRTC_Unique_By_Key_Copy(*keys_in, *values_in, *keys_out, *values_out); - if (res != uint32_t(-1)) - return PyLong_FromUnsignedLong((unsigned long)res); - else - Py_RETURN_NONE; - } - else - { - uint32_t res = TRTC_Unique_By_Key_Copy(*keys_in, *values_in, *keys_out, *values_out, *binary_pred); - if (res != uint32_t(-1)) - return PyLong_FromUnsignedLong((unsigned long)res); - else - Py_RETURN_NONE; - } -} - diff --git a/python/viewbuf_to_python.hpp b/python/viewbuf_to_python.hpp deleted file mode 100644 index 682c37b..0000000 --- a/python/viewbuf_to_python.hpp +++ /dev/null @@ -1,38 +0,0 @@ -#ifndef _viewbuf_to_python_hpp_ -#define _viewbuf_to_python_hpp_ - -#include -#include "DeviceViewable.h" - -PyObject* PyValue_FromViewBuf(const ViewBuf& buf, const char* type) -{ - std::string s_type = type; - if (s_type == "int8_t") - return PyLong_FromLong((long)(*(int8_t*)buf.data())); - else if (s_type == "uint8_t") - return PyLong_FromUnsignedLong((unsigned long)(*(uint8_t*)buf.data())); - else if (s_type == "int16_t") - return PyLong_FromLong((long)(*(int16_t*)buf.data())); - else if (s_type == "uint16_t") - return PyLong_FromUnsignedLong((unsigned long)(*(uint16_t*)buf.data())); - else if (s_type == "int32_t") - return PyLong_FromLong((long)(*(int32_t*)buf.data())); - else if (s_type == "uint32_t") - return PyLong_FromUnsignedLong((unsigned long)(*(uint32_t*)buf.data())); - else if (s_type == "int64_t") - return PyLong_FromLongLong((long long)(*(int64_t*)buf.data())); - else if (s_type == "uint64_t") - return PyLong_FromUnsignedLongLong((unsigned long long)(*(uint64_t*)buf.data())); - else if (s_type == "float") - return PyFloat_FromDouble((double)(*(float*)buf.data())); - else if (s_type == "double") - return PyFloat_FromDouble(*(double*)buf.data()); - else if (s_type == "bool") - return PyBool_FromLong(*(bool*)buf.data()?1:0); - - char str[64]; - sprintf(str, "[Device-viewable object, type: %s]", type); - return PyUnicode_FromString(str); -} - -#endif