[arrayfire] 02/75: FEAT: Adding functions exposing Array internals
Ghislain Vaillant
ghisvail-guest at moszumanska.debian.org
Mon Feb 29 08:01:06 UTC 2016
This is an automated email from the git hooks/post-receive script.
ghisvail-guest pushed a commit to branch dfsg-clean
in repository arrayfire.
commit e0879cb37f9223d68004f65b984183eb83e8ddc7
Author: Pavan Yalamanchili <pavan at arrayfire.com>
Date: Fri Feb 5 01:44:16 2016 -0500
FEAT: Adding functions exposing Array internals
- af_create_array_with_strides
- af_get_strides
- af_get_offset
- af_get_raw_ptr
- af_is_linear
- af_is_owner
---
include/af/internal.h | 62 +++++++++++++++
src/api/c/internal.cpp | 181 +++++++++++++++++++++++++++++++++++++++++++
src/api/cpp/internal.cpp | 63 +++++++++++++++
src/api/unified/internal.cpp | 54 +++++++++++++
src/backend/cpu/Array.cpp | 18 +++++
src/backend/cpu/Array.hpp | 5 ++
src/backend/cuda/Array.cpp | 21 +++++
src/backend/cuda/Array.hpp | 4 +
src/backend/opencl/Array.cpp | 21 +++++
src/backend/opencl/Array.hpp | 4 +
10 files changed, 433 insertions(+)
diff --git a/include/af/internal.h b/include/af/internal.h
new file mode 100644
index 0000000..fdd0158
--- /dev/null
+++ b/include/af/internal.h
@@ -0,0 +1,62 @@
+/*******************************************************
+ * Copyright (c) 2016, ArrayFire
+ * All rights reserved.
+ *
+ * This file is distributed under 3-clause BSD license.
+ * The complete license agreement can be obtained at:
+ * http://arrayfire.com/licenses/BSD-3-Clause
+ ********************************************************/
+
+#pragma once
+#include <af/defines.h>
+#include <af/dim4.hpp>
+
+#ifdef __cplusplus
+namespace af
+{
+ class array;
+
+ AFAPI array createArray(const void *data, const dim_t offset,
+ const dim4 dims, const dim4 strides,
+ const af::dtype ty,
+ const af::source location);
+
+ AFAPI dim4 getStrides(const array &in);
+
+ AFAPI dim_t getOffset(const array &in);
+
+ AFAPI void *getRawPtr(const array &in);
+
+ AFAPI bool isLinear(const array &in);
+
+ AFAPI bool isOwner(const array &in);
+}
+#endif
+
+#ifdef __cplusplus
+extern "C"
+{
+#endif
+
+ AFAPI af_err af_create_array_with_strides(af_array *arr,
+ const void *data,
+ const dim_t offset,
+ const unsigned ndims,
+ const dim_t *const dims,
+ const dim_t *const strides,
+ const af_dtype ty,
+ const af_source location);
+
+ AFAPI af_err af_get_strides(dim_t *s0, dim_t *s1, dim_t *s2, dim_t *s3, const af_array arr);
+
+ AFAPI af_err af_get_offset(dim_t *offset, const af_array arr);
+
+ AFAPI af_err af_get_raw_ptr(void **ptr, const af_array arr);
+
+ AFAPI af_err af_is_linear(bool *result, const af_array arr);
+
+ AFAPI af_err af_is_owner(bool *result, const af_array arr);
+
+#ifdef __cplusplus
+}
+#endif
diff --git a/src/api/c/internal.cpp b/src/api/c/internal.cpp
new file mode 100644
index 0000000..d086d43
--- /dev/null
+++ b/src/api/c/internal.cpp
@@ -0,0 +1,181 @@
+/*******************************************************
+ * Copyright (c) 2016, ArrayFire
+ * All rights reserved.
+ *
+ * This file is distributed under 3-clause BSD license.
+ * The complete license agreement can be obtained at:
+ * http://arrayfire.com/licenses/BSD-3-Clause
+ ********************************************************/
+
+#include <af/dim4.hpp>
+#include <af/device.h>
+#include <af/version.h>
+#include <af/internal.h>
+#include <backend.hpp>
+#include <platform.hpp>
+#include <Array.hpp>
+#include <handle.hpp>
+#include "err_common.hpp"
+#include <cstring>
+
+using namespace detail;
+
+af_err af_create_array_with_strides(af_array *arr,
+ const void *data,
+ const dim_t offset,
+ const unsigned ndims,
+ const dim_t *const dims_,
+ const dim_t *const strides_,
+ const af_dtype ty,
+ const af_source location)
+{
+ try {
+
+ ARG_ASSERT(2, offset >= 0);
+ ARG_ASSERT(3, ndims >=1 && ndims <= 4);
+ ARG_ASSERT(4, dims_ != NULL);
+ ARG_ASSERT(5, strides_ != NULL);
+ ARG_ASSERT(5, strides_[0] == 1);
+
+ dim4 dims(ndims, dims_);
+ dim4 strides(ndims, strides_);
+
+ bool isdev = location == afDevice;
+
+ af_array res;
+ AF_CHECK(af_init());
+
+ switch (ty) {
+ case f32: res = getHandle(Array<float >(dims, strides, offset, (float *)data, isdev)); break;
+ case f64: res = getHandle(Array<double >(dims, strides, offset, (double *)data, isdev)); break;
+ case c32: res = getHandle(Array<cfloat >(dims, strides, offset, (cfloat *)data, isdev)); break;
+ case c64: res = getHandle(Array<cdouble>(dims, strides, offset, (cdouble *)data, isdev)); break;
+ case u32: res = getHandle(Array<uint >(dims, strides, offset, (uint *)data, isdev)); break;
+ case s32: res = getHandle(Array<int >(dims, strides, offset, (int *)data, isdev)); break;
+ case u64: res = getHandle(Array<uintl >(dims, strides, offset, (uintl *)data, isdev)); break;
+ case s64: res = getHandle(Array<intl >(dims, strides, offset, (intl *)data, isdev)); break;
+ case u16: res = getHandle(Array<ushort >(dims, strides, offset, (ushort *)data, isdev)); break;
+ case s16: res = getHandle(Array<short >(dims, strides, offset, (short *)data, isdev)); break;
+ case b8 : res = getHandle(Array<char >(dims, strides, offset, (char *)data, isdev)); break;
+ case u8 : res = getHandle(Array<uchar >(dims, strides, offset, (uchar *)data, isdev)); break;
+ default: TYPE_ERROR(6, ty);
+ }
+
+ std::swap(*arr, res);
+ }
+ CATCHALL;
+ return AF_SUCCESS;
+}
+
+af_err af_get_strides(dim_t *s0, dim_t *s1, dim_t *s2, dim_t *s3, const af_array in)
+{
+ try {
+ ArrayInfo info = getInfo(in);
+ *s0 = info.strides()[0];
+ *s1 = info.strides()[1];
+ *s2 = info.strides()[2];
+ *s3 = info.strides()[3];
+ }
+ CATCHALL
+ return AF_SUCCESS;
+}
+
+af_err af_get_offset(dim_t *offset, const af_array arr)
+{
+ try {
+
+ dim_t res = 0;
+
+ af_dtype ty = getInfo(arr).getType();
+
+ switch (ty) {
+ case f32: res = getArray<float >(arr).getOffset(); break;
+ case f64: res = getArray<double >(arr).getOffset(); break;
+ case c32: res = getArray<cfloat >(arr).getOffset(); break;
+ case c64: res = getArray<cdouble>(arr).getOffset(); break;
+ case u32: res = getArray<uint >(arr).getOffset(); break;
+ case s32: res = getArray<int >(arr).getOffset(); break;
+ case u64: res = getArray<uintl >(arr).getOffset(); break;
+ case s64: res = getArray<intl >(arr).getOffset(); break;
+ case u16: res = getArray<ushort >(arr).getOffset(); break;
+ case s16: res = getArray<short >(arr).getOffset(); break;
+ case b8 : res = getArray<char >(arr).getOffset(); break;
+ case u8 : res = getArray<uchar >(arr).getOffset(); break;
+ default: TYPE_ERROR(6, ty);
+ }
+
+ std::swap(*offset, res);
+ }
+ CATCHALL;
+ return AF_SUCCESS;
+
+}
+
+af_err af_get_raw_ptr(void **ptr, const af_array arr)
+{
+ try {
+
+ void *res = NULL;
+
+ af_dtype ty = getInfo(arr).getType();
+
+ switch (ty) {
+ case f32: res = (void *)getArray<float >(arr).get(); break;
+ case f64: res = (void *)getArray<double >(arr).get(); break;
+ case c32: res = (void *)getArray<cfloat >(arr).get(); break;
+ case c64: res = (void *)getArray<cdouble>(arr).get(); break;
+ case u32: res = (void *)getArray<uint >(arr).get(); break;
+ case s32: res = (void *)getArray<int >(arr).get(); break;
+ case u64: res = (void *)getArray<uintl >(arr).get(); break;
+ case s64: res = (void *)getArray<intl >(arr).get(); break;
+ case u16: res = (void *)getArray<ushort >(arr).get(); break;
+ case s16: res = (void *)getArray<short >(arr).get(); break;
+ case b8 : res = (void *)getArray<char >(arr).get(); break;
+ case u8 : res = (void *)getArray<uchar >(arr).get(); break;
+ default: TYPE_ERROR(6, ty);
+ }
+
+ std::swap(*ptr, res);
+ }
+ CATCHALL;
+ return AF_SUCCESS;
+}
+
+af_err af_is_linear(bool *result, const af_array arr)
+{
+ try {
+ *result = getInfo(arr).isLinear();
+ }
+ CATCHALL
+ return AF_SUCCESS;
+}
+
+af_err af_is_owner(bool *result, const af_array arr)
+{
+ try {
+
+ bool res = false;
+
+ af_dtype ty = getInfo(arr).getType();
+
+ switch (ty) {
+ case f32: res = (void *)getArray<float >(arr).isOwner(); break;
+ case f64: res = (void *)getArray<double >(arr).isOwner(); break;
+ case c32: res = (void *)getArray<cfloat >(arr).isOwner(); break;
+ case c64: res = (void *)getArray<cdouble>(arr).isOwner(); break;
+ case u32: res = (void *)getArray<uint >(arr).isOwner(); break;
+ case s32: res = (void *)getArray<int >(arr).isOwner(); break;
+ case u64: res = (void *)getArray<uintl >(arr).isOwner(); break;
+ case s64: res = (void *)getArray<intl >(arr).isOwner(); break;
+ case u16: res = (void *)getArray<ushort >(arr).isOwner(); break;
+ case s16: res = (void *)getArray<short >(arr).isOwner(); break;
+ case b8 : res = (void *)getArray<char >(arr).isOwner(); break;
+ case u8 : res = (void *)getArray<uchar >(arr).isOwner(); break;
+ default: TYPE_ERROR(6, ty);
+ }
+
+ std::swap(*result, res);
+ }
+ CATCHALL;
+ return AF_SUCCESS;
+}
diff --git a/src/api/cpp/internal.cpp b/src/api/cpp/internal.cpp
new file mode 100644
index 0000000..f26f9f8
--- /dev/null
+++ b/src/api/cpp/internal.cpp
@@ -0,0 +1,63 @@
+/*******************************************************
+ * Copyright (c) 2016, ArrayFire
+ * All rights reserved.
+ *
+ * This file is distributed under 3-clause BSD license.
+ * The complete license agreement can be obtained at:
+ * http://arrayfire.com/licenses/BSD-3-Clause
+ ********************************************************/
+
+#include <af/internal.h>
+#include <af/array.h>
+#include "error.hpp"
+
+namespace af
+{
+ array createArray(const void *data, const dim_t offset,
+ const dim4 dims, const dim4 strides,
+ const af::dtype ty,
+ const af::source location)
+ {
+ af_array res;
+ AF_THROW(af_create_array_with_strides(&res, data, offset,
+ dims.ndims(), dims.get(), strides.get(),
+ ty, location));
+ return array(res);
+ }
+
+ dim4 getStrides(const array &in)
+ {
+ dim_t s0, s1, s2, s3;
+ AF_THROW(af_get_strides(&s0, &s1, &s2, &s3, in.get()));
+ return dim4(s0, s1, s2, s3);
+ }
+
+ dim_t getOffset(const array &in)
+ {
+ dim_t offset;
+ AF_THROW(af_get_offset(&offset, in.get()));
+ return offset;
+ }
+
+ void *getRawPtr(const array &in)
+ {
+ void *ptr = NULL;
+ AF_THROW(af_get_raw_ptr(&ptr, in.get()));
+ return ptr;
+ }
+
+ bool isLinear(const array &in)
+ {
+ bool is_linear = false;
+ AF_THROW(af_is_linear(&is_linear, in.get()));
+ return is_linear;
+ }
+
+ bool isOwner(const array &in)
+ {
+ bool is_owner = false;
+ AF_THROW(af_is_owner(&is_owner, in.get()));
+ return is_owner;
+ }
+
+}
diff --git a/src/api/unified/internal.cpp b/src/api/unified/internal.cpp
new file mode 100644
index 0000000..7c223e7
--- /dev/null
+++ b/src/api/unified/internal.cpp
@@ -0,0 +1,54 @@
+/*******************************************************
+ * Copyright (c) 2016, ArrayFire
+ * All rights reserved.
+ *
+ * This file is distributed under 3-clause BSD license.
+ * The complete license agreement can be obtained at:
+ * http://arrayfire.com/licenses/BSD-3-Clause
+ ********************************************************/
+
+#include <af/internal.h>
+#include "symbol_manager.hpp"
+
+
+af_err af_create_array_with_strides(af_array *arr,
+ const void *data,
+ const dim_t offset,
+ const unsigned ndims,
+ const dim_t *const dims_,
+ const dim_t *const strides_,
+ const af_dtype ty,
+ const af_source location)
+{
+ return CALL(arr, data, offset, ndims, dims_, strides_, ty, location);
+}
+
+af_err af_get_strides(dim_t *s0, dim_t *s1, dim_t *s2, dim_t *s3, const af_array in)
+{
+ CHECK_ARRAYS(in);
+ return CALL(s0, s1, s2, s3, in);
+}
+
+af_err af_get_offset(dim_t *offset, const af_array arr)
+{
+ CHECK_ARRAYS(arr);
+ return CALL(offset, arr);
+}
+
+af_err af_get_raw_ptr(void **ptr, const af_array arr)
+{
+ CHECK_ARRAYS(arr);
+ return CALL(ptr, arr);
+}
+
+af_err af_is_linear(bool *result, const af_array arr)
+{
+ CHECK_ARRAYS(arr);
+ return CALL(result, arr);
+}
+
+af_err af_is_owner(bool *result, const af_array arr)
+{
+ CHECK_ARRAYS(arr);
+ return CALL(result, arr);
+}
diff --git a/src/backend/cpu/Array.cpp b/src/backend/cpu/Array.cpp
index 6d51f63..fbb8e10 100644
--- a/src/backend/cpu/Array.cpp
+++ b/src/backend/cpu/Array.cpp
@@ -68,6 +68,21 @@ Array<T>::Array(const Array<T>& parent, const dim4 &dims, const dim4 &offsets, c
ready(true), owner(false)
{ }
+template<typename T>
+Array<T>::Array(af::dim4 dims, af::dim4 strides, dim_t offset_,
+ const T * const in_data, bool is_device) :
+ info(getActiveDeviceId(), dims, af::dim4(offset_), strides, (af_dtype)dtype_traits<T>::af_type),
+ data(is_device ? (T*)in_data : memAlloc<T>(info.elements()), memFree<T>),
+ data_dims(dims),
+ node(),
+ offset(offset_),
+ ready(true),
+ owner(true)
+{
+ if (!is_device) {
+ std::copy(in_data, in_data + dims.elements(), data.get());
+ }
+}
template<typename T>
void Array<T>::eval()
@@ -240,6 +255,9 @@ writeDeviceDataArray(Array<T> &arr, const void * const data, const size_t bytes)
template void Array<T>::eval() const; \
template Array<T>::Array(af::dim4 dims, const T * const in_data, \
bool is_device, bool copy_device); \
+ template Array<T>::Array(af::dim4 dims, af::dim4 strides, dim_t offset, \
+ const T * const in_data, \
+ bool is_device); \
template TNJ::Node_ptr Array<T>::getNode() const; \
template void writeHostDataArray<T> (Array<T> &arr, const T * const data, const size_t bytes); \
template void writeDeviceDataArray<T> (Array<T> &arr, const void * const data, const size_t bytes); \
diff --git a/src/backend/cpu/Array.hpp b/src/backend/cpu/Array.hpp
index 9cd154e..891d867 100644
--- a/src/backend/cpu/Array.hpp
+++ b/src/backend/cpu/Array.hpp
@@ -106,12 +106,17 @@ namespace cpu
Array() = default;
Array(dim4 dims);
+
explicit Array(dim4 dims, const T * const in_data, bool is_device, bool copy_device=false);
Array(const Array<T>& parnt, const dim4 &dims, const dim4 &offset, const dim4 &stride);
explicit Array(af::dim4 dims, TNJ::Node_ptr n);
public:
+
+ Array(af::dim4 dims, af::dim4 strides, dim_t offset,
+ const T * const in_data, bool is_device = false);
+
void resetInfo(const af::dim4& dims) { info.resetInfo(dims); }
void resetDims(const af::dim4& dims) { info.resetDims(dims); }
void modDims(const af::dim4 &newDims) { info.modDims(newDims); }
diff --git a/src/backend/cuda/Array.cpp b/src/backend/cuda/Array.cpp
index 48bee65..366d8e2 100644
--- a/src/backend/cuda/Array.cpp
+++ b/src/backend/cuda/Array.cpp
@@ -86,6 +86,24 @@ namespace cuda
{
}
+ template<typename T>
+ Array<T>::Array(af::dim4 dims, af::dim4 strides, dim_t offset_,
+ const T * const in_data, bool is_device) :
+ info(getActiveDeviceId(), dims, af::dim4(offset_), strides, (af_dtype)dtype_traits<T>::af_type),
+ data(is_device ? (T*)in_data : memAlloc<T>(info.elements()), memFree<T>),
+ data_dims(dims),
+ node(),
+ offset(offset_),
+ ready(true),
+ owner(true)
+ {
+ if (!is_device) {
+ cudaStream_t stream = getStream(getActiveDeviceId());
+ CUDA_CHECK(cudaMemcpyAsync(data.get(), in_data, info.elements() * sizeof(T),
+ cudaMemcpyHostToDevice, stream));
+ CUDA_CHECK(cudaStreamSynchronize(stream));
+ }
+ }
template<typename T>
void Array<T>::eval()
@@ -275,6 +293,9 @@ namespace cuda
bool copy); \
template void destroyArray<T> (Array<T> *A); \
template Array<T> createNodeArray<T> (const dim4 &size, JIT::Node_ptr node); \
+ template Array<T>::Array(af::dim4 dims, af::dim4 strides, dim_t offset, \
+ const T * const in_data, \
+ bool is_device); \
template Array<T>::Array(af::dim4 dims, const T * const in_data, \
bool is_device, bool copy_device); \
template Array<T>::~Array (); \
diff --git a/src/backend/cuda/Array.hpp b/src/backend/cuda/Array.hpp
index ad4396b..b8832db 100644
--- a/src/backend/cuda/Array.hpp
+++ b/src/backend/cuda/Array.hpp
@@ -103,12 +103,16 @@ namespace cuda
bool owner;
Array(af::dim4 dims);
+
explicit Array(af::dim4 dims, const T * const in_data, bool is_device = false, bool copy_device = false);
Array(const Array<T>& parnt, const dim4 &dims, const dim4 &offset, const dim4 &stride);
Array(Param<T> &tmp);
Array(af::dim4 dims, JIT::Node_ptr n);
public:
+ Array(af::dim4 dims, af::dim4 strides, dim_t offset,
+ const T * const in_data, bool is_device = false);
+
void resetInfo(const af::dim4& dims) { info.resetInfo(dims); }
void resetDims(const af::dim4& dims) { info.resetDims(dims); }
void modDims(const af::dim4 &newDims) { info.modDims(newDims); }
diff --git a/src/backend/opencl/Array.cpp b/src/backend/opencl/Array.cpp
index 178be5b..f41b2c7 100644
--- a/src/backend/opencl/Array.cpp
+++ b/src/backend/opencl/Array.cpp
@@ -99,6 +99,24 @@ namespace opencl
{
}
+ template<typename T>
+ Array<T>::Array(af::dim4 dims, af::dim4 strides, dim_t offset_,
+ const T * const in_data, bool is_device) :
+ info(getActiveDeviceId(), dims, af::dim4(offset_), strides, (af_dtype)dtype_traits<T>::af_type),
+ data(is_device ?
+ (new cl::Buffer((cl_mem)in_data)) :
+ (bufferAlloc(info.elements() * sizeof(T))), bufferFree),
+ data_dims(dims),
+ node(),
+ offset(offset_),
+ ready(true),
+ owner(true)
+ {
+ if (!is_device) {
+ getQueue().enqueueWriteBuffer(*data.get(), CL_TRUE, 0, sizeof(T) * info.elements(), in_data);
+ }
+ }
+
template<typename T>
void Array<T>::eval()
@@ -308,6 +326,9 @@ namespace opencl
bool copy); \
template void destroyArray<T> (Array<T> *A); \
template Array<T> createNodeArray<T> (const dim4 &size, JIT::Node_ptr node); \
+ template Array<T>::Array(af::dim4 dims, af::dim4 strides, dim_t offset, \
+ const T * const in_data, \
+ bool is_device); \
template Array<T>::Array(af::dim4 dims, cl_mem mem, size_t src_offset, bool copy); \
template Array<T>::~Array (); \
template Node_ptr Array<T>::getNode() const; \
diff --git a/src/backend/opencl/Array.hpp b/src/backend/opencl/Array.hpp
index 4c8c05a..d1a4d97 100644
--- a/src/backend/opencl/Array.hpp
+++ b/src/backend/opencl/Array.hpp
@@ -95,6 +95,7 @@ namespace opencl
bool owner;
Array(af::dim4 dims);
+
Array(const Array<T>& parnt, const dim4 &dims, const dim4 &offset, const dim4 &stride);
Array(Param &tmp);
explicit Array(af::dim4 dims, JIT::Node_ptr n);
@@ -103,6 +104,9 @@ namespace opencl
public:
+ Array(af::dim4 dims, af::dim4 strides, dim_t offset,
+ const T * const in_data, bool is_device = false);
+
void resetInfo(const af::dim4& dims) { info.resetInfo(dims); }
void resetDims(const af::dim4& dims) { info.resetDims(dims); }
void modDims(const af::dim4 &newDims) { info.modDims(newDims); }
--
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/debian-science/packages/arrayfire.git
More information about the debian-science-commits
mailing list