Skip to content

Commit e4e4c18

Browse files
author
Pradeep
committed
Proper error handling added to transpose
transpose unit tests have been modified to catch updated error messages
1 parent 3a3efc1 commit e4e4c18

9 files changed

Lines changed: 86 additions & 85 deletions

File tree

‎src/backend/cpu/transpose.cpp‎

Lines changed: 42 additions & 42 deletions
Original file line numberDiff line numberDiff line change
@@ -11,59 +11,59 @@ using af::dim4;
1111
namespace cpu
1212
{
1313

14-
static inline unsigned getIdx(const dim4 &strides,
15-
int i, int j = 0, int k = 0, int l = 0)
16-
{
17-
return (l * strides[3] +
18-
k * strides[2] +
19-
j * strides[1] +
20-
i );
21-
}
14+
static inline unsigned getIdx(const dim4 &strides,
15+
int i, int j = 0, int k = 0, int l = 0)
16+
{
17+
return (l * strides[3] +
18+
k * strides[2] +
19+
j * strides[1] +
20+
i );
21+
}
2222

23-
template<typename T>
24-
Array<T> * transpose(const Array<T> &in)
25-
{
26-
const dim4 inDims = in.dims();
23+
template<typename T>
24+
Array<T> * transpose(const Array<T> &in)
25+
{
26+
const dim4 inDims = in.dims();
2727

28-
dim4 outDims = dim4(inDims[1],inDims[0],inDims[2],inDims[3]);
28+
dim4 outDims = dim4(inDims[1],inDims[0],inDims[2],inDims[3]);
2929

30-
// create an array with first two dimensions swapped
31-
Array<T>* out = createEmptyArray<T>(outDims);
30+
// create an array with first two dimensions swapped
31+
Array<T>* out = createEmptyArray<T>(outDims);
3232

33-
// get data pointers for input and output Arrays
34-
T* outData = out->get();
35-
const T* inData = in.get();
33+
// get data pointers for input and output Arrays
34+
T* outData = out->get();
35+
const T* inData = in.get();
3636

37-
for (uint k=0; k<outDims[2]; ++k) {
38-
// Outermost loop handles batch mode
39-
// if input has no data along third dimension
40-
// this loop runs only once
41-
for (uint j=0; j<outDims[1]; ++j) {
42-
for (uint i=0; i<outDims[0]; ++i) {
43-
// calculate array indices based on offsets and strides
44-
// the helper getIdx takes care of indices
45-
int inIdx = getIdx(in.strides(),j,i,k);
46-
int outIdx = getIdx(out->strides(),i,j,k);
47-
outData[outIdx] = inData[inIdx];
48-
}
37+
for (uint k=0; k<outDims[2]; ++k) {
38+
// Outermost loop handles batch mode
39+
// if input has no data along third dimension
40+
// this loop runs only once
41+
for (uint j=0; j<outDims[1]; ++j) {
42+
for (uint i=0; i<outDims[0]; ++i) {
43+
// calculate array indices based on offsets and strides
44+
// the helper getIdx takes care of indices
45+
int inIdx = getIdx(in.strides(),j,i,k);
46+
int outIdx = getIdx(out->strides(),i,j,k);
47+
outData[outIdx] = inData[inIdx];
4948
}
50-
// outData and inData pointers doesn't need to be
51-
// offset as the getIdx function is taking care
52-
// of the batch parameter
5349
}
54-
return out;
50+
// outData and inData pointers doesn't need to be
51+
// offset as the getIdx function is taking care
52+
// of the batch parameter
5553
}
54+
return out;
55+
}
5656

5757
#define INSTANTIATE(T)\
5858
template Array<T> * transpose(const Array<T> &in);
5959

60-
INSTANTIATE(float)
61-
INSTANTIATE(cfloat)
62-
INSTANTIATE(double)
63-
INSTANTIATE(cdouble)
64-
INSTANTIATE(char)
65-
INSTANTIATE(int)
66-
INSTANTIATE(uint)
67-
INSTANTIATE(uchar)
60+
INSTANTIATE(float )
61+
INSTANTIATE(cfloat )
62+
INSTANTIATE(double )
63+
INSTANTIATE(cdouble)
64+
INSTANTIATE(char )
65+
INSTANTIATE(int )
66+
INSTANTIATE(uint )
67+
INSTANTIATE(uchar )
6868

6969
}

‎src/backend/cpu/transpose.hpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,7 @@
33
namespace cpu
44
{
55

6-
template<typename T>
7-
Array<T> * transpose(const Array<T> &in);
6+
template<typename T>
7+
Array<T> * transpose(const Array<T> &in);
88

99
}

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

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2,9 +2,11 @@
22
#include <backend.hpp>
33
#include <dispatch.hpp>
44
#include <Param.hpp>
5+
#include <debug_cuda.hpp>
56

67
namespace cuda
78
{
9+
810
namespace kernel
911
{
1012

@@ -84,6 +86,10 @@ namespace kernel
8486
(transpose< T, true >)<<< blocks,threads >>>(out, in, blk_x);
8587
else
8688
(transpose< T, false>)<<< blocks,threads >>>(out, in, blk_x);
89+
90+
POST_LAUNCH_CHECK();
8791
}
92+
8893
}
94+
8995
}

‎src/backend/cuda/transpose.cu‎

Lines changed: 18 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -8,31 +8,31 @@ using af::dim4;
88
namespace cuda
99
{
1010

11-
template<typename T>
12-
Array<T> * transpose(const Array<T> &in)
13-
{
14-
const dim4 inDims = in.dims();
15-
const dim4 inStrides= in.strides();
11+
template<typename T>
12+
Array<T> * transpose(const Array<T> &in)
13+
{
14+
const dim4 inDims = in.dims();
15+
const dim4 inStrides= in.strides();
1616

17-
dim4 outDims = dim4(inDims[1],inDims[0],inDims[2],inDims[3]);
17+
dim4 outDims = dim4(inDims[1],inDims[0],inDims[2],inDims[3]);
1818

19-
Array<T>* out = createEmptyArray<T>(outDims);
19+
Array<T>* out = createEmptyArray<T>(outDims);
2020

21-
kernel::transpose<T>(*out, in, inDims.ndims());
21+
kernel::transpose<T>(*out, in, inDims.ndims());
2222

23-
return out;
24-
}
23+
return out;
24+
}
2525

2626
#define INSTANTIATE(T)\
2727
template Array<T> * transpose(const Array<T> &in);
2828

29-
INSTANTIATE(float)
30-
INSTANTIATE(cfloat)
31-
INSTANTIATE(double)
32-
INSTANTIATE(cdouble)
33-
INSTANTIATE(char)
34-
INSTANTIATE(int)
35-
INSTANTIATE(uint)
36-
INSTANTIATE(uchar)
29+
INSTANTIATE(float )
30+
INSTANTIATE(cfloat )
31+
INSTANTIATE(double )
32+
INSTANTIATE(cdouble)
33+
INSTANTIATE(char )
34+
INSTANTIATE(int )
35+
INSTANTIATE(uint )
36+
INSTANTIATE(uchar )
3737

3838
}

‎src/backend/cuda/transpose.hpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,7 @@
33
namespace cuda
44
{
55

6-
template<typename T>
7-
Array<T> * transpose(const Array<T> &in);
6+
template<typename T>
7+
Array<T> * transpose(const Array<T> &in);
88

99
}

‎src/backend/opencl/transpose.cpp‎

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -21,13 +21,13 @@ Array<T> * transpose(const Array<T> &in)
2121
#define INSTANTIATE(T)\
2222
template Array<T> * transpose(const Array<T> &in);
2323

24-
INSTANTIATE(float)
25-
INSTANTIATE(cfloat)
26-
INSTANTIATE(double)
24+
INSTANTIATE(float )
25+
INSTANTIATE(cfloat )
26+
INSTANTIATE(double )
2727
INSTANTIATE(cdouble)
28-
INSTANTIATE(char)
29-
INSTANTIATE(int)
30-
INSTANTIATE(uint)
31-
INSTANTIATE(uchar)
28+
INSTANTIATE(char )
29+
INSTANTIATE(int )
30+
INSTANTIATE(uint )
31+
INSTANTIATE(uchar )
3232

3333
}

‎src/backend/opencl/transpose.hpp‎

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -2,6 +2,8 @@
22

33
namespace opencl
44
{
5-
template<typename T>
6-
Array<T> * transpose(const Array<T> &in);
5+
6+
template<typename T>
7+
Array<T> * transpose(const Array<T> &in);
8+
79
}

‎src/backend/transpose.cpp‎

Lines changed: 4 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -17,16 +17,12 @@ static inline af_array transpose(const af_array in)
1717

1818
af_err af_transpose(af_array *out, af_array in)
1919
{
20-
af_err ret = AF_ERR_RUNTIME;
21-
2220
try {
2321
ArrayInfo info = getInfo(in);
2422
af_dtype type = info.getType();
2523
af::dim4 dims = info.dims();
2624

27-
if (dims.ndims()>3) {
28-
return AF_ERR_ARG;
29-
}
25+
DIM_ASSERT(1, (dims.ndims()<=3));
3026

3127
if (dims[0]==1 || dims[1]==1) {
3228
// for a vector OR a batch of vectors
@@ -45,14 +41,11 @@ af_err af_transpose(af_array *out, af_array in)
4541
case s32: output = transpose<int>(in); break;
4642
case u32: output = transpose<uint>(in); break;
4743
case u8 : output = transpose<uchar>(in); break;
48-
default : ret = AF_ERR_NOT_SUPPORTED; break;
49-
}
50-
if (ret!=AF_ERR_NOT_SUPPORTED) {
51-
std::swap(*out,output);
52-
ret = AF_SUCCESS;
44+
default : TYPE_ERROR(1, type);
5345
}
46+
std::swap(*out,output);
5447
}
5548
CATCHALL;
5649

57-
return ret;
50+
return AF_SUCCESS;
5851
}

‎test/transpose.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -157,7 +157,7 @@ TYPED_TEST(Transpose,InvalidArgs)
157157
af::dim4 newDims(5,5,2,2);
158158
ASSERT_EQ(AF_SUCCESS, af_create_array(&inArray, &(in[0].front()), newDims.ndims(), newDims.get(), (af_dtype) af::dtype_traits<TypeParam>::af_type));
159159

160-
ASSERT_EQ(AF_ERR_ARG, af_transpose(&outArray,inArray));
160+
ASSERT_EQ(AF_ERR_SIZE, af_transpose(&outArray,inArray));
161161
}
162162

163163
TYPED_TEST(Transpose,SubRef)

0 commit comments

Comments
 (0)