Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Half fixes for ATen #2070

Closed
wants to merge 2 commits into from
Closed

Half fixes for ATen #2070

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
740cb19
Half fixes for ATen
csarofeen Jul 12, 2017
98c3147
Had to include both cuda,cuda_runtime .h files to get versioning corr…
csarofeen Jul 13, 2017
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Jump to
Jump to file
Failed to load files.
Diff view
Diff view
27 changes: 25 additions & 2 deletions torch/lib/ATen/Half.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,6 +2,7 @@

#include<stdint.h>
#ifdef AT_CUDA_ENABLED
#include <cuda.h>
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#endif
Expand All @@ -24,7 +25,15 @@ template<typename To, typename From> To convert(From f) {
typedef struct AT_ALIGN(2) {
unsigned short x;
#ifdef AT_CUDA_ENABLED
operator half() { return half { x }; }
#if CUDA_VERSION < 9000
operator half() { return half{ x }; }
#else
operator half() {
__half_raw x_raw;
x_raw.x = x;
return half(x_raw);
}
#endif
#endif
operator double();
} Half;
Expand All @@ -41,11 +50,25 @@ inline Half::operator double() {
template<> half convert(double d);
#endif


template<typename To, typename From>
static inline To HalfFix(From h) {
return To { h.x };
}

#ifdef AT_CUDA_ENABLED
#if CUDA_VERSION >= 9000
template<>
inline __half HalfFix<__half, Half>(Half h) {
__half_raw raw;
raw.x = h.x;
return __half { raw };
}

template<>
inline Half HalfFix<Half, __half>(__half h) {
__half_raw raw(h);
return Half { raw.x };
}
#endif
#endif
}
9 changes: 8 additions & 1 deletion torch/lib/ATen/Scalar.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,7 +24,14 @@ template<> int64_t convert(Half f) {

#ifdef AT_CUDA_ENABLED
template<> half convert(double d) {
return half { convert<Half,double>(d).x };

#if CUDA_VERSION < 9000
return half {convert<Half,double>(d).x};
#else
__half_raw raw;
raw.x = convert<Half,double>(d).x;
return half {raw};
#endif
}
#endif

Expand Down
7 changes: 6 additions & 1 deletion torch/lib/ATen/Scalar.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,7 +27,12 @@ class Scalar {
#ifdef AT_CUDA_ENABLED
Scalar(half vv)
: tag(Tag::HAS_d) {
v.d = convert<double,Half>(Half{vv.x});
#if CUDA_VERSION < 9000
v.d = convert<double, Half>(Half{vv.x});
#else
__half_raw vv_raw(vv);
v.d = convert<double,Half>(Half{vv_raw.x});
#endif
}
#endif

Expand Down