Skip to content

Commit 4b22852

Browse files
zasdfgbnmzasdfgbnm
committed
save
1 parent 030b7b5 commit 4b22852

File tree

3 files changed

+70
-128
lines changed

3 files changed

+70
-128
lines changed

‎aten/src/ATen/test/cuda_bfloat16_test.cu

Lines changed: 3 additions & 64 deletions
Original file line numberDiff line numberDiff line change
@@ -2,85 +2,24 @@
22

33
#include <ATen/ATen.h>
44
#include <ATen/cuda/CUDAContext.h>
5-
#include <ATen/cuda/NumericLimits.cuh>
65

76
#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
87
#include <cuda.h>
98
#include <cuda_bf16.h>
109
#include <cuda_runtime.h>
10+
#include <ATen/test/cuda_lowfp_test.cuh>
1111

1212
#include <assert.h>
1313

1414
using namespace at;
1515

16-
__device__ void test(){
17-
// test bfloat16 construction and implicit conversions in device
18-
assert(BFloat16(3) == BFloat16(3.0f));
19-
assert(static_cast<BFloat16>(3.0f) == BFloat16(3.0f));
20-
// there is no float <=> __nv_bfloat16 implicit conversion
21-
assert(static_cast<BFloat16>(3.0f) == 3.0f);
16+
__global__ void kernel(){
17+
test<BFloat16>();
2218

2319
__nv_bfloat16 a = __float2bfloat16(3.0f);
2420
__nv_bfloat16 b = __float2bfloat16(2.0f);
2521
__nv_bfloat16 c = a - BFloat16(b);
2622
assert(static_cast<BFloat16>(c) == BFloat16(1.0));
27-
28-
// asserting if the functions used on
29-
// bfloat16 types give almost equivalent results when using
30-
// functions on double.
31-
// The purpose of these asserts are to test the device side
32-
// bfloat16 API for the common mathematical functions.
33-
// Note: When calling std math functions from device, don't
34-
// use the std namespace, but just "::" so that the function
35-
// gets resolved from nvcc math_functions.hpp
36-
37-
float threshold = 0.00001;
38-
assert(::abs(::lgamma(BFloat16(10.0)) - ::lgamma(10.0f)) <= threshold);
39-
assert(::abs(::exp(BFloat16(1.0)) - ::exp(1.0f)) <= threshold);
40-
assert(::abs(::log(BFloat16(1.0)) - ::log(1.0f)) <= threshold);
41-
assert(::abs(::log10(BFloat16(1000.0)) - ::log10(1000.0f)) <= threshold);
42-
assert(::abs(::log1p(BFloat16(0.0)) - ::log1p(0.0f)) <= threshold);
43-
assert(::abs(::log2(BFloat16(1000.0)) - ::log2(1000.0f)) <= threshold);
44-
assert(::abs(::expm1(BFloat16(1.0)) - ::expm1(1.0f)) <= threshold);
45-
assert(::abs(::cos(BFloat16(0.0)) - ::cos(0.0f)) <= threshold);
46-
assert(::abs(::sin(BFloat16(0.0)) - ::sin(0.0f)) <= threshold);
47-
assert(::abs(::sqrt(BFloat16(100.0)) - ::sqrt(100.0f)) <= threshold);
48-
assert(::abs(::ceil(BFloat16(2.4)) - ::ceil(2.4f)) <= threshold);
49-
assert(::abs(::floor(BFloat16(2.7)) - ::floor(2.7f)) <= threshold);
50-
assert(::abs(::trunc(BFloat16(2.7)) - ::trunc(2.7f)) <= threshold);
51-
assert(::abs(::acos(BFloat16(-1.0)) - ::acos(-1.0f)) <= threshold);
52-
assert(::abs(::cosh(BFloat16(1.0)) - ::cosh(1.0f)) <= threshold);
53-
assert(::abs(::acosh(BFloat16(1.0)) - ::acosh(1.0f)) <= threshold);
54-
assert(::abs(::acosh(BFloat16(1.0)) - ::acosh(1.0f)) <= threshold);
55-
assert(::abs(::asinh(BFloat16(1.0)) - ::asinh(1.0f)) <= threshold);
56-
assert(::abs(::atanh(BFloat16(0.5)) - ::atanh(0.5f)) <= threshold);
57-
assert(::abs(::asin(BFloat16(1.0)) - ::asin(1.0f)) <= threshold);
58-
assert(::abs(::sinh(BFloat16(1.0)) - ::sinh(1.0f)) <= threshold);
59-
assert(::abs(::asinh(BFloat16(1.0)) - ::asinh(1.0f)) <= threshold);
60-
assert(::abs(::tan(BFloat16(0.0)) - ::tan(0.0f)) <= threshold);
61-
assert(::abs(::atan(BFloat16(1.0)) - ::atan(1.0f)) <= threshold);
62-
assert(::abs(::tanh(BFloat16(1.0)) - ::tanh(1.0f)) <= threshold);
63-
assert(::abs(::erf(BFloat16(10.0)) - ::erf(10.0f)) <= threshold);
64-
assert(::abs(::erfc(BFloat16(10.0)) - ::erfc(10.0f)) <= threshold);
65-
assert(::abs(::abs(BFloat16(-3.0)) - ::abs(-3.0f)) <= threshold);
66-
assert(::abs(::round(BFloat16(2.3)) - ::round(2.3f)) <= threshold);
67-
assert(::abs(::pow(BFloat16(2.0), BFloat16(10.0)) - ::pow(2.0f, 10.0f)) <= threshold);
68-
assert(
69-
::abs(::atan2(BFloat16(7.0), BFloat16(0.0)) - ::atan2(7.0f, 0.0f)) <= threshold);
70-
// note: can't use namespace on isnan and isinf in device code
71-
#ifdef _MSC_VER
72-
// Windows requires this explicit conversion. The reason is unclear
73-
// related issue with clang: https://reviews.llvm.org/D37906
74-
assert(::abs(::isnan((float)BFloat16(0.0)) - ::isnan(0.0f)) <= threshold);
75-
assert(::abs(::isinf((float)BFloat16(0.0)) - ::isinf(0.0f)) <= threshold);
76-
#else
77-
assert(::abs(::isnan(BFloat16(0.0)) - ::isnan(0.0f)) <= threshold);
78-
assert(::abs(::isinf(BFloat16(0.0)) - ::isinf(0.0f)) <= threshold);
79-
#endif
80-
}
81-
82-
__global__ void kernel(){
83-
test();
8423
}
8524

8625
void launch_function(){

‎aten/src/ATen/test/cuda_half_test.cu

Lines changed: 4 additions & 64 deletions
Original file line numberDiff line numberDiff line change
@@ -2,83 +2,23 @@
22

33
#include <ATen/ATen.h>
44
#include <ATen/cuda/CUDAContext.h>
5-
#include <ATen/cuda/NumericLimits.cuh>
5+
66
#include <cuda.h>
77
#include <cuda_fp16.h>
88
#include <cuda_runtime.h>
9+
#include <ATen/test/cuda_lowfp_test.cuh>
910

1011
#include <assert.h>
1112

1213
using namespace at;
1314

14-
__device__ void test(){
15-
// test half construction and implicit conversions in device
16-
assert(Half(3) == Half(3.0f));
17-
assert(static_cast<Half>(3.0f) == Half(3.0f));
18-
// there is no float <=> __half implicit conversion
19-
assert(static_cast<Half>(3.0f) == 3.0f);
15+
__global__ void kernel(){
16+
test<Half>();
2017

2118
__half a = __float2half(3.0f);
2219
__half b = __float2half(2.0f);
2320
__half c = a - Half(b);
2421
assert(static_cast<Half>(c) == Half(1.0));
25-
26-
// asserting if the functions used on
27-
// half types give almost equivalent results when using
28-
// functions on double.
29-
// The purpose of these asserts are to test the device side
30-
// half API for the common mathematical functions.
31-
// Note: When calling std math functions from device, don't
32-
// use the std namespace, but just "::" so that the function
33-
// gets resolved from nvcc math_functions.hpp
34-
35-
float threshold = 0.00001;
36-
assert(::abs(::lgamma(Half(10.0)) - ::lgamma(10.0f)) <= threshold);
37-
assert(::abs(::exp(Half(1.0)) - ::exp(1.0f)) <= threshold);
38-
assert(::abs(::log(Half(1.0)) - ::log(1.0f)) <= threshold);
39-
assert(::abs(::log10(Half(1000.0)) - ::log10(1000.0f)) <= threshold);
40-
assert(::abs(::log1p(Half(0.0)) - ::log1p(0.0f)) <= threshold);
41-
assert(::abs(::log2(Half(1000.0)) - ::log2(1000.0f)) <= threshold);
42-
assert(::abs(::expm1(Half(1.0)) - ::expm1(1.0f)) <= threshold);
43-
assert(::abs(::cos(Half(0.0)) - ::cos(0.0f)) <= threshold);
44-
assert(::abs(::sin(Half(0.0)) - ::sin(0.0f)) <= threshold);
45-
assert(::abs(::sqrt(Half(100.0)) - ::sqrt(100.0f)) <= threshold);
46-
assert(::abs(::ceil(Half(2.4)) - ::ceil(2.4f)) <= threshold);
47-
assert(::abs(::floor(Half(2.7)) - ::floor(2.7f)) <= threshold);
48-
assert(::abs(::trunc(Half(2.7)) - ::trunc(2.7f)) <= threshold);
49-
assert(::abs(::acos(Half(-1.0)) - ::acos(-1.0f)) <= threshold);
50-
assert(::abs(::cosh(Half(1.0)) - ::cosh(1.0f)) <= threshold);
51-
assert(::abs(::acosh(Half(1.0)) - ::acosh(1.0f)) <= threshold);
52-
assert(::abs(::acosh(Half(1.0)) - ::acosh(1.0f)) <= threshold);
53-
assert(::abs(::asinh(Half(1.0)) - ::asinh(1.0f)) <= threshold);
54-
assert(::abs(::atanh(Half(0.5)) - ::atanh(0.5f)) <= threshold);
55-
assert(::abs(::asin(Half(1.0)) - ::asin(1.0f)) <= threshold);
56-
assert(::abs(::sinh(Half(1.0)) - ::sinh(1.0f)) <= threshold);
57-
assert(::abs(::asinh(Half(1.0)) - ::asinh(1.0f)) <= threshold);
58-
assert(::abs(::tan(Half(0.0)) - ::tan(0.0f)) <= threshold);
59-
assert(::abs(::atan(Half(1.0)) - ::atan(1.0f)) <= threshold);
60-
assert(::abs(::tanh(Half(1.0)) - ::tanh(1.0f)) <= threshold);
61-
assert(::abs(::erf(Half(10.0)) - ::erf(10.0f)) <= threshold);
62-
assert(::abs(::erfc(Half(10.0)) - ::erfc(10.0f)) <= threshold);
63-
assert(::abs(::abs(Half(-3.0)) - ::abs(-3.0f)) <= threshold);
64-
assert(::abs(::round(Half(2.3)) - ::round(2.3f)) <= threshold);
65-
assert(::abs(::pow(Half(2.0), Half(10.0)) - ::pow(2.0f, 10.0f)) <= threshold);
66-
assert(
67-
::abs(::atan2(Half(7.0), Half(0.0)) - ::atan2(7.0f, 0.0f)) <= threshold);
68-
// note: can't use namespace on isnan and isinf in device code
69-
#ifdef _MSC_VER
70-
// Windows requires this explicit conversion. The reason is unclear
71-
// related issue with clang: https://reviews.llvm.org/D37906
72-
assert(::abs(::isnan((float)Half(0.0)) - ::isnan(0.0f)) <= threshold);
73-
assert(::abs(::isinf((float)Half(0.0)) - ::isinf(0.0f)) <= threshold);
74-
#else
75-
assert(::abs(::isnan(Half(0.0)) - ::isnan(0.0f)) <= threshold);
76-
assert(::abs(::isinf(Half(0.0)) - ::isinf(0.0f)) <= threshold);
77-
#endif
78-
}
79-
80-
__global__ void kernel(){
81-
test();
8222
}
8323

8424
void launch_function(){

‎aten/src/ATen/test/cuda_lowfp_test.cuh

Lines changed: 63 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,63 @@
1+
#pragma once
2+
3+
template<template T>
4+
__device__ void test(){
5+
// test half construction and implicit conversions in device
6+
assert(T(3) == T(3.0f));
7+
assert(static_cast<T>(3.0f) == T(3.0f));
8+
// there is no float <=> __half/__nv_bfloat16 implicit conversion
9+
assert(static_cast<T>(3.0f) == 3.0f);
10+
11+
// asserting if the functions used on
12+
// half types give almost equivalent results when using
13+
// functions on double.
14+
// The purpose of these asserts are to test the device side
15+
// half API for the common mathematical functions.
16+
// Note: When calling std math functions from device, don't
17+
// use the std namespace, but just "::" so that the function
18+
// gets resolved from nvcc math_functions.hpp
19+
20+
float threshold = 0.00001;
21+
assert(::abs(::lgamma(T(10.0)) - ::lgamma(10.0f)) <= threshold);
22+
assert(::abs(::exp(T(1.0)) - ::exp(1.0f)) <= threshold);
23+
assert(::abs(::log(T(1.0)) - ::log(1.0f)) <= threshold);
24+
assert(::abs(::log10(T(1000.0)) - ::log10(1000.0f)) <= threshold);
25+
assert(::abs(::log1p(T(0.0)) - ::log1p(0.0f)) <= threshold);
26+
assert(::abs(::log2(T(1000.0)) - ::log2(1000.0f)) <= threshold);
27+
assert(::abs(::expm1(T(1.0)) - ::expm1(1.0f)) <= threshold);
28+
assert(::abs(::cos(T(0.0)) - ::cos(0.0f)) <= threshold);
29+
assert(::abs(::sin(T(0.0)) - ::sin(0.0f)) <= threshold);
30+
assert(::abs(::sqrt(T(100.0)) - ::sqrt(100.0f)) <= threshold);
31+
assert(::abs(::ceil(T(2.4)) - ::ceil(2.4f)) <= threshold);
32+
assert(::abs(::floor(T(2.7)) - ::floor(2.7f)) <= threshold);
33+
assert(::abs(::trunc(T(2.7)) - ::trunc(2.7f)) <= threshold);
34+
assert(::abs(::acos(T(-1.0)) - ::acos(-1.0f)) <= threshold);
35+
assert(::abs(::cosh(T(1.0)) - ::cosh(1.0f)) <= threshold);
36+
assert(::abs(::acosh(T(1.0)) - ::acosh(1.0f)) <= threshold);
37+
assert(::abs(::acosh(T(1.0)) - ::acosh(1.0f)) <= threshold);
38+
assert(::abs(::asinh(T(1.0)) - ::asinh(1.0f)) <= threshold);
39+
assert(::abs(::atanh(T(0.5)) - ::atanh(0.5f)) <= threshold);
40+
assert(::abs(::asin(T(1.0)) - ::asin(1.0f)) <= threshold);
41+
assert(::abs(::sinh(T(1.0)) - ::sinh(1.0f)) <= threshold);
42+
assert(::abs(::asinh(T(1.0)) - ::asinh(1.0f)) <= threshold);
43+
assert(::abs(::tan(T(0.0)) - ::tan(0.0f)) <= threshold);
44+
assert(::abs(::atan(T(1.0)) - ::atan(1.0f)) <= threshold);
45+
assert(::abs(::tanh(T(1.0)) - ::tanh(1.0f)) <= threshold);
46+
assert(::abs(::erf(T(10.0)) - ::erf(10.0f)) <= threshold);
47+
assert(::abs(::erfc(T(10.0)) - ::erfc(10.0f)) <= threshold);
48+
assert(::abs(::abs(T(-3.0)) - ::abs(-3.0f)) <= threshold);
49+
assert(::abs(::round(T(2.3)) - ::round(2.3f)) <= threshold);
50+
assert(::abs(::pow(T(2.0), T(10.0)) - ::pow(2.0f, 10.0f)) <= threshold);
51+
assert(
52+
::abs(::atan2(T(7.0), T(0.0)) - ::atan2(7.0f, 0.0f)) <= threshold);
53+
// note: can't use namespace on isnan and isinf in device code
54+
#ifdef _MSC_VER
55+
// Windows requires this explicit conversion. The reason is unclear
56+
// related issue with clang: https://reviews.llvm.org/D37906
57+
assert(::abs(::isnan((float)T(0.0)) - ::isnan(0.0f)) <= threshold);
58+
assert(::abs(::isinf((float)T(0.0)) - ::isinf(0.0f)) <= threshold);
59+
#else
60+
assert(::abs(::isnan(T(0.0)) - ::isnan(0.0f)) <= threshold);
61+
assert(::abs(::isinf(T(0.0)) - ::isinf(0.0f)) <= threshold);
62+
#endif
63+
}

0 commit comments

Comments
 (0)