[arrayfire] 267/408: Style changes to code in unwrap
Ghislain Vaillant
ghisvail-guest at moszumanska.debian.org
Mon Sep 21 19:12:11 UTC 2015
This is an automated email from the git hooks/post-receive script.
ghisvail-guest pushed a commit to branch debian/sid
in repository arrayfire.
commit 78b9b99ded6e54d5143e6af3081b41ba1ce17fce
Author: Pavan Yalamanchili <pavan at arrayfire.com>
Date: Fri Aug 14 14:34:09 2015 -0400
Style changes to code in unwrap
---
src/backend/cuda/kernel/unwrap.hpp | 19 +++++++---------
src/backend/opencl/kernel/unwrap.cl | 11 ++++------
src/backend/opencl/kernel/unwrap.hpp | 42 ++++++++++++++++++++----------------
3 files changed, 36 insertions(+), 36 deletions(-)
diff --git a/src/backend/cuda/kernel/unwrap.hpp b/src/backend/cuda/kernel/unwrap.hpp
index 4083ba2..e012340 100644
--- a/src/backend/cuda/kernel/unwrap.hpp
+++ b/src/backend/cuda/kernel/unwrap.hpp
@@ -20,7 +20,7 @@ namespace cuda
///////////////////////////////////////////////////////////////////////////
// Resize Kernel
///////////////////////////////////////////////////////////////////////////
- template<typename T, int threads>
+ template<typename T, int TX>
__global__
void unwrap_kernel(Param<T> out, CParam<T> in,
const dim_t wx, const dim_t wy, const dim_t sx, const dim_t sy,
@@ -54,13 +54,11 @@ namespace cuda
T* optr = out.ptr + cOut + colId * out.strides[1];
const T* iptr = in.ptr + cIn;
- bool cond = false;
- if(spx >= 0 && spx + wx < in.dims[0] && spy >= 0 && spy + wy < in.dims[1])
- cond = true;
+ bool cond = (spx >= 0 && spx + wx < in.dims[0] && spy >= 0 && spy + wy < in.dims[1]);
for(int i = 0; i < repsPerColumn; i++) {
// Compute output index local to column
- const dim_t colIndex = i * threads + threadIdx.x;
+ const dim_t colIndex = i * TX + threadIdx.x;
if(colIndex >= out.dims[0])
return;
@@ -75,12 +73,12 @@ namespace cuda
const dim_t outIdx = (y * wx + x) * out.strides[0];
// Copy
+ T val = scalar<T>(0.0);
if(cond || (xpad >= 0 && xpad < in.dims[0] && ypad >= 0 && ypad < in.dims[1])) {
- const dim_t inIdx = ypad * in.strides[1] + xpad * in.strides[0];
- optr[outIdx] = iptr[inIdx];
- } else {
- optr[outIdx] = scalar<T>(0.0);
+ const dim_t inIdx = ypad * in.strides[1] + xpad;
+ val = iptr[inIdx];
}
+ optr[outIdx] = val;
}
}
@@ -101,9 +99,8 @@ namespace cuda
dim3 blocks(divup(out.dims[1], threads.y), out.dims[2] * out.dims[3]);
CUDA_LAUNCH((unwrap_kernel<T, TX>), blocks, threads,
- out, in, wx, wy, sx, sy, px, py, nx, repsPerColumn);
+ out, in, wx, wy, sx, sy, px, py, nx, repsPerColumn);
POST_LAUNCH_CHECK();
}
}
}
-
diff --git a/src/backend/opencl/kernel/unwrap.cl b/src/backend/opencl/kernel/unwrap.cl
index 8a15f0e..61aab1a 100644
--- a/src/backend/opencl/kernel/unwrap.cl
+++ b/src/backend/opencl/kernel/unwrap.cl
@@ -54,10 +54,7 @@ void unwrap_kernel(__global T *d_out, const KParam out,
__global T* optr = d_out + cOut + colId * out.strides[1];
__global const T* iptr = d_in + cIn + in.offset;
- bool cond = false;
- if(spx >= 0 && spx + wx < in.dims[0] && spy >= 0 && spy + wy < in.dims[1])
- cond = true;
-
+ bool cond = (spx >= 0 && spx + wx < in.dims[0] && spy >= 0 && spy + wy < in.dims[1]);
for(int i = 0; i < repsPerColumn; i++) {
// Compute output index local to column
@@ -76,11 +73,11 @@ void unwrap_kernel(__global T *d_out, const KParam out,
const dim_t outIdx = (y * wx + x) * out.strides[0];
// Copy
+ T val = ZERO;
if(cond || (xpad >= 0 && xpad < in.dims[0] && ypad >= 0 && ypad < in.dims[1])) {
const dim_t inIdx = ypad * in.strides[1] + xpad * in.strides[0];
- optr[outIdx] = iptr[inIdx];
- } else {
- set_scalar(optr[outIdx], 0);
+ val = iptr[inIdx];
}
+ optr[outIdx] = val;
}
}
diff --git a/src/backend/opencl/kernel/unwrap.hpp b/src/backend/opencl/kernel/unwrap.hpp
index d9f7d49..51c3297 100644
--- a/src/backend/opencl/kernel/unwrap.hpp
+++ b/src/backend/opencl/kernel/unwrap.hpp
@@ -17,6 +17,9 @@
#include <dispatch.hpp>
#include <Param.hpp>
#include <debug_opencl.hpp>
+#include <type_util.hpp>
+#include <math.hpp>
+#include "config.hpp"
using cl::Buffer;
using cl::Program;
@@ -42,27 +45,30 @@ namespace opencl
int device = getActiveDeviceId();
std::call_once( compileFlags[device], [device] () {
- std::ostringstream options;
- options << " -D T=" << dtype_traits<T>::getName();
- options << " -D TX=" << TX;
- if((af_dtype) dtype_traits<T>::af_type == c32 ||
- (af_dtype) dtype_traits<T>::af_type == c64) {
- options << " -D CPLX=1";
- } else {
- options << " -D CPLX=0";
- }
+ ToNum<T> toNum;
+ std::ostringstream options;
+ options << " -D ZERO=" << toNum(scalar<T>(0));
+ options << " -D T=" << dtype_traits<T>::getName();
+ options << " -D TX=" << TX;
- if (std::is_same<T, double>::value ||
- std::is_same<T, cdouble>::value) {
- options << " -D USE_DOUBLE";
- }
+ if((af_dtype) dtype_traits<T>::af_type == c32 ||
+ (af_dtype) dtype_traits<T>::af_type == c64) {
+ options << " -D CPLX=1";
+ } else {
+ options << " -D CPLX=0";
+ }
- Program prog;
- buildProgram(prog, unwrap_cl, unwrap_cl_len, options.str());
- unwrapProgs[device] = new Program(prog);
- unwrapKernels[device] = new Kernel(*unwrapProgs[device], "unwrap_kernel");
- });
+ if (std::is_same<T, double>::value ||
+ std::is_same<T, cdouble>::value) {
+ options << " -D USE_DOUBLE";
+ }
+
+ Program prog;
+ buildProgram(prog, unwrap_cl, unwrap_cl_len, options.str());
+ unwrapProgs[device] = new Program(prog);
+ unwrapKernels[device] = new Kernel(*unwrapProgs[device], "unwrap_kernel");
+ });
auto unwrapOp = make_kernel<Buffer, const KParam, const Buffer, const KParam,
const dim_t, const dim_t, const dim_t, const dim_t,
--
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