FazBrowse GitHub Viewer | Trending |
URL:
| Home
Tools: [Download Repo ZIP]   [Original HTTPS Page]

Update example function code flow for better readability · arrayfire/arrayfire@a17be56 · GitHub

  1. arrayfire
  2. arrayfire

Repository navigation

Commit a17be56

Browse files
9prady9
committed
Update example function code flow for better readability
1 parent d0d87ab commit a17be56

13 files changed

Lines changed: 122 additions & 54 deletions

‎include/af/util.h‎

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -128,7 +128,7 @@ namespace af
128128
#endif
129129

130130
// Purpose of Addition: "How to add Function" documentation
131-
AFAPI array exampleFunction(const array& in, const af_someenum_t param);
131+
AFAPI array exampleFunction(const array& a, const array& b, const af_someenum_t param);
132132

133133
#if AF_API_VERSION >= 34
134134
///
@@ -252,7 +252,9 @@ extern "C" {
252252
#endif
253253

254254
// Purpose of Addition: "How to add Function" documentation
255-
AFAPI af_err af_example_function(af_array* out, const af_array in, const af_someenum_t param);
255+
AFAPI af_err af_example_function(af_array* out,
256+
const af_array a, const af_array b,
257+
const af_someenum_t param);
256258

257259
///
258260
/// Get the version information of the library

‎src/api/c/exampleFunction.cpp‎

Lines changed: 15 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -35,7 +35,7 @@ using namespace detail; // detail is an alias to appropriate backend
3535
// change this
3636

3737
template<typename T>
38-
af_array example(const af_array& in, const af_someenum_t& param)
38+
af_array example(const af_array& a, const af_array& b, const af_someenum_t& param)
3939
{
4040
// getArray<T> function is defined in handle.hpp
4141
// and it returns backend specific Array, namely one of the following
@@ -45,25 +45,23 @@ af_array example(const af_array& in, const af_someenum_t& param)
4545
// getHandle<T> function is defined in handle.hpp takes one of the
4646
// above backend specific detail::Array<T> and returns the
4747
// universal array handle af_array
48-
return getHandle<T>( exampleFunction(getArray<T>(in), param) );
48+
return getHandle<T>( exampleFunction(getArray<T>(a), getArray<T>(b), param) );
4949
}
5050

51-
af_err af_example_function(af_array* out,
52-
const af_array in,
53-
const af_someenum_t param)
51+
af_err af_example_function(af_array* out, const af_array a, const af_array b, const af_someenum_t param)
5452
{
5553
try {
5654
af_array output = 0;
57-
ArrayInfo info = getInfo(in); // ArrayInfo is the base class which
55+
ArrayInfo info = getInfo(a); // ArrayInfo is the base class which
5856
// each backend specific Array inherits
5957
// This class stores the basic array meta-data
6058
// such as type of data, dimensions,
6159
// offsets and strides. This class is declared
6260
// in src/backend/ArrayInfo.hpp
61+
ArrayInfo info2 = getInfo(b);
6362

6463
af::dim4 dims = info.dims();
6564

66-
6765
ARG_ASSERT(2, (dims.ndims()>=0 && dims.ndims()<=3));
6866
// defined in err_common.hpp
6967
// there are other useful Macros
@@ -72,16 +70,18 @@ af_err af_example_function(af_array* out,
7270

7371
af_dtype type = info.getType();
7472

73+
ARG_ASSERT(2, (info2.getType() == type));
74+
7575
switch(type) { // Based on the data type, call backend specific
7676
// implementation
77-
case f64: output = example<double >(in, param); break;
78-
case f32: output = example<float >(in, param); break;
79-
case s32: output = example<int >(in, param); break;
80-
case u32: output = example<uint >(in, param); break;
81-
case u8: output = example<uchar >(in, param); break;
82-
case b8: output = example<char >(in, param); break;
83-
case c32: output = example<cfloat >(in, param); break;
84-
case c64: output = example<cdouble>(in, param); break;
77+
case f64: output = example<double >(a, b, param); break;
78+
case f32: output = example<float >(a, b, param); break;
79+
case s32: output = example<int >(a, b, param); break;
80+
case u32: output = example<uint >(a, b, param); break;
81+
case u8: output = example<uchar >(a, b, param); break;
82+
case b8: output = example<char >(a, b, param); break;
83+
case c32: output = example<cfloat >(a, b, param); break;
84+
case c64: output = example<cdouble>(a, b, param); break;
8585
default : TYPE_ERROR(1, type); // Another helpful macro from err_common.hpp
8686
// that helps throw type based error messages
8787
}

‎src/api/cpp/exampleFunction.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -18,13 +18,13 @@
1818
namespace af
1919
{
2020

21-
array exampleFunction(const array& in, const af_someenum_t p)
21+
array exampleFunction(const array& a, const array& b, const af_someenum_t p)
2222
{
2323
// create a temporary af_array handle
2424
af_array temp = 0;
2525

2626
// call C-API function
27-
AF_THROW( af_example_function(&temp, in.get(), p) );
27+
AF_THROW( af_example_function(&temp, a.get(), b.get(), p) );
2828

2929
// array::get() returns af_array handle for the corresponding cpp af::array
3030
return array(temp);

‎src/api/unified/util.cpp‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -51,8 +51,8 @@ af_err af_array_to_string(char **output, const char *exp, const af_array arr,
5151
return CALL(output, exp, arr, precision, transpose);
5252
}
5353

54-
af_err af_example_function(af_array* out, const af_array in, const af_someenum_t param)
54+
af_err af_example_function(af_array* out, const af_array a, const af_array b, const af_someenum_t param)
5555
{
56-
CHECK_ARRAYS(in);
57-
return CALL(out, in, param);
56+
CHECK_ARRAYS(a, b);
57+
return CALL(out, a, b, param);
5858
}

‎src/backend/cpu/exampleFunction.cpp‎

Lines changed: 34 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -22,14 +22,15 @@ namespace cpu
2222
{
2323

2424
template<typename T>
25-
Array<T> exampleFunction(const Array<T> &in, const af_someenum_t method)
25+
Array<T> exampleFunction(const Array<T> &a, const Array<T> &b, const af_someenum_t method)
2626
{
27-
in.eval(); // All input Arrays should call eval mandatorily
27+
a.eval(); // All input Arrays should call eval mandatorily
2828
// in CPU backend function implementations. Since
2929
// the cpu fns are asynchronous launches, any Arrays
3030
// that are either views/JIT nodes needs to evaluated
3131
// before they are passed onto functions that are
3232
// enqueued onto the queues.
33+
b.eval();
3334

3435
dim4 outputDims; // this should be '= in.dims();' in most cases
3536
// but would definitely depend on the type of
@@ -42,27 +43,52 @@ Array<T> exampleFunction(const Array<T> &in, const af_someenum_t method)
4243
// file to know what are the different types you
4344
// can create.
4445

45-
//dim4 in_dims = in.dims(); // you can retrieve dimensions
46+
dim4 aDims = a.dims(); // you can retrieve dimensions
47+
dim4 bDims = b.dims();
48+
dim4 oDims = out.dims();
4649

47-
//dim_t in_offset = in.getOffset(); // you can retrieve the offset - used when given array
50+
dim_t aOffset = a.getOffset(); // you can retrieve the offset - used when given array
4851
// is an sub-array pointing to some other array and
4952
// doesn't have memory of its own
53+
dim_t bOffset = b.getOffset();
5054

51-
//dim4 in_strides = in.strides(); // you can retrieve strides
55+
dim4 aStrides = a.strides(); // you can retrieve strides
56+
dim4 bStrides = b.strides();
57+
dim4 oStrides = out.strides();
5258

53-
//const T* src = in.get(); // cpu::Array<T>::get returns the pointer to the
59+
const T* src1 = a.get(); // cpu::Array<T>::get returns the pointer to the
5460
// memory allocated for that Array
61+
const T* src2 = b.get(); // cpu::Array<T>::get returns the pointer to the
62+
// memory allocated for that Array
63+
src1 += aOffset;
64+
src2 += bOffset;
5565

56-
//T* dst = out.get();
66+
T* dst = out.get();
5767

5868
// Implement your algorithm and write results to dst
69+
for(int j=0; j<aDims[1]; ++j) {
70+
for (int i=0; i<aDims[0]; ++i) {
71+
72+
int src1Idx = i*aStrides[0] + j*aStrides[1];
73+
int src2Idx = i*bStrides[0] + j*bStrides[1];
74+
int dstIdx = i*oStrides[0] + j*oStrides[1];
75+
76+
// kernel algorithm goes here
77+
switch(method) {
78+
case 1: dst[dstIdx] = src1[src1Idx] + src2[src2Idx]; break;
79+
case 2: dst[dstIdx] = src1[src1Idx] - src2[src2Idx]; break;
80+
case 3: dst[dstIdx] = src1[src1Idx] * src2[src2Idx]; break;
81+
case 4: dst[dstIdx] = src1[src1Idx] / src2[src2Idx]; break;
82+
}
83+
}
84+
}
5985

6086
return out; // return the result
6187
}
6288

6389

6490
#define INSTANTIATE(T) \
65-
template Array<T> exampleFunction<T>(const Array<T> &in, const af_someenum_t method);
91+
template Array<T> exampleFunction<T>(const Array<T> &a, const Array<T> &b, const af_someenum_t method);
6692

6793
// INSTANTIATIONS for all the types which
6894
// are present in the switch case statement

‎src/backend/cpu/exampleFunction.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -12,6 +12,6 @@
1212
namespace cpu
1313
{
1414
template<typename T>
15-
Array<T> exampleFunction(const Array<T> &in, const af_someenum_t method);
15+
Array<T> exampleFunction(const Array<T> &a, const Array<T> &b, const af_someenum_t method);
1616
}
1717

‎src/backend/cuda/exampleFunction.cu‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -27,7 +27,7 @@ namespace cuda
2727
{
2828

2929
template<typename T>
30-
Array<T> exampleFunction(const Array<T> &in, const af_someenum_t method)
30+
Array<T> exampleFunction(const Array<T> &a, const Array<T> &b, const af_someenum_t method)
3131
{
3232
dim4 outputDims; // this should be '= in.dims();' in most cases
3333
// but would definitely depend on the type of
@@ -41,14 +41,14 @@ Array<T> exampleFunction(const Array<T> &in, const af_someenum_t method)
4141
// can create.
4242

4343
// Relay the actual computation to CUDA kernel wrapper
44-
kernel::exampleFunc<T>(out, in, method);
44+
kernel::exampleFunc<T>(out, a, b, method);
4545

4646
return out; // return the result
4747
}
4848

4949

5050
#define INSTANTIATE(T) \
51-
template Array<T> exampleFunction<T>(const Array<T> &in, const af_someenum_t method);
51+
template Array<T> exampleFunction<T>(const Array<T> &a, const Array<T> &b, const af_someenum_t method);
5252

5353
// INSTANTIATIONS for all the types which
5454
// are present in the switch case statement

‎src/backend/cuda/exampleFunction.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -12,5 +12,5 @@
1212
namespace cuda
1313
{
1414
template<typename T>
15-
Array<T> exampleFunction(const Array<T> &in, const af_someenum_t method);
15+
Array<T> exampleFunction(const Array<T> &a, const Array<T> &b, const af_someenum_t method);
1616
}

‎src/backend/cuda/kernel/exampleFunction.hpp‎

Lines changed: 27 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -35,20 +35,41 @@ static const unsigned TY = 16; // Kernel Launch Config Values
3535

3636
template<typename T>
3737
__global__
38-
void exampleFuncKernel(Param<T> out, CParam<T> in, const af_someenum_t p)
38+
void exampleFuncKernel(Param<T> c, CParam<T> a, CParam<T> b, const af_someenum_t p)
3939
{
40-
// kernel implementation goes here
40+
// get current thread global identifiers along required dimensions
41+
int i = blockDim.x * blockIdx.x + threadIdx.x;
42+
int j = blockDim.y * blockIdx.y + threadIdx.y;
43+
44+
if ( i<a.dims[0] && j<a.dims[1] ) {
45+
// if needed use strides array to compute linear index of arrays
46+
int src1Idx = i*a.strides[0] + j*a.strides[1];
47+
int src2Idx = i*b.strides[0] + j*b.strides[1];
48+
int dstIdx = i*c.strides[0] + j*c.strides[1];
49+
50+
T* dst = c.ptr;
51+
const T* src1 = a.ptr;
52+
const T* src2 = b.ptr;
53+
54+
// kernel algorithm goes here
55+
switch(p) {
56+
case 1: dst[dstIdx] = src1[src1Idx] + src2[src2Idx]; break;
57+
case 2: dst[dstIdx] = src1[src1Idx] - src2[src2Idx]; break;
58+
case 3: dst[dstIdx] = src1[src1Idx] * src2[src2Idx]; break;
59+
case 4: dst[dstIdx] = src1[src1Idx] / src2[src2Idx]; break;
60+
}
61+
}
4162
}
4263

4364

4465
template <typename T> // CUDA kernel wrapper function
45-
void exampleFunc(Param<T> out, CParam<T> in, const af_someenum_t p)
66+
void exampleFunc(Param<T> c, CParam<T> a, CParam<T> b, const af_someenum_t p)
4667
{
4768

4869
dim3 threads(TX, TY, 1); // set your cuda launch config for blocks
4970

50-
int blk_x = divup(out.dims[0], threads.x);
51-
int blk_y = divup(out.dims[1], threads.y);
71+
int blk_x = divup(c.dims[0], threads.x);
72+
int blk_y = divup(c.dims[1], threads.y);
5273
dim3 blocks(blk_x, blk_y); // set your opencl launch config for grid
5374

5475
// launch your kernel
@@ -61,7 +82,7 @@ void exampleFunc(Param<T> out, CParam<T> in, const af_someenum_t p)
6182
//
6283
// CUDA_LAUNCH_SMEM takes in an additional parameter, size of shared memory, after
6384
// threads paramters, which are then followed by kernel parameters
64-
CUDA_LAUNCH((exampleFuncKernel<T>), blocks, threads, out, in, p);
85+
CUDA_LAUNCH((exampleFuncKernel<T>), blocks, threads, c, a, b, p);
6586

6687
POST_LAUNCH_CHECK(); // Macro for post kernel launch checks
6788
// these checks are carried ONLY IN DEBUG mode

‎src/backend/opencl/exampleFunction.cpp‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -27,7 +27,7 @@ namespace opencl
2727
{
2828

2929
template<typename T>
30-
Array<T> exampleFunction(const Array<T> &in, const af_someenum_t method)
30+
Array<T> exampleFunction(const Array<T> &a, const Array<T> &b, const af_someenum_t method)
3131
{
3232
dim4 outputDims; // this should be '= in.dims();' in most cases
3333
// but would definitely depend on the type of
@@ -41,14 +41,14 @@ Array<T> exampleFunction(const Array<T> &in, const af_someenum_t method)
4141
// can create.
4242

4343
// Relay the actual computation to OpenCL kernel wrapper
44-
kernel::exampleFunc<T>(out, in, method);
44+
kernel::exampleFunc<T>(out, a, b, method);
4545

4646
return out; // return the result
4747
}
4848

4949

5050
#define INSTANTIATE(T) \
51-
template Array<T> exampleFunction<T>(const Array<T> &in, const af_someenum_t method);
51+
template Array<T> exampleFunction<T>(const Array<T> &a, const Array<T> &b, const af_someenum_t method);
5252

5353
// INSTANTIATIONS for all the types which
5454
// are present in the switch case statement

0 commit comments

Comments
 (0)
Back | FazBrowse Home | New Git URL