AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

Add missing interop with volatile. Fix unit tests.

Просмотр исходного кода 
[ROCm/hip commit: 0108819e2a]
Этот коммит содержится в:
Alex Voicu
2018-05-31 15:27:12 +01:00
родитель fa1952a008
Коммит 90e7799bf4
3 изменённых файлов: 124 добавлений и 45 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
projects/hip/include/hip/hcc_detail/hip_fp16.h
+45 -5
Просмотреть файл
@@ -58,6 +58,11 @@ THE SOFTWARE.
#include "hip_vector_types.h"
#include "host_defines.h"
namespace std
{
template<> struct is_floating_point<_Float16> : std::true_type {};
}
template<bool cond, typename T = void>
using Enable_if_t = typename std::enable_if<cond, T>::type;
@@ -111,6 +116,32 @@ THE SOFTWARE.
data = x.data;
return *this;
}
__host__ __device__
volatile __half& operator=(const __half_raw& x) volatile
{
data = x.data;
return *this;
}
volatile __half& operator=(const volatile __half_raw& x) volatile
{
data = x.data;
return *this;
}
__half& operator=(__half_raw&& x)
{
data = x.data;
return *this;
}
volatile __half& operator=(__half_raw&& x) volatile
{
data = x.data;
return *this;
}
volatile __half& operator=(volatile __half_raw&& x) volatile
{
data = x.data;
return *this;
}
#if !defined(__HIP_NO_HALF_CONVERSIONS__)
template<
typename T,
@@ -182,18 +213,27 @@ THE SOFTWARE.
// ACCESSORS
#if !defined(__HIP_NO_HALF_CONVERSIONS__)
__host__ __device__
operator decltype(data)() const { return data; }
__host__ __device__
operator float() const { return data; }
template<
typename T,
Enable_if_t<
std::is_floating_point<T>{} &&
!std::is_same<T, double>{}>* = nullptr>
operator T() const { return data; }
#endif
__host__ __device__
operator __half_raw() const { return __half_raw{data}; }
__host__ __device__
operator volatile __half_raw() const volatile
{
return __half_raw{data};
}
// ACCESSORS - DEVICE ONLY
#if !defined(__HIP_NO_HALF_CONVERSIONS__)
template<
typename T, Enable_if_t<std::is_integral<T>{}>* = nullptr>
__device__
operator bool() const { return data; }
operator T() const { return data; }
#endif
#if !defined(__HIP_NO_HALF_OPERATORS__)
projects/hip/tests/src/deviceLib/hipTestHalf.cpp
+12 -13
Просмотреть файл
@@ -31,10 +31,7 @@ THE SOFTWARE.
#if __HIP_ARCH_GFX803__ || __HIP_ARCH_GFX900__ || __HIP_ARCH_GFX906__
__global__
__attribute__((optnone))
void __halfMath(bool* result) {
__half a{1};
void __halfMath(bool* result, __half a) {
result[0] = __heq(__hadd(a, __half{1}), __half{2});
result[0] = __heq(__hadd_sat(a, __half{1}), __half{1}) && result[0];
result[0] = __heq(__hfma(a, __half{2}, __half{3}), __half{5}) && result[0];
@@ -56,11 +53,9 @@ bool to_bool(const __half2& x)
}
__global__
__attribute__((optnone))
void __half2Math(bool* result) {
__half2 a{1, 1};
result[0] = to_bool(__heq2(__hadd2(a, __half2{1, 1}), __half2{2, 2}));
void __half2Math(bool* result, __half2 a) {
result[0] =
to_bool(__heq2(__hadd2(a, __half2{1, 1}), __half2{2, 2}));
result[0] = to_bool(__heq2(__hadd2_sat(a, __half2{1, 1}), __half2{1, 1})) &&
result[0];
result[0] = to_bool(__heq2(
@@ -79,12 +74,14 @@ void __half2Math(bool* result) {
result[0];
}
__global__ void kernel_hisnan(__half* input, int* output) {
__global__
void kernel_hisnan(__half* input, int* output) {
int tx = threadIdx.x;
output[tx] = __hisnan(input[tx]);
}
__global__ void kernel_hisinf(__half* input, int* output) {
__global__
void kernel_hisinf(__half* input, int* output) {
int tx = threadIdx.x;
output[tx] = __hisinf(input[tx]);
}
@@ -235,13 +232,15 @@ int main() {
hipHostMalloc(&result, sizeof(result));
result[0] = false;
hipLaunchKernelGGL(__halfMath, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, result);
hipLaunchKernelGGL(
__halfMath, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, result, __half{1});
hipDeviceSynchronize();
if (!result[0]) { failed("Failed __half tests."); }
result[0] = false;
hipLaunchKernelGGL(__half2Math, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, result);
hipLaunchKernelGGL(
__half2Math, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, result, __half2{1, 1});
hipDeviceSynchronize();
if (!result[0]) { failed("Failed __half2 tests."); }
projects/hip/tests/src/deviceLib/hipTestNativeHalf.cpp
+67 -27
Просмотреть файл
@@ -28,28 +28,60 @@ THE SOFTWARE.
#include "test_common.h"
#include <type_traits>
using namespace std;
#if __HIP_ARCH_GFX803__ || __HIP_ARCH_GFX900__ || __HIP_ARCH_GFX906__
__global__
__attribute__((optnone))
void __halfTest(bool* result) {
void __halfTest(bool* result, __half a) {
// Construction
__half a{1}; result[0] = __heq(a, 1);
a = __half{1.0f}; result[0] = __heq(a, 1) && result[0];
a = __half{1.0}; result[0] = __heq(a, 1) && result[0];
a = __half{static_cast<unsigned short>(1)};
result[0] = __heq(a, 1) && result[0];
a = __half{static_cast<short>(1)}; result[0] = __heq(a, 1) && result[0];
a = __half{1u}; result[0] = __heq(a, 1) && result[0];
a = __half{1ul}; result[0] = __heq(a, 1) && result[0];
a = __half{1l}; result[0] = __heq(a, 1) && result[0];
a = __half{1ll}; result[0] = __heq(a, 1) && result[0];
a = __half{1ull}; result[0] = __heq(a, 1) && result[0];
static_assert(is_default_constructible<__half>{}, "");
static_assert(is_copy_constructible<__half>{}, "");
static_assert(is_move_constructible<__half>{}, "");
static_assert(is_constructible<__half, float>{}, "");
static_assert(is_constructible<__half, double>{}, "");
static_assert(is_constructible<__half, unsigned short>{}, "");
static_assert(is_constructible<__half, short>{}, "");
static_assert(is_constructible<__half, unsigned int>{}, "");
static_assert(is_constructible<__half, int>{}, "");
static_assert(is_constructible<__half, unsigned long>{}, "");
static_assert(is_constructible<__half, long>{}, "");
static_assert(is_constructible<__half, long long>{}, "");
static_assert(is_constructible<__half, unsigned long long>{}, "");
static_assert(is_constructible<__half, __half_raw>{}, "");
// Assignment
a = 0.0f; result[0] = __heq(a, 0) && result[0];
a = 1.0; result[0] = __heq(a, 1) && result[0];
a = __half_raw{2}; result[0] = __heq(a, 2) && result[0];
static_assert(is_copy_assignable<__half>{}, "");
static_assert(is_move_assignable<__half>{}, "");
static_assert(is_assignable<__half, float>{}, "");
static_assert(is_assignable<__half, double>{}, "");
static_assert(is_assignable<__half, unsigned short>{}, "");
static_assert(is_assignable<__half, short>{}, "");
static_assert(is_assignable<__half, unsigned int>{}, "");
static_assert(is_assignable<__half, int>{}, "");
static_assert(is_assignable<__half, unsigned long>{}, "");
static_assert(is_assignable<__half, long>{}, "");
static_assert(is_assignable<__half, long long>{}, "");
static_assert(is_assignable<__half, unsigned long long>{}, "");
static_assert(is_assignable<__half, __half_raw>{}, "");
static_assert(is_assignable<__half, volatile __half_raw&>{}, "");
static_assert(is_assignable<__half, volatile __half_raw&&>{}, "");
// Conversion
static_assert(is_convertible<__half, float>{}, "");
static_assert(is_convertible<__half, unsigned short>{}, "");
static_assert(is_convertible<__half, short>{}, "");
static_assert(is_convertible<__half, unsigned int>{}, "");
static_assert(is_convertible<__half, int>{}, "");
static_assert(is_convertible<__half, unsigned long>{}, "");
static_assert(is_convertible<__half, long>{}, "");
static_assert(is_convertible<__half, long long>{}, "");
static_assert(is_convertible<__half, bool>{}, "");
static_assert(is_convertible<__half, unsigned long long>{}, "");
static_assert(is_convertible<__half, __half_raw>{}, "");
static_assert(is_convertible<__half, volatile __half_raw>{}, "");
// Nullary
result[0] = __heq(a, +a) && result[0];
@@ -83,17 +115,23 @@ bool to_bool(const __half2& x)
return r.data.x != 0 && r.data.y != 0;
}
__global__
__attribute__((optnone))
void __half2Test(bool* result) {
void __half2Test(bool* result, __half2 a) {
// Construction
__half2 a{1};
result[0] = to_bool(__heq2(a, 1));
a = __half2{__half{1}, __half{1}};
result[0] = to_bool(__heq2(a, {1, 1})) && result[0];
static_assert(is_default_constructible<__half2>{}, "");
static_assert(is_copy_constructible<__half2>{}, "");
static_assert(is_move_constructible<__half2>{}, "");
static_assert(is_constructible<__half2, __half, __half>{}, "");
static_assert(is_constructible<__half2, __half2_raw>{}, "");
// Assignment
a = __half2_raw{2}; result[0] = to_bool(__heq2(a, {2, 2})) && result[0];
static_assert(is_copy_assignable<__half2>{}, "");
static_assert(is_move_assignable<__half2>{}, "");
static_assert(is_assignable<__half2, __half2_raw>{}, "");
// Conversion
static_assert(is_convertible<__half2, __half2_raw>{}, "");
// Nullary
result[0] = to_bool(__heq2(a, +a)) && result[0];
@@ -126,14 +164,16 @@ int main() {
bool* result{nullptr};
hipHostMalloc(&result, 1);
result[0] = false;
hipLaunchKernelGGL(__halfTest, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, result);
result[0] = true;
hipLaunchKernelGGL(
__halfTest, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, result, __half{1});
hipDeviceSynchronize();
if (!result[0]) { failed("Failed __half tests."); }
result[0] = false;
hipLaunchKernelGGL(__half2Test, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, result);
result[0] = true;
hipLaunchKernelGGL(
__half2Test, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, result, __half2{1, 1});
hipDeviceSynchronize();
if (!result[0]) { failed("Failed __half2 tests."); }