Skip to content

Commit

Permalink
Replaces MSHADOW_HALF2_INLINE with MSHADOW_XINLINE (apache#249)
Browse files Browse the repository at this point in the history
  • Loading branch information
ap-hynninen authored and piiswrong committed May 16, 2017
1 parent 1d633cb commit c037b06
Showing 1 changed file with 12 additions and 14 deletions.
26 changes: 12 additions & 14 deletions mshadow/half2.h
Original file line number Diff line number Diff line change
Expand Up @@ -11,10 +11,8 @@
#if (defined(__CUDACC__) && __CUDA_ARCH__ >= 530 && MSHADOW_USE_CUDA && CUDA_VERSION >= 7050)
#define MSHADOW_CUDA_HALF2 1
#include <cuda_fp16.h>
#define MSHADOW_HALF2_INLINE MSHADOW_FORCE_INLINE __device__
#else
#define MSHADOW_CUDA_HALF2 0
#define MSHADOW_HALF2_INLINE MSHADOW_FORCE_INLINE
#endif

/*! \brief namespace for mshadow */
Expand All @@ -36,18 +34,18 @@ class half2_t {
half_t half_t2[2];
#endif

MSHADOW_HALF2_INLINE half2_t() {}
MSHADOW_XINLINE half2_t() {}

#if MSHADOW_CUDA_HALF2
MSHADOW_HALF2_INLINE explicit half2_t(half2 a) : half2_(a) {}
MSHADOW_XINLINE explicit half2_t(half2 a) : half2_(a) {}
#else
MSHADOW_HALF2_INLINE explicit half2_t(half_t a, half_t b) {
MSHADOW_XINLINE explicit half2_t(half_t a, half_t b) {
half_t2[0] = a;
half_t2[1] = b;
}
#endif

MSHADOW_HALF2_INLINE explicit half2_t(int a) {
MSHADOW_XINLINE explicit half2_t(int a) {
#if MSHADOW_CUDA_HALF2
half2_ = __half2half2(__int2half_rz(a));
#else
Expand All @@ -56,19 +54,19 @@ class half2_t {
#endif
}

MSHADOW_HALF2_INLINE half2_t operator+() {
MSHADOW_XINLINE half2_t operator+() {
return *this;
}

MSHADOW_HALF2_INLINE half2_t operator-() {
MSHADOW_XINLINE half2_t operator-() {
#if MSHADOW_CUDA_HALF2
return half2_t(__hneg2(half2_));
#else
return half2_t(-half_t2[0], -half_t2[1]);
#endif
}

MSHADOW_HALF2_INLINE half2_t operator=(const half2_t& a) {
MSHADOW_XINLINE half2_t operator=(const half2_t& a) {
#if MSHADOW_CUDA_HALF2
half2_ = a.half2_;
#else
Expand All @@ -85,39 +83,39 @@ class half2_t {
};

/*! \brief overloaded + operator for half2_t */
MSHADOW_HALF2_INLINE half2_t operator+(half2_t a, half2_t b) {
MSHADOW_XINLINE half2_t operator+(half2_t a, half2_t b) {
#if MSHADOW_CUDA_HALF2
return half2_t(__hadd2(a.half2_, b.half2_));
#else
return half2_t(a.half_t2[0] + b.half_t2[0], a.half_t2[1] + b.half_t2[1]);
#endif
}
/*! \brief overloaded - operator for half2_t */
MSHADOW_HALF2_INLINE half2_t operator-(half2_t a, half2_t b) {
MSHADOW_XINLINE half2_t operator-(half2_t a, half2_t b) {
#if MSHADOW_CUDA_HALF2
return half2_t(__hsub2(a.half2_, b.half2_));
#else
return half2_t(a.half_t2[0] - b.half_t2[0], a.half_t2[1] - b.half_t2[1]);
#endif
}
/*! \brief overloaded * operator for half2_t */
MSHADOW_HALF2_INLINE half2_t operator*(half2_t a, half2_t b) {
MSHADOW_XINLINE half2_t operator*(half2_t a, half2_t b) {
#if MSHADOW_CUDA_HALF2
return half2_t(__hmul2(a.half2_, b.half2_));
#else
return half2_t(a.half_t2[0] * b.half_t2[0], a.half_t2[1] * b.half_t2[1]);
#endif
}
/*! \brief overloaded / operator for half2_t */
MSHADOW_HALF2_INLINE half2_t operator/(half2_t a, half2_t b) {
MSHADOW_XINLINE half2_t operator/(half2_t a, half2_t b) {
#if MSHADOW_CUDA_HALF2
return half2_t(h2div(a.half2_, b.half2_));
#else
return half2_t(a.half_t2[0] / b.half_t2[0], a.half_t2[1] / b.half_t2[1]);
#endif
}
/*! \brief overloaded == operator for half2_t */
MSHADOW_HALF2_INLINE bool operator==(half2_t a, half2_t b) {
MSHADOW_XINLINE bool operator==(half2_t a, half2_t b) {
#if MSHADOW_CUDA_HALF2
return __hbeq2(a.half2_, b.half2_);
#else
Expand Down

0 comments on commit c037b06

Please sign in to comment.