test_gelu.cu 4.18 KB
Newer Older
Przemek Tredak's avatar
Przemek Tredak committed
1
/*************************************************************************
2
 * Copyright (c) 2022-2023, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
Przemek Tredak's avatar
Przemek Tredak committed
3
4
5
6
7
 *
 * See LICENSE for license information.
 ************************************************************************/

#include <cmath>
Tim Moon's avatar
Tim Moon committed
8
#include <cstring>
Przemek Tredak's avatar
Przemek Tredak committed
9
10
#include <memory>
#include <iomanip>
Tim Moon's avatar
Tim Moon committed
11
#include <iostream>
Przemek Tredak's avatar
Przemek Tredak committed
12
13
#include <random>
#include <type_traits>
Tim Moon's avatar
Tim Moon committed
14
15
16
17
18
19

#include <cuda_bf16.h>
#include <cuda_runtime.h>
#include <gtest/gtest.h>

#include <transformer_engine/activation.h>
Przemek Tredak's avatar
Przemek Tredak committed
20
21
22
23
24
25
26
#include "../test_common.h"

using namespace transformer_engine;

template <typename IT, typename OT, typename CT>
void compute_ref_gelu_cast(const IT *input_h,
                           OT *output_h,
27
                           const CT scale,
Przemek Tredak's avatar
Przemek Tredak committed
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
                           CT *amax_h,
                           const size_t N,
                           const size_t H) {
  CT amax  = 0.;

  for (size_t i = 0; i < N; i++) {
    for (size_t j = 0; j < H; j++) {
      CT elt = CT(input_h[i * H + j]);
      elt = 0.5f * elt * (1.0f + tanhf(0.79788456F * elt *
                                       (1.0f + 0.044715f * elt * elt)));
      output_h[i * H + j] = OT(scale * elt);
      amax = std::abs(elt) > amax ? std::abs(elt) : amax;
    }
  }

  *amax_h = amax;
}

template <typename IType, typename OType>
void performTestGelu(const size_t N, const size_t H) {
  using namespace test;

  DType itype = TypeInfo<IType>::dtype;
  DType otype = TypeInfo<OType>::dtype;

  Tensor input({ N, H }, itype);
  Tensor output({ N, H }, otype);

56
57
  fillUniform(&input);
  setRandomScale(&output);
Przemek Tredak's avatar
Przemek Tredak committed
58
59
60

  std::unique_ptr<OType[]> ref_output = std::make_unique<OType[]>(N*H);

61
  nvte_gelu(input.data(), output.data(), 0);
Przemek Tredak's avatar
Przemek Tredak committed
62
63
64

  float ref_amax;
  compute_ref_gelu_cast(input.cpu_dptr<IType>(), ref_output.get(),
65
                        output.scale(), &ref_amax, N, H);
Przemek Tredak's avatar
Przemek Tredak committed
66
67
68
69
70
71
72

  cudaDeviceSynchronize();
  auto err = cudaGetLastError();
  ASSERT_EQ(err, cudaSuccess) << cudaGetErrorString(err);

  if (otype == DType::kFloat8E4M3 || otype == DType::kFloat8E5M2) {
    auto [atol_amax, rtol_amax] = getTolerances(DType::kFloat32);
73
    compareResults("amax", output.amax(), ref_amax, atol_amax, rtol_amax);
Przemek Tredak's avatar
Przemek Tredak committed
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
  }
  auto [atol, rtol] = getTolerances(otype);
  compareResults("output_gelu", output, ref_output.get(), atol, rtol);
}

class GELUTestSuite : public ::testing::TestWithParam<std::tuple<transformer_engine::DType,
                                                                 transformer_engine::DType,
                                                                 std::pair<size_t, size_t>>> {};

TEST_P(GELUTestSuite, TestGELU) {
    using namespace transformer_engine;
    using namespace test;

    const DType input_type = std::get<0>(GetParam());
    const DType output_type = std::get<1>(GetParam());
    const auto size = std::get<2>(GetParam());

    TRANSFORMER_ENGINE_TYPE_SWITCH_ALL(input_type, InputType,
      TRANSFORMER_ENGINE_TYPE_SWITCH_ALL(output_type, OutputType,
        performTestGelu<InputType, OutputType>(size.first, size.second);
      );
    );
}

namespace {

std::vector<std::pair<size_t, size_t>> gelu_test_cases = {{2048, 12288},
                                                          {768, 1024},
                                                          {256, 65536},
                                                          {65536, 128},
                                                          {256, 256},
                                                          {257, 259},
                                                          {128, 128+1}};

}  // namespace

INSTANTIATE_TEST_SUITE_P(
    OperatorTest,
    GELUTestSuite,
    ::testing::Combine(
        ::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
        ::testing::ValuesIn(test::all_fp_types),
        ::testing::ValuesIn(gelu_test_cases)),
    [](const testing::TestParamInfo<GELUTestSuite::ParamType>& info) {
      std::string name = test::typeName(std::get<0>(info.param)) + "X" +
                         test::typeName(std::get<1>(info.param)) + "X" +
                         std::to_string(std::get<2>(info.param).first) + "X" +
                         std::to_string(std::get<2>(info.param).second);
      return name;
    });