Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 2 additions & 1 deletion examples/machine_learning/kmeans.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -61,7 +61,8 @@ array new_means(array data, array clusters, int k) {
void kmeans(array &means, array &clusters, const array in, int k,
int iter = 100) {
unsigned n = in.dims(0); // Num of data points
unsigned d = in.dims(2); // Num of features (will only be 1 in spider image example)
unsigned d =
in.dims(2); // Num of features (will only be 1 in spider image example)

// reshape input
array data = in * 0;
Expand Down
2 changes: 1 addition & 1 deletion examples/pde/bhrt.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -830,7 +830,7 @@ struct AccretionDisk : public Object {

af::array get_color(const af::array& ray_begin,
const af::array& ray_end) const override {
auto pair = intersect(ray_begin, ray_end);
auto pair = intersect(ray_begin, ray_end);
af::array hit = pair.first;
af::array pos = pair.second;

Expand Down
1 change: 0 additions & 1 deletion src/api/c/confidence_connected.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -45,7 +45,6 @@ using std::swap;
template<typename T>
Array<T> pointList(const Array<T>& in, const Array<uint>& x,
const Array<uint>& y) {

// TODO: Temporary Fix, must fix handling subarrays upstream
// Array<T> has to be a basic array, to be accepted as af_index
Array<uint> x_ = (x.getOffset() == 0 && x.isLinear()) ? x : copyArray(x);
Expand Down
12 changes: 4 additions & 8 deletions src/api/c/hist.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -69,8 +69,8 @@ fg_chart setup_histogram(fg_window const window, const af_array in,
T freqMax =
getScalar<T>(detail::reduce_all<af_max_t, T, T>(histogramInput));

// For histogram, xMin and xMax should always be the first
// and last bin respectively and should not be rounded
// For histogram, xMin and xMax should always be the first
// and last bin respectively and should not be rounded
if (xMin == 0 && xMax == 0 && yMin == 0 && yMax == 0) {
// No previous limits. Set without checking
xMin = static_cast<float>(minval);
Expand All @@ -79,12 +79,8 @@ fg_chart setup_histogram(fg_window const window, const af_array in,
// For histogram, always set yMin to 0.
yMin = 0;
} else {
if (xMin > minval) {
xMin = static_cast<float>(minval);
}
if (xMax < maxval) {
xMax = static_cast<float>(maxval);
}
if (xMin > minval) { xMin = static_cast<float>(minval); }
if (xMax < maxval) { xMax = static_cast<float>(maxval); }
if (yMax < freqMax) {
yMax = static_cast<float>(step_round(freqMax, true));
}
Expand Down
10 changes: 6 additions & 4 deletions src/api/c/index.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -178,11 +178,12 @@ af_err af_lookup(af_array* out, const af_array in, const af_array indices,
ARG_ASSERT(2, (idxType != b8));

af_array output = 0;
af_array idx = 0;
af_array idx = 0;

if (!idxInfo.isColumn()) {
// Force a deep copy to flatten the array and handle subarrays of not column vector arrays correctly
AF_CHECK(af_copy_array(&idx, indices));
// Force a deep copy to flatten the array and handle subarrays of
// not column vector arrays correctly
AF_CHECK(af_copy_array(&idx, indices));
} else {
idx = indices;
}
Expand All @@ -204,7 +205,8 @@ af_err af_lookup(af_array* out, const af_array in, const af_array indices,
std::swap(*out, output);

if (idx != indices) {
AF_CHECK(af_release_array(idx)); // Release indices array if a copy has been made
AF_CHECK(af_release_array(
idx)); // Release indices array if a copy has been made
}
}
CATCHALL;
Expand Down
8 changes: 5 additions & 3 deletions src/api/c/norm.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -39,8 +39,8 @@ using detail::scalar;

template<typename T>
using normReductionResult =
typename std::conditional<std::is_same<T, arrayfire::common::half>::value, float,
T>::type;
typename std::conditional<std::is_same<T, arrayfire::common::half>::value,
float, T>::type;

template<typename T>
double matrixNorm(const Array<T> &A, double p) {
Expand Down Expand Up @@ -141,7 +141,9 @@ af_err af_norm(double *out, const af_array in, const af_norm_type type,
case f64: *out = norm<double>(in, type, p, q); break;
case c32: *out = norm<cfloat>(in, type, p, q); break;
case c64: *out = norm<cdouble>(in, type, p, q); break;
case f16: *out = norm<arrayfire::common::half>(in, type, p, q); break;
case f16:
*out = norm<arrayfire::common::half>(in, type, p, q);
break;
default: TYPE_ERROR(1, i_type);
}
}
Expand Down
2 changes: 2 additions & 0 deletions src/backend/common/cast.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -30,6 +30,7 @@ namespace common {
/// Cast operations that will be removed
/// outer -> inner -> outer
///
// clang-format off
/// inner cast
/// f32 f64 c32 c64 s32 u32 s8 u8 b8 s64 u64 s16 u16 f16
/// f32 x x x x x
Expand All @@ -46,6 +47,7 @@ namespace common {
/// t s16 x x x x x x x x x x x
/// u16 x x x x x x x x x x x
/// f16 x x x x x
// clang-format on
///
/// \param[in] outer The type of the second cast and the child of the
/// previous cast
Expand Down
3 changes: 1 addition & 2 deletions src/backend/common/err_common.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -93,8 +93,7 @@ int ArgumentError::getArgIndex() const noexcept { return argIndex; }
SupportError::SupportError(const char *const func, const char *const file,
const int line, const char *const back,
const char *const message, stacktrace st)
: AfError(func, file, line, message, AF_ERR_NOT_SUPPORTED,
std::move(st))
: AfError(func, file, line, message, AF_ERR_NOT_SUPPORTED, std::move(st))
, backend(back) {}

const string &SupportError::getBackendName() const noexcept { return backend; }
Expand Down
4 changes: 1 addition & 3 deletions src/backend/common/half.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -881,9 +881,7 @@ AF_CONSTEXPR T half2int(native_half_t value) {
else AF_IF_CONSTEXPR(std::is_same<T, int>::value) {
return __half2int_rn(value);
}
else {
return __half2uint_rn(value);
}
else { return __half2uint_rn(value); }
#elif defined(AF_ONEAPI)
return static_cast<T>(value);
#else
Expand Down
6 changes: 3 additions & 3 deletions src/backend/common/jit/BufferNodeBase.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -122,12 +122,12 @@ class BufferNodeBase : public common::Node {

virtual void modDims(const af::dim4 &newDim) override {
af::dim4 strides(1, 1, 1, 1);
for(dim_t i = 1; i < 4; ++i) {
for (dim_t i = 1; i < 4; ++i) {
strides[i] = strides[i - 1] * newDim[i - 1];
}

for(dim_t i = 0; i < 4; ++i) {
m_param.dims[i] = newDim[i];
for (dim_t i = 0; i < 4; ++i) {
m_param.dims[i] = newDim[i];
m_param.strides[i] = strides[i];
}
}
Expand Down
4 changes: 1 addition & 3 deletions src/backend/common/jit/Node.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -312,9 +312,7 @@ class Node {
}
virtual std::unique_ptr<Node> clone() = 0;

virtual void modDims(const af::dim4 &newDim) {
UNUSED(newDim);
}
virtual void modDims(const af::dim4 &newDim) { UNUSED(newDim); }

#ifdef AF_CPU
template<typename U>
Expand Down
15 changes: 7 additions & 8 deletions src/backend/common/moddims.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -18,30 +18,28 @@ using detail::Array;
using detail::copyArray;
using detail::createNodeArray;

using std::make_shared;
using std::shared_ptr;
using std::array;
using arrayfire::common::Node;
using arrayfire::common::Node_ptr;
using std::array;
using std::make_shared;
using std::shared_ptr;
using std::vector;

namespace arrayfire {
namespace common {

Node_ptr copyModdims(const Node_ptr &in, const af::dim4 &newDim) {

Node_ptr out = in->clone();
for(int i = 0; i < in->kMaxChildren && in->m_children[i] != nullptr; ++i) {
for (int i = 0; i < in->kMaxChildren && in->m_children[i] != nullptr; ++i) {
out->m_children[i] = copyModdims(in->m_children[i], newDim);
}
if(out->isBuffer()) out->modDims(newDim);
if (out->isBuffer()) out->modDims(newDim);

return out;
}

template<typename T>
Array<T> moddimOp(const Array<T> &in, af::dim4 outDim) {

const auto &node = in.getNode();

NodeIterator<> it(node.get());
Expand All @@ -55,7 +53,8 @@ Array<T> moddimOp(const Array<T> &in, af::dim4 outDim) {
}
if (all_linear == false) in.eval();

Array<T> out = createNodeArray<T>(outDim, copyModdims(in.getNode(), outDim));
Array<T> out =
createNodeArray<T>(outDim, copyModdims(in.getNode(), outDim));

return out;
}
Expand Down
2 changes: 1 addition & 1 deletion src/backend/cpu/binary.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -89,7 +89,7 @@ LOGIC_CPLX_FN(double, af_or_t, ||)

template<typename T>
static T __mod(T lhs, T rhs) {
return lhs % rhs; // Same as other backends
return lhs % rhs; // Same as other backends
}

template<typename T>
Expand Down
8 changes: 4 additions & 4 deletions src/backend/cpu/err_cpu.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -9,8 +9,8 @@

#include <common/err_common.hpp>

#define CPU_NOT_SUPPORTED(message) \
do { \
throw SupportError(__AF_FUNC__, __AF_FILENAME__, __LINE__, "CPU", \
message, boost::stacktrace::stacktrace()); \
#define CPU_NOT_SUPPORTED(message) \
do { \
throw SupportError(__AF_FUNC__, __AF_FILENAME__, __LINE__, "CPU", \
message, boost::stacktrace::stacktrace()); \
} while (0)
2 changes: 1 addition & 1 deletion src/backend/cpu/ireduce.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -64,7 +64,7 @@ T ireduce_all(unsigned *loc, const Array<T> &in) {
af::dim4 dims = in.dims();
af::dim4 strides = in.strides();
const T *inPtr = in.get();
dim_t idx = 0;
dim_t idx = 0;

kernel::MinMaxOp<op, T> Op(inPtr[0], 0);

Expand Down
7 changes: 3 additions & 4 deletions src/backend/cpu/jit/BufferNode.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -178,16 +178,15 @@ class BufferNode : public TNode<T> {

virtual void modDims(const af::dim4 &newDim) override {
af::dim4 strides(1, 1, 1, 1);
for(dim_t i = 1; i < 4; ++i) {
for (dim_t i = 1; i < 4; ++i) {
strides[i] = strides[i - 1] * newDim[i - 1];
}

for(dim_t i = 0; i < 4; ++i) {
m_dims[i] = newDim[i];
for (dim_t i = 0; i < 4; ++i) {
m_dims[i] = newDim[i];
m_strides[i] = strides[i];
}
}

};

} // namespace jit
Expand Down
2 changes: 1 addition & 1 deletion src/backend/cuda/device_manager.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -514,7 +514,7 @@ void DeviceManager::checkCudaVsDriverVersion() {
debugRuntimeCheck(getLogger(), runtime, driver);

int runtime_major = runtime / 1000;
int driver_major = driver / 1000;
int driver_major = driver / 1000;

if (runtime_major > driver_major) {
string msg =
Expand Down
8 changes: 4 additions & 4 deletions src/backend/cuda/err_cuda.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -12,10 +12,10 @@
#include <common/err_common.hpp>
#include <stdio.h>

#define CUDA_NOT_SUPPORTED(message) \
do { \
throw SupportError(__AF_FUNC__, __AF_FILENAME__, __LINE__, "CUDA", \
message, boost::stacktrace::stacktrace()); \
#define CUDA_NOT_SUPPORTED(message) \
do { \
throw SupportError(__AF_FUNC__, __AF_FILENAME__, __LINE__, "CUDA", \
message, boost::stacktrace::stacktrace()); \
} while (0)

#define CU_CHECK(fn) \
Expand Down
3 changes: 1 addition & 2 deletions src/backend/cuda/kernel/ireduce.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -167,8 +167,7 @@ T ireduce_all(uint *idx, CParam<T> in) {

bool is_linear = (in.strides[0] == 1);
for (int k = 1; k < 4; k++) {
is_linear &=
(in.strides[k] == (in.strides[k - 1] * in.dims[k - 1]));
is_linear &= (in.strides[k] == (in.strides[k - 1] * in.dims[k - 1]));
}

// FIXME: Use better heuristics to get to the optimum number
Expand Down
10 changes: 4 additions & 6 deletions src/backend/cuda/kernel/reduce_by_key.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -66,8 +66,8 @@ __global__ void test_needs_reduction(int *needs_another_reduction,

if (tid < n) { k = keys_in.ptr[tid]; }

int update_key = (k == shfl_down_sync(k, 1)) &&
(tid < (n - 1)) && ((threadIdx.x % 32) < 31);
int update_key = (k == shfl_down_sync(k, 1)) && (tid < (n - 1)) &&
((threadIdx.x % 32) < 31);
int remaining_updates = any_sync(update_key);

__syncthreads();
Expand Down Expand Up @@ -267,8 +267,7 @@ __global__ static void reduce_blocks_by_key(int *reduced_block_sizes,
int eq_check, update_key;
#pragma unroll
for (int delta = 1; delta < 32; delta <<= 1) {
eq_check =
(unique_id == shfl_down_sync(unique_id, delta));
eq_check = (unique_id == shfl_down_sync(unique_id, delta));

// checks if this thread should perform a reduction
update_key =
Expand Down Expand Up @@ -494,8 +493,7 @@ __global__ static void reduce_blocks_dim_by_key(
int eq_check, update_key;
#pragma unroll
for (int delta = 1; delta < 32; delta <<= 1) {
eq_check =
(unique_id == shfl_down_sync(unique_id, delta));
eq_check = (unique_id == shfl_down_sync(unique_id, delta));

// checks if this thread should perform a reduction
update_key =
Expand Down
3 changes: 1 addition & 2 deletions src/backend/cuda/kernel/shfl_intrinsics.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -65,8 +65,7 @@ inline __device__ cfloat shfl_down_sync(cfloat var, int delta) {
}
// specialization for cdouble
template<>
inline __device__ cdouble shfl_down_sync(cdouble var,
int delta) {
inline __device__ cdouble shfl_down_sync(cdouble var, int delta) {
#if (CUDA_VERSION >= 9000)
cdouble res = {__shfl_down_sync(FULL_MASK, var.x, delta),
__shfl_down_sync(FULL_MASK, var.y, delta)};
Expand Down
9 changes: 4 additions & 5 deletions src/backend/cuda/qr.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -71,12 +71,11 @@ template<typename T>
struct mqr_buf_func_def_t {
using mqr_buf_func_def = cusolverStatus_t (*)(cusolverDnHandle_t,
cublasSideMode_t,
cublasOperation_t, int, int, int,
const T *, int, const T *, T *,
int, int *);
cublasOperation_t, int, int,
int, const T *, int,
const T *, T *, int, int *);
};


#define QR_FUNC_DEF(FUNC) \
template<typename T> \
typename FUNC##_func_def_t<T>::FUNC##_func_def FUNC##_func(); \
Expand Down Expand Up @@ -165,7 +164,7 @@ void qr(Array<T> &q, Array<T> &r, Array<T> &t, const Array<T> &in) {

CUSOLVER_CHECK(mqr_buf_func<T>()(
solverDnHandle(), CUBLAS_SIDE_LEFT, CUBLAS_OP_N, q.dims()[0],
q.dims()[1], min(M, N), in_copy.get(), in_copy.strides()[1], t.get(),
q.dims()[1], min(M, N), in_copy.get(), in_copy.strides()[1], t.get(),
q.get(), q.strides()[1], &lwork));

workspace = memAlloc<T>(lwork);
Expand Down
2 changes: 1 addition & 1 deletion src/backend/cuda/types.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -84,7 +84,7 @@ inline const char *shortname<char>(bool caps) {
}
template<>
inline const char *shortname<schar>(bool caps) {
return caps ? "A" : "a"; // TODO
return caps ? "A" : "a"; // TODO
}
template<>
inline const char *shortname<uchar>(bool caps) {
Expand Down
2 changes: 1 addition & 1 deletion src/backend/oneapi/binary.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -8,8 +8,8 @@
********************************************************/

#pragma once
#include <optypes.hpp>
#include <common/half.hpp>
#include <optypes.hpp>

using arrayfire::common::half;

Expand Down
7 changes: 3 additions & 4 deletions src/backend/oneapi/copy.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -223,10 +223,9 @@ T getScalar(const Array<T> &in) {
auto in_get = in.get();
getQueue()
.submit([&](sycl::handler &h) {
auto acc_in =
in_get->template get_access<sycl::access::mode::read>(
h, sycl::range{1},
sycl::id{static_cast<uintl>(in.getOffset())});
auto acc_in = in_get->template get_access<sycl::access::mode::read>(
h, sycl::range{1},
sycl::id{static_cast<uintl>(in.getOffset())});
h.copy(acc_in, &retVal);
})
.wait();
Expand Down
Loading
Loading