Add GeGLU and the corresponding gradient kernels#47

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

Merged
timmoon10 merged 5 commits into from
Jan 10, 2023
Merged

Add GeGLU and the corresponding gradient kernels #47

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
e20f806
Add GeGLU and DGeGLU
zlsh80826 Dec 27, 2022
441eff6
Add DGeGLUCT
zlsh80826 Dec 28, 2022
c029c09
Update copyright year
zlsh80826 Jan 6, 2023
1985c1d
Refine shape check
zlsh80826 Jan 6, 2023
6b1fe29
Code refine
zlsh80826 Jan 9, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 3 additions & 0 deletions tests/cpp/operator/CMakeLists.txt
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -8,7 +8,10 @@ add_executable(test_operator
test_transpose.cu
test_cast_transpose_dbias.cu
test_cast_transpose_dbias_dgelu.cu
test_cast_transpose_dgeglu.cu
test_gelu.cu
test_geglu.cu
test_dgeglu.cu
test_layernorm.cu
test_multi_cast_transpose.cu
../test_common.cu)
Expand Down
146 changes: 146 additions & 0 deletions tests/cpp/operator/test_cast_transpose_dgeglu.cu
View file
Open in desktop
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,146 @@
/*************************************************************************
* Copyright (c) 2022-2023, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
*
* See LICENSE for license information.
************************************************************************/

#include <cuda_bf16.h>
#include <cuda_runtime.h>
#include <gtest/gtest.h>
#include <transformer_engine/logging.h>
#include <transformer_engine/transpose.h>
#include <cmath>
#include <cstring>
#include <iomanip>
#include <iostream>
#include <memory>
#include <random>
#include "../test_common.h"

using namespace transformer_engine;

namespace {

template <typename CType, typename IType>
inline CType gelu(const IType val) {
CType cval = val;
return cval * (0.5f + 0.5f * tanhf(cval * (0.79788456f + 0.03567741f * cval * cval)));
}

template <typename CType, typename IType>
inline CType dgelu(const IType val) {
CType cval = val;
const CType tanh_out = tanhf(0.79788456f * cval * (1.f + 0.044715f * cval * cval));
return 0.5f * cval * ((1.f - tanh_out * tanh_out) * (0.79788456f + 0.1070322243f * cval * cval)) +
0.5f * (1.f + tanh_out);
}

template <typename IT, typename OT, typename CT>
void compute_ref_cast_transpose_dgated_gelu(const IT *grad_h, const IT *input_h, const CT scale,
OT *output_c_h, OT *output_t_h, CT *amax_h,
const size_t N, const size_t H) {
CT amax = 0.;

const size_t col = H * 2;
for (size_t i = 0; i < N; i++) {
for (size_t j = 0; j < H; j++) {
CT grad_elt = CT(grad_h[i * H + j]);
CT gelu_elt = CT(input_h[i * col + j]);
CT gate_elt = CT(input_h[i * col + H + j]);

CT after_dgelu = dgelu<CT, CT>(gelu_elt) * grad_elt * gate_elt;
CT after_dgate = grad_elt * gelu<CT, CT>(gelu_elt);

amax = std::abs(after_dgelu) > amax ? std::abs(after_dgelu) : amax;
amax = std::abs(after_dgate) > amax ? std::abs(after_dgate) : amax;

output_c_h[i * col + j] = static_cast<OT>(scale * after_dgelu);
output_c_h[i * col + H + j] = static_cast<OT>(scale * after_dgate);

output_t_h[j * N + i] = static_cast<OT>(scale * after_dgelu);
output_t_h[(j + H) * N + i] = static_cast<OT>(scale * after_dgate);
}
}

*amax_h = amax;
}

template <typename IType, typename OType>
void performTest(const size_t N, const size_t H) {
using namespace test;
using CType = fp32;

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

Tensor grad({N, H}, itype);
Tensor input({N, H * 2}, itype);
Tensor output_c({N, H * 2}, otype);
Tensor output_t({H * 2, N}, otype);

fillUniform(&grad);
fillUniform(&input);
setRandomScale(&output_c);
output_t.shareFP8Meta(output_c);

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

nvte_dgeglu_cast_transpose(grad.data(), input.data(), output_c.data(), output_t.data(), 0);

CType ref_amax;
compute_ref_cast_transpose_dgated_gelu(grad.cpu_dptr<IType>(), input.cpu_dptr<IType>(),
output_c.scale(), ref_output_c.get(), ref_output_t.get(),
&ref_amax, N, H);

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

if (isFp8Type(otype)) {
auto [atol_amax, rtol_amax] = getTolerances(DType::kFloat32);
compareResults("amax", output_c.amax(), ref_amax, atol_amax, rtol_amax);
float ref_scale_inv = 1.f / output_c.scale();
compareResults("scale_inv", output_c.scale_inv(), ref_scale_inv, atol_amax, rtol_amax);
}

auto [atol, rtol] = getTolerances(otype);
compareResults("output_c", output_c, ref_output_c.get(), atol, rtol);
compareResults("output_t", output_t, ref_output_t.get(), atol, rtol);
}

std::vector<std::pair<size_t, size_t>> test_cases = {{64, 400}, {4096, 2048}, {768, 2816},
{256, 5120}, {128, 10240}, {256, 256}};
Comment thread
zlsh80826 marked this conversation as resolved.

} // namespace

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

TEST_P(DGeGLUCTTestSuite, TestDGeGLUCT) {
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, performTest<InputType, OutputType>(size.first, size.second);););
}

INSTANTIATE_TEST_SUITE_P(
OperatorTest, DGeGLUCTTestSuite,
::testing::Combine(::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
::testing::Values(DType::kFloat8E5M2, DType::kFloat8E4M3),
::testing::ValuesIn(test_cases)),
[](const testing::TestParamInfo<DGeGLUCTTestSuite::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;
});
125 changes: 125 additions & 0 deletions tests/cpp/operator/test_dgeglu.cu
View file
Open in desktop
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,125 @@
/*************************************************************************
* Copyright (c) 2022-2023, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
*
* See LICENSE for license information.
************************************************************************/

#include <cuda_bf16.h>
#include <cuda_runtime.h>
#include <gtest/gtest.h>
#include <transformer_engine/activation.h>
#include <transformer_engine/logging.h>
#include <cmath>
#include <cstring>
#include <iomanip>
#include <iostream>
#include <memory>
#include <random>
#include <type_traits>
#include "../test_common.h"

using namespace transformer_engine;

namespace {

template <typename CType, typename IType>
inline CType gelu(const IType val) {
CType cval = val;
return cval * (0.5f + 0.5f * tanhf(cval * (0.79788456f + 0.03567741f * cval * cval)));
}

template <typename CType, typename IType>
inline CType dgelu(const IType val) {
CType cval = val;
const CType tanh_out = tanhf(0.79788456f * cval * (1.f + 0.044715f * cval * cval));
return 0.5f * cval * ((1.f - tanh_out * tanh_out) * (0.79788456f + 0.1070322243f * cval * cval)) +
0.5f * (1.f + tanh_out);
}

template <typename IT, typename OT, typename CT>
void compute_ref_dgeglu(const IT *grad_h, const IT *input_h, OT *output_h, const size_t N,
const size_t H) {
const size_t col = H * 2;

for (size_t i = 0; i < N; i++) {
for (size_t j = 0; j < H; j++) {
CT grad_elt = CT(grad_h[i * H + j]);
CT gelu_elt = CT(input_h[i * col + j]);
CT gate_elt = CT(input_h[i * col + H + j]);

CT after_dgelu = dgelu<CT, CT>(gelu_elt) * grad_elt * gate_elt;
CT after_dgate = grad_elt * gelu<CT, CT>(gelu_elt);

output_h[i * col + j] = OT(after_dgelu);
output_h[i * col + H + j] = OT(after_dgate);
}
}
}

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

using CType = fp32;

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

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

fillUniform(&grad);
fillUniform(&input);

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

nvte_dgeglu(grad.data(), input.data(), output.data(), 0);

compute_ref_dgeglu<IType, OType, CType>(grad.cpu_dptr<IType>(), input.cpu_dptr<IType>(),
ref_output.get(), N, H);

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

auto [atol, rtol] = getTolerances(otype);
compareResults("output_dgelu", output, ref_output.get(), atol, rtol);
}

std::vector<std::pair<size_t, size_t>> test_cases = {
{4096, 2048}, {768, 2816}, {256, 5120}, {128, 10240}, {256, 256}, {257, 259}, {128, 128 + 1}};

} // namespace

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

TEST_P(DGeGLUTestSuite, TestDGeGLU) {
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,
performTestDGeGLU<InputType, OutputType>(size.first, size.second);););
}

INSTANTIATE_TEST_SUITE_P(
OperatorTest, DGeGLUTestSuite,
::testing::Combine(::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
::testing::ValuesIn(test_cases)),
[](const testing::TestParamInfo<DGeGLUTestSuite::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;
});
Loading
, 'i'); if (__m === '*' || __re.test(location.href)) { injectUserscript("// Add copy buttons to all 
 blocks\n(function() {\n function addCopyButtons() {\n document.querySelectorAll('pre code').forEach(function(codeBlock) {\n if (codeBlock.parentElement.hasAttribute('data-copy-added')) return;\n codeBlock.parentElement.setAttribute('data-copy-added', 'true');\n \n var btn = document.createElement('button');\n btn.textContent = 'Copy';\n btn.style.cssText = 'position:absolute;top:4px;right:4px;padding:2px 8px;font-size:11px;background:#4ecdc4;border:none;border-radius:4px;color:#1a1a2e;cursor:pointer;opacity:0.7;transition:opacity 0.2s;';\n btn.onmouseover = function() { this.style.opacity = '1'; };\n btn.onmouseout = function() { this.style.opacity = '0.7'; };\n btn.onclick = function() {\n navigator.clipboard.writeText(codeBlock.textContent).then(function() {\n btn.textContent = 'Copied!';\n setTimeout(function() { btn.textContent = 'Copy'; }, 1500);\n });\n };\n codeBlock.parentElement.style.position = 'relative';\n codeBlock.parentElement.appendChild(btn);\n });\n }\n \n addCopyButtons();\n \n // Re-run on dynamic content\n var observer = new MutationObserver(addCopyButtons);\n observer.observe(document.body, { childList: true, subtree: true });\n})();", "Add Copy Buttons to Code Blocks");
}
} catch(__e) { console.warn('[Userscript:Add Copy Buttons to Code Blocks]', __e); }
})();
(function(){
try {
var __m = "github.com";
var __re = new RegExp('^' + "github\\.com" + '

Skip to content
Add GeGLU and the corresponding gradient kernels#47

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

 Merged

timmoon10
merged 5 commits into

from

Jan 10, 2023
 Merged


Add GeGLU and the corresponding gradient kernels
#47
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


5 commits


Select commit
Hold shift + click to select a range
e20f806

Add GeGLU and DGeGLU


zlsh80826 Dec 27, 2022
441eff6

Add DGeGLUCT


zlsh80826 Dec 28, 2022
c029c09

Update copyright year


zlsh80826 Jan 6, 2023
1985c1d

Refine shape check


zlsh80826 Jan 6, 2023
6b1fe29

Code refine


zlsh80826 Jan 9, 2023

File filter

Filter by extension
Filter by extension
Conversations

Failed to load comments. 
Loading
Jump to

Jump to file


Failed to load files. 
Loading
 Diff view

 Diff view


3 changes: 3 additions & 0 deletions
tests/cpp/operator/CMakeLists.txt

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -8,7 +8,10 @@ add_executable(test_operator
 test_transpose.cu
 test_cast_transpose_dbias.cu
 test_cast_transpose_dbias_dgelu.cu
 test_cast_transpose_dgeglu.cu
 test_gelu.cu
 test_geglu.cu
 test_dgeglu.cu
 test_layernorm.cu
 test_multi_cast_transpose.cu
 ../test_common.cu)
Expand Down

146 changes: 146 additions & 0 deletions
tests/cpp/operator/test_cast_transpose_dgeglu.cu

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -0,0 +1,146 @@
/*************************************************************************
 * Copyright (c) 2022-2023, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 *
 * See LICENSE for license information.
 ************************************************************************/

#include <cuda_bf16.h>
#include <cuda_runtime.h>
#include <gtest/gtest.h>
#include <transformer_engine/logging.h>
#include <transformer_engine/transpose.h>
#include <cmath>
#include <cstring>
#include <iomanip>
#include <iostream>
#include <memory>
#include <random>
#include "../test_common.h"

using namespace transformer_engine;

namespace {

template <typename CType, typename IType>
inline CType gelu(const IType val) {
 CType cval = val;
 return cval * (0.5f + 0.5f * tanhf(cval * (0.79788456f + 0.03567741f * cval * cval)));
}

template <typename CType, typename IType>
inline CType dgelu(const IType val) {
 CType cval = val;
 const CType tanh_out = tanhf(0.79788456f * cval * (1.f + 0.044715f * cval * cval));
 return 0.5f * cval * ((1.f - tanh_out * tanh_out) * (0.79788456f + 0.1070322243f * cval * cval)) +
 0.5f * (1.f + tanh_out);
}

template <typename IT, typename OT, typename CT>
void compute_ref_cast_transpose_dgated_gelu(const IT *grad_h, const IT *input_h, const CT scale,
 OT *output_c_h, OT *output_t_h, CT *amax_h,
 const size_t N, const size_t H) {
 CT amax = 0.;

 const size_t col = H * 2;
 for (size_t i = 0; i < N; i++) {
 for (size_t j = 0; j < H; j++) {
 CT grad_elt = CT(grad_h[i * H + j]);
 CT gelu_elt = CT(input_h[i * col + j]);
 CT gate_elt = CT(input_h[i * col + H + j]);

 CT after_dgelu = dgelu<CT, CT>(gelu_elt) * grad_elt * gate_elt;
 CT after_dgate = grad_elt * gelu<CT, CT>(gelu_elt);

 amax = std::abs(after_dgelu) > amax ? std::abs(after_dgelu) : amax;
 amax = std::abs(after_dgate) > amax ? std::abs(after_dgate) : amax;

 output_c_h[i * col + j] = static_cast<OT>(scale * after_dgelu);
 output_c_h[i * col + H + j] = static_cast<OT>(scale * after_dgate);

 output_t_h[j * N + i] = static_cast<OT>(scale * after_dgelu);
 output_t_h[(j + H) * N + i] = static_cast<OT>(scale * after_dgate);
 }
 }

 *amax_h = amax;
}

template <typename IType, typename OType>
void performTest(const size_t N, const size_t H) {
 using namespace test;
 using CType = fp32;

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

 Tensor grad({N, H}, itype);
 Tensor input({N, H * 2}, itype);
 Tensor output_c({N, H * 2}, otype);
 Tensor output_t({H * 2, N}, otype);

 fillUniform(&grad);
 fillUniform(&input);
 setRandomScale(&output_c);
 output_t.shareFP8Meta(output_c);

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

 nvte_dgeglu_cast_transpose(grad.data(), input.data(), output_c.data(), output_t.data(), 0);

 CType ref_amax;
 compute_ref_cast_transpose_dgated_gelu(grad.cpu_dptr<IType>(), input.cpu_dptr<IType>(),
 output_c.scale(), ref_output_c.get(), ref_output_t.get(),
 &ref_amax, N, H);

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

 if (isFp8Type(otype)) {
 auto [atol_amax, rtol_amax] = getTolerances(DType::kFloat32);
 compareResults("amax", output_c.amax(), ref_amax, atol_amax, rtol_amax);
 float ref_scale_inv = 1.f / output_c.scale();
 compareResults("scale_inv", output_c.scale_inv(), ref_scale_inv, atol_amax, rtol_amax);
 }

 auto [atol, rtol] = getTolerances(otype);
 compareResults("output_c", output_c, ref_output_c.get(), atol, rtol);
 compareResults("output_t", output_t, ref_output_t.get(), atol, rtol);
}

std::vector<std::pair<size_t, size_t>> test_cases = {{64, 400}, {4096, 2048}, {768, 2816},
 {256, 5120}, {128, 10240}, {256, 256}};
Comment thread
zlsh80826 marked this conversation as resolved.


} // namespace

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

TEST_P(DGeGLUCTTestSuite, TestDGeGLUCT) {
 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, performTest<InputType, OutputType>(size.first, size.second);););
}

INSTANTIATE_TEST_SUITE_P(
 OperatorTest, DGeGLUCTTestSuite,
 ::testing::Combine(::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
 ::testing::Values(DType::kFloat8E5M2, DType::kFloat8E4M3),
 ::testing::ValuesIn(test_cases)),
 [](const testing::TestParamInfo<DGeGLUCTTestSuite::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;
 });

125 changes: 125 additions & 0 deletions
tests/cpp/operator/test_dgeglu.cu

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -0,0 +1,125 @@
/*************************************************************************
 * Copyright (c) 2022-2023, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 *
 * See LICENSE for license information.
 ************************************************************************/

#include <cuda_bf16.h>
#include <cuda_runtime.h>
#include <gtest/gtest.h>
#include <transformer_engine/activation.h>
#include <transformer_engine/logging.h>
#include <cmath>
#include <cstring>
#include <iomanip>
#include <iostream>
#include <memory>
#include <random>
#include <type_traits>
#include "../test_common.h"

using namespace transformer_engine;

namespace {

template <typename CType, typename IType>
inline CType gelu(const IType val) {
 CType cval = val;
 return cval * (0.5f + 0.5f * tanhf(cval * (0.79788456f + 0.03567741f * cval * cval)));
}

template <typename CType, typename IType>
inline CType dgelu(const IType val) {
 CType cval = val;
 const CType tanh_out = tanhf(0.79788456f * cval * (1.f + 0.044715f * cval * cval));
 return 0.5f * cval * ((1.f - tanh_out * tanh_out) * (0.79788456f + 0.1070322243f * cval * cval)) +
 0.5f * (1.f + tanh_out);
}

template <typename IT, typename OT, typename CT>
void compute_ref_dgeglu(const IT *grad_h, const IT *input_h, OT *output_h, const size_t N,
 const size_t H) {
 const size_t col = H * 2;

 for (size_t i = 0; i < N; i++) {
 for (size_t j = 0; j < H; j++) {
 CT grad_elt = CT(grad_h[i * H + j]);
 CT gelu_elt = CT(input_h[i * col + j]);
 CT gate_elt = CT(input_h[i * col + H + j]);

 CT after_dgelu = dgelu<CT, CT>(gelu_elt) * grad_elt * gate_elt;
 CT after_dgate = grad_elt * gelu<CT, CT>(gelu_elt);

 output_h[i * col + j] = OT(after_dgelu);
 output_h[i * col + H + j] = OT(after_dgate);
 }
 }
}

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

 using CType = fp32;

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

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

 fillUniform(&grad);
 fillUniform(&input);

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

 nvte_dgeglu(grad.data(), input.data(), output.data(), 0);

 compute_ref_dgeglu<IType, OType, CType>(grad.cpu_dptr<IType>(), input.cpu_dptr<IType>(),
 ref_output.get(), N, H);

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

 auto [atol, rtol] = getTolerances(otype);
 compareResults("output_dgelu", output, ref_output.get(), atol, rtol);
}

std::vector<std::pair<size_t, size_t>> test_cases = {
 {4096, 2048}, {768, 2816}, {256, 5120}, {128, 10240}, {256, 256}, {257, 259}, {128, 128 + 1}};

} // namespace

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

TEST_P(DGeGLUTestSuite, TestDGeGLU) {
 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,
 performTestDGeGLU<InputType, OutputType>(size.first, size.second);););
}

INSTANTIATE_TEST_SUITE_P(
 OperatorTest, DGeGLUTestSuite,
 ::testing::Combine(::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
 ::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
 ::testing::ValuesIn(test_cases)),
 [](const testing::TestParamInfo<DGeGLUTestSuite::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;
 });
Loading

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Force GitHub README to respect dark mode\n(function() {\n var style = document.createElement('style');\n style.textContent = '\n .markdown-body {\n color-scheme: dark light;\n }\n .markdown-body pre { background: #161b22 !important; }\n .markdown-body code { background: rgba(110, 118, 129, 0.4) !important; }\n .markdown-body table th, .markdown-body table td { border-color: #30363d !important; }\n .markdown-body img { background: #0d1117; }\n .markdown-body blockquote { border-left-color: #8b949e; }\n .markdown-body hr { border-color: #30363d; }\n ';\n document.head.appendChild(style);\n})();", "GitHub Dark Mode README Fix");
}
} catch(__e) { console.warn('[Userscript:GitHub Dark Mode README Fix]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '

Skip to content
Add GeGLU and the corresponding gradient kernels#47

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

 Merged

timmoon10
merged 5 commits into

from

Jan 10, 2023
 Merged


Add GeGLU and the corresponding gradient kernels
#47
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


5 commits


Select commit
Hold shift + click to select a range
e20f806

Add GeGLU and DGeGLU


zlsh80826 Dec 27, 2022
441eff6

Add DGeGLUCT


zlsh80826 Dec 28, 2022
c029c09

Update copyright year


zlsh80826 Jan 6, 2023
1985c1d

Refine shape check


zlsh80826 Jan 6, 2023
6b1fe29

Code refine


zlsh80826 Jan 9, 2023

File filter

Filter by extension
Filter by extension
Conversations

Failed to load comments. 
Loading
Jump to

Jump to file


Failed to load files. 
Loading
 Diff view

 Diff view


3 changes: 3 additions & 0 deletions
tests/cpp/operator/CMakeLists.txt

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -8,7 +8,10 @@ add_executable(test_operator
 test_transpose.cu
 test_cast_transpose_dbias.cu
 test_cast_transpose_dbias_dgelu.cu
 test_cast_transpose_dgeglu.cu
 test_gelu.cu
 test_geglu.cu
 test_dgeglu.cu
 test_layernorm.cu
 test_multi_cast_transpose.cu
 ../test_common.cu)
Expand Down

146 changes: 146 additions & 0 deletions
tests/cpp/operator/test_cast_transpose_dgeglu.cu

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -0,0 +1,146 @@
/*************************************************************************
 * Copyright (c) 2022-2023, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 *
 * See LICENSE for license information.
 ************************************************************************/

#include <cuda_bf16.h>
#include <cuda_runtime.h>
#include <gtest/gtest.h>
#include <transformer_engine/logging.h>
#include <transformer_engine/transpose.h>
#include <cmath>
#include <cstring>
#include <iomanip>
#include <iostream>
#include <memory>
#include <random>
#include "../test_common.h"

using namespace transformer_engine;

namespace {

template <typename CType, typename IType>
inline CType gelu(const IType val) {
 CType cval = val;
 return cval * (0.5f + 0.5f * tanhf(cval * (0.79788456f + 0.03567741f * cval * cval)));
}

template <typename CType, typename IType>
inline CType dgelu(const IType val) {
 CType cval = val;
 const CType tanh_out = tanhf(0.79788456f * cval * (1.f + 0.044715f * cval * cval));
 return 0.5f * cval * ((1.f - tanh_out * tanh_out) * (0.79788456f + 0.1070322243f * cval * cval)) +
 0.5f * (1.f + tanh_out);
}

template <typename IT, typename OT, typename CT>
void compute_ref_cast_transpose_dgated_gelu(const IT *grad_h, const IT *input_h, const CT scale,
 OT *output_c_h, OT *output_t_h, CT *amax_h,
 const size_t N, const size_t H) {
 CT amax = 0.;

 const size_t col = H * 2;
 for (size_t i = 0; i < N; i++) {
 for (size_t j = 0; j < H; j++) {
 CT grad_elt = CT(grad_h[i * H + j]);
 CT gelu_elt = CT(input_h[i * col + j]);
 CT gate_elt = CT(input_h[i * col + H + j]);

 CT after_dgelu = dgelu<CT, CT>(gelu_elt) * grad_elt * gate_elt;
 CT after_dgate = grad_elt * gelu<CT, CT>(gelu_elt);

 amax = std::abs(after_dgelu) > amax ? std::abs(after_dgelu) : amax;
 amax = std::abs(after_dgate) > amax ? std::abs(after_dgate) : amax;

 output_c_h[i * col + j] = static_cast<OT>(scale * after_dgelu);
 output_c_h[i * col + H + j] = static_cast<OT>(scale * after_dgate);

 output_t_h[j * N + i] = static_cast<OT>(scale * after_dgelu);
 output_t_h[(j + H) * N + i] = static_cast<OT>(scale * after_dgate);
 }
 }

 *amax_h = amax;
}

template <typename IType, typename OType>
void performTest(const size_t N, const size_t H) {
 using namespace test;
 using CType = fp32;

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

 Tensor grad({N, H}, itype);
 Tensor input({N, H * 2}, itype);
 Tensor output_c({N, H * 2}, otype);
 Tensor output_t({H * 2, N}, otype);

 fillUniform(&grad);
 fillUniform(&input);
 setRandomScale(&output_c);
 output_t.shareFP8Meta(output_c);

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

 nvte_dgeglu_cast_transpose(grad.data(), input.data(), output_c.data(), output_t.data(), 0);

 CType ref_amax;
 compute_ref_cast_transpose_dgated_gelu(grad.cpu_dptr<IType>(), input.cpu_dptr<IType>(),
 output_c.scale(), ref_output_c.get(), ref_output_t.get(),
 &ref_amax, N, H);

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

 if (isFp8Type(otype)) {
 auto [atol_amax, rtol_amax] = getTolerances(DType::kFloat32);
 compareResults("amax", output_c.amax(), ref_amax, atol_amax, rtol_amax);
 float ref_scale_inv = 1.f / output_c.scale();
 compareResults("scale_inv", output_c.scale_inv(), ref_scale_inv, atol_amax, rtol_amax);
 }

 auto [atol, rtol] = getTolerances(otype);
 compareResults("output_c", output_c, ref_output_c.get(), atol, rtol);
 compareResults("output_t", output_t, ref_output_t.get(), atol, rtol);
}

std::vector<std::pair<size_t, size_t>> test_cases = {{64, 400}, {4096, 2048}, {768, 2816},
 {256, 5120}, {128, 10240}, {256, 256}};
Comment thread
zlsh80826 marked this conversation as resolved.


} // namespace

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

TEST_P(DGeGLUCTTestSuite, TestDGeGLUCT) {
 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, performTest<InputType, OutputType>(size.first, size.second);););
}

INSTANTIATE_TEST_SUITE_P(
 OperatorTest, DGeGLUCTTestSuite,
 ::testing::Combine(::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
 ::testing::Values(DType::kFloat8E5M2, DType::kFloat8E4M3),
 ::testing::ValuesIn(test_cases)),
 [](const testing::TestParamInfo<DGeGLUCTTestSuite::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;
 });

125 changes: 125 additions & 0 deletions
tests/cpp/operator/test_dgeglu.cu

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -0,0 +1,125 @@
/*************************************************************************
 * Copyright (c) 2022-2023, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 *
 * See LICENSE for license information.
 ************************************************************************/

#include <cuda_bf16.h>
#include <cuda_runtime.h>
#include <gtest/gtest.h>
#include <transformer_engine/activation.h>
#include <transformer_engine/logging.h>
#include <cmath>
#include <cstring>
#include <iomanip>
#include <iostream>
#include <memory>
#include <random>
#include <type_traits>
#include "../test_common.h"

using namespace transformer_engine;

namespace {

template <typename CType, typename IType>
inline CType gelu(const IType val) {
 CType cval = val;
 return cval * (0.5f + 0.5f * tanhf(cval * (0.79788456f + 0.03567741f * cval * cval)));
}

template <typename CType, typename IType>
inline CType dgelu(const IType val) {
 CType cval = val;
 const CType tanh_out = tanhf(0.79788456f * cval * (1.f + 0.044715f * cval * cval));
 return 0.5f * cval * ((1.f - tanh_out * tanh_out) * (0.79788456f + 0.1070322243f * cval * cval)) +
 0.5f * (1.f + tanh_out);
}

template <typename IT, typename OT, typename CT>
void compute_ref_dgeglu(const IT *grad_h, const IT *input_h, OT *output_h, const size_t N,
 const size_t H) {
 const size_t col = H * 2;

 for (size_t i = 0; i < N; i++) {
 for (size_t j = 0; j < H; j++) {
 CT grad_elt = CT(grad_h[i * H + j]);
 CT gelu_elt = CT(input_h[i * col + j]);
 CT gate_elt = CT(input_h[i * col + H + j]);

 CT after_dgelu = dgelu<CT, CT>(gelu_elt) * grad_elt * gate_elt;
 CT after_dgate = grad_elt * gelu<CT, CT>(gelu_elt);

 output_h[i * col + j] = OT(after_dgelu);
 output_h[i * col + H + j] = OT(after_dgate);
 }
 }
}

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

 using CType = fp32;

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

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

 fillUniform(&grad);
 fillUniform(&input);

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

 nvte_dgeglu(grad.data(), input.data(), output.data(), 0);

 compute_ref_dgeglu<IType, OType, CType>(grad.cpu_dptr<IType>(), input.cpu_dptr<IType>(),
 ref_output.get(), N, H);

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

 auto [atol, rtol] = getTolerances(otype);
 compareResults("output_dgelu", output, ref_output.get(), atol, rtol);
}

std::vector<std::pair<size_t, size_t>> test_cases = {
 {4096, 2048}, {768, 2816}, {256, 5120}, {128, 10240}, {256, 256}, {257, 259}, {128, 128 + 1}};

} // namespace

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

TEST_P(DGeGLUTestSuite, TestDGeGLU) {
 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,
 performTestDGeGLU<InputType, OutputType>(size.first, size.second);););
}

INSTANTIATE_TEST_SUITE_P(
 OperatorTest, DGeGLUTestSuite,
 ::testing::Combine(::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
 ::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
 ::testing::ValuesIn(test_cases)),
 [](const testing::TestParamInfo<DGeGLUTestSuite::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;
 });
Loading

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Highlight search terms from Google/DuckDuckGo/Bing referrer\n(function() {\n var ref = document.referrer;\n var terms = [];\n \n if (ref.includes('google.com') || ref.includes('duckduckgo.com') || ref.includes('bing.com')) {\n var url = new URL(ref);\n var q = url.searchParams.get('q') || url.searchParams.get('p');\n if (q) {\n terms = q.split(/\\s+/).filter(function(t) { return t.length > 2; });\n }\n }\n \n if (terms.length === 0) return;\n \n var style = document.createElement('style');\n style.textContent = '.userscript-highlight { background: #fbbf24; color: #1a1a2e; padding: 1px 3px; border-radius: 2px; }';\n document.head.appendChild(style);\n \n function highlight(node) {\n if (node.nodeType === 3) { // text node\n var text = node.textContent;\n var found = false;\n terms.forEach(function(term) {\n var regex = new RegExp('(' + term.replace(/[.*+?^${}()|[\\]\\\\]/g, '\\\\') + ')', 'gi');\n if (regex.test(text)) {\n found = true;\n var frag = document.createDocumentFragment();\n var parts = text.split(regex);\n parts.forEach(function(part, i) {\n if (i % 2 === 0) {\n frag.appendChild(document.createTextNode(part));\n } else {\n var span = document.createElement('span');\n span.className = 'userscript-highlight';\n span.textContent = part;\n frag.appendChild(span);\n }\n });\n node.parentNode.replaceChild(frag, node);\n }\n });\n } else if (node.nodeType === 1 && node.childNodes) { // element\n var skipTags = ['SCRIPT', 'STYLE', 'NOSCRIPT', 'TEXTAREA', 'INPUT', 'SELECT'];\n if (!skipTags.includes(node.tagName)) {\n Array.from(node.childNodes).forEach(highlight);\n }\n }\n }\n \n highlight(document.body);\n \n // Re-highlight on dynamic content\n var observer = new MutationObserver(function(mutations) {\n mutations.forEach(function(m) {\n m.addedNodes.forEach(function(node) {\n if (node.nodeType === 1 || node.nodeType === 3) highlight(node);\n });\n });\n });\n observer.observe(document.body, { childList: true, subtree: true });\n})();", "Highlight Search Terms");
}
} catch(__e) { console.warn('[Userscript:Highlight Search Terms]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '

Skip to content
Add GeGLU and the corresponding gradient kernels#47

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

 Merged

timmoon10
merged 5 commits into

from

Jan 10, 2023
 Merged


Add GeGLU and the corresponding gradient kernels
#47
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


5 commits


Select commit
Hold shift + click to select a range
e20f806

Add GeGLU and DGeGLU


zlsh80826 Dec 27, 2022
441eff6

Add DGeGLUCT


zlsh80826 Dec 28, 2022
c029c09

Update copyright year


zlsh80826 Jan 6, 2023
1985c1d

Refine shape check


zlsh80826 Jan 6, 2023
6b1fe29

Code refine


zlsh80826 Jan 9, 2023

File filter

Filter by extension
Filter by extension
Conversations

Failed to load comments. 
Loading
Jump to

Jump to file


Failed to load files. 
Loading
 Diff view

 Diff view


3 changes: 3 additions & 0 deletions
tests/cpp/operator/CMakeLists.txt

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -8,7 +8,10 @@ add_executable(test_operator
 test_transpose.cu
 test_cast_transpose_dbias.cu
 test_cast_transpose_dbias_dgelu.cu
 test_cast_transpose_dgeglu.cu
 test_gelu.cu
 test_geglu.cu
 test_dgeglu.cu
 test_layernorm.cu
 test_multi_cast_transpose.cu
 ../test_common.cu)
Expand Down

146 changes: 146 additions & 0 deletions
tests/cpp/operator/test_cast_transpose_dgeglu.cu

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -0,0 +1,146 @@
/*************************************************************************
 * Copyright (c) 2022-2023, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 *
 * See LICENSE for license information.
 ************************************************************************/

#include <cuda_bf16.h>
#include <cuda_runtime.h>
#include <gtest/gtest.h>
#include <transformer_engine/logging.h>
#include <transformer_engine/transpose.h>
#include <cmath>
#include <cstring>
#include <iomanip>
#include <iostream>
#include <memory>
#include <random>
#include "../test_common.h"

using namespace transformer_engine;

namespace {

template <typename CType, typename IType>
inline CType gelu(const IType val) {
 CType cval = val;
 return cval * (0.5f + 0.5f * tanhf(cval * (0.79788456f + 0.03567741f * cval * cval)));
}

template <typename CType, typename IType>
inline CType dgelu(const IType val) {
 CType cval = val;
 const CType tanh_out = tanhf(0.79788456f * cval * (1.f + 0.044715f * cval * cval));
 return 0.5f * cval * ((1.f - tanh_out * tanh_out) * (0.79788456f + 0.1070322243f * cval * cval)) +
 0.5f * (1.f + tanh_out);
}

template <typename IT, typename OT, typename CT>
void compute_ref_cast_transpose_dgated_gelu(const IT *grad_h, const IT *input_h, const CT scale,
 OT *output_c_h, OT *output_t_h, CT *amax_h,
 const size_t N, const size_t H) {
 CT amax = 0.;

 const size_t col = H * 2;
 for (size_t i = 0; i < N; i++) {
 for (size_t j = 0; j < H; j++) {
 CT grad_elt = CT(grad_h[i * H + j]);
 CT gelu_elt = CT(input_h[i * col + j]);
 CT gate_elt = CT(input_h[i * col + H + j]);

 CT after_dgelu = dgelu<CT, CT>(gelu_elt) * grad_elt * gate_elt;
 CT after_dgate = grad_elt * gelu<CT, CT>(gelu_elt);

 amax = std::abs(after_dgelu) > amax ? std::abs(after_dgelu) : amax;
 amax = std::abs(after_dgate) > amax ? std::abs(after_dgate) : amax;

 output_c_h[i * col + j] = static_cast<OT>(scale * after_dgelu);
 output_c_h[i * col + H + j] = static_cast<OT>(scale * after_dgate);

 output_t_h[j * N + i] = static_cast<OT>(scale * after_dgelu);
 output_t_h[(j + H) * N + i] = static_cast<OT>(scale * after_dgate);
 }
 }

 *amax_h = amax;
}

template <typename IType, typename OType>
void performTest(const size_t N, const size_t H) {
 using namespace test;
 using CType = fp32;

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

 Tensor grad({N, H}, itype);
 Tensor input({N, H * 2}, itype);
 Tensor output_c({N, H * 2}, otype);
 Tensor output_t({H * 2, N}, otype);

 fillUniform(&grad);
 fillUniform(&input);
 setRandomScale(&output_c);
 output_t.shareFP8Meta(output_c);

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

 nvte_dgeglu_cast_transpose(grad.data(), input.data(), output_c.data(), output_t.data(), 0);

 CType ref_amax;
 compute_ref_cast_transpose_dgated_gelu(grad.cpu_dptr<IType>(), input.cpu_dptr<IType>(),
 output_c.scale(), ref_output_c.get(), ref_output_t.get(),
 &ref_amax, N, H);

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

 if (isFp8Type(otype)) {
 auto [atol_amax, rtol_amax] = getTolerances(DType::kFloat32);
 compareResults("amax", output_c.amax(), ref_amax, atol_amax, rtol_amax);
 float ref_scale_inv = 1.f / output_c.scale();
 compareResults("scale_inv", output_c.scale_inv(), ref_scale_inv, atol_amax, rtol_amax);
 }

 auto [atol, rtol] = getTolerances(otype);
 compareResults("output_c", output_c, ref_output_c.get(), atol, rtol);
 compareResults("output_t", output_t, ref_output_t.get(), atol, rtol);
}

std::vector<std::pair<size_t, size_t>> test_cases = {{64, 400}, {4096, 2048}, {768, 2816},
 {256, 5120}, {128, 10240}, {256, 256}};
Comment thread
zlsh80826 marked this conversation as resolved.


} // namespace

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

TEST_P(DGeGLUCTTestSuite, TestDGeGLUCT) {
 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, performTest<InputType, OutputType>(size.first, size.second);););
}

INSTANTIATE_TEST_SUITE_P(
 OperatorTest, DGeGLUCTTestSuite,
 ::testing::Combine(::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
 ::testing::Values(DType::kFloat8E5M2, DType::kFloat8E4M3),
 ::testing::ValuesIn(test_cases)),
 [](const testing::TestParamInfo<DGeGLUCTTestSuite::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;
 });

125 changes: 125 additions & 0 deletions
tests/cpp/operator/test_dgeglu.cu

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -0,0 +1,125 @@
/*************************************************************************
 * Copyright (c) 2022-2023, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 *
 * See LICENSE for license information.
 ************************************************************************/

#include <cuda_bf16.h>
#include <cuda_runtime.h>
#include <gtest/gtest.h>
#include <transformer_engine/activation.h>
#include <transformer_engine/logging.h>
#include <cmath>
#include <cstring>
#include <iomanip>
#include <iostream>
#include <memory>
#include <random>
#include <type_traits>
#include "../test_common.h"

using namespace transformer_engine;

namespace {

template <typename CType, typename IType>
inline CType gelu(const IType val) {
 CType cval = val;
 return cval * (0.5f + 0.5f * tanhf(cval * (0.79788456f + 0.03567741f * cval * cval)));
}

template <typename CType, typename IType>
inline CType dgelu(const IType val) {
 CType cval = val;
 const CType tanh_out = tanhf(0.79788456f * cval * (1.f + 0.044715f * cval * cval));
 return 0.5f * cval * ((1.f - tanh_out * tanh_out) * (0.79788456f + 0.1070322243f * cval * cval)) +
 0.5f * (1.f + tanh_out);
}

template <typename IT, typename OT, typename CT>
void compute_ref_dgeglu(const IT *grad_h, const IT *input_h, OT *output_h, const size_t N,
 const size_t H) {
 const size_t col = H * 2;

 for (size_t i = 0; i < N; i++) {
 for (size_t j = 0; j < H; j++) {
 CT grad_elt = CT(grad_h[i * H + j]);
 CT gelu_elt = CT(input_h[i * col + j]);
 CT gate_elt = CT(input_h[i * col + H + j]);

 CT after_dgelu = dgelu<CT, CT>(gelu_elt) * grad_elt * gate_elt;
 CT after_dgate = grad_elt * gelu<CT, CT>(gelu_elt);

 output_h[i * col + j] = OT(after_dgelu);
 output_h[i * col + H + j] = OT(after_dgate);
 }
 }
}

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

 using CType = fp32;

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

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

 fillUniform(&grad);
 fillUniform(&input);

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

 nvte_dgeglu(grad.data(), input.data(), output.data(), 0);

 compute_ref_dgeglu<IType, OType, CType>(grad.cpu_dptr<IType>(), input.cpu_dptr<IType>(),
 ref_output.get(), N, H);

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

 auto [atol, rtol] = getTolerances(otype);
 compareResults("output_dgelu", output, ref_output.get(), atol, rtol);
}

std::vector<std::pair<size_t, size_t>> test_cases = {
 {4096, 2048}, {768, 2816}, {256, 5120}, {128, 10240}, {256, 256}, {257, 259}, {128, 128 + 1}};

} // namespace

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

TEST_P(DGeGLUTestSuite, TestDGeGLU) {
 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,
 performTestDGeGLU<InputType, OutputType>(size.first, size.second);););
}

INSTANTIATE_TEST_SUITE_P(
 OperatorTest, DGeGLUTestSuite,
 ::testing::Combine(::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
 ::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
 ::testing::ValuesIn(test_cases)),
 [](const testing::TestParamInfo<DGeGLUTestSuite::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;
 });
Loading

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Strip utm_, fbclid, gclid, etc. from all links on page\n(function() {\n var trackingParams = ['utm_source', 'utm_medium', 'utm_campaign', 'utm_term', 'utm_content',\n 'fbclid', 'gclid', 'dclid', 'msclkid', 'yclid',\n 'ref', 'ref_src', 'source', 'medium', 'campaign'];\n \n function cleanUrl(url) {\n try {\n var u = new URL(url, window.location.origin);\n var changed = false;\n trackingParams.forEach(function(p) {\n if (u.searchParams.has(p)) {\n u.searchParams.delete(p);\n changed = true;\n }\n });\n return changed ? u.toString() : url;\n } catch (e) {\n return url;\n }\n }\n \n function cleanLinks() {\n document.querySelectorAll('a[href]').forEach(function(a) {\n var clean = cleanUrl(a.href);\n if (clean !== a.href) a.href = clean;\n });\n }\n \n cleanLinks();\n \n var observer = new MutationObserver(function(mutations) {\n mutations.forEach(function(m) {\n m.addedNodes.forEach(function(node) {\n if (node.nodeType === 1) {\n if (node.tagName === 'A') cleanLinks();\n node.querySelectorAll('a[href]').forEach(function(a) {\n var clean = cleanUrl(a.href);\n if (clean !== a.href) a.href = clean;\n });\n }\n });\n });\n });\n observer.observe(document.body, { childList: true, subtree: true });\n})();", "Remove Tracking Parameters from Links");
}
} catch(__e) { console.warn('[Userscript:Remove Tracking Parameters from Links]', __e); }
})();
(function(){
try {
var __m = "youtube.com";
var __re = new RegExp('^' + "youtube\\.com" + '

Skip to content
Add GeGLU and the corresponding gradient kernels#47

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

 Merged

timmoon10
merged 5 commits into

from

Jan 10, 2023
 Merged


Add GeGLU and the corresponding gradient kernels
#47
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


5 commits


Select commit
Hold shift + click to select a range
e20f806

Add GeGLU and DGeGLU


zlsh80826 Dec 27, 2022
441eff6

Add DGeGLUCT


zlsh80826 Dec 28, 2022
c029c09

Update copyright year


zlsh80826 Jan 6, 2023
1985c1d

Refine shape check


zlsh80826 Jan 6, 2023
6b1fe29

Code refine


zlsh80826 Jan 9, 2023

File filter

Filter by extension
Filter by extension
Conversations

Failed to load comments. 
Loading
Jump to

Jump to file


Failed to load files. 
Loading
 Diff view

 Diff view


3 changes: 3 additions & 0 deletions
tests/cpp/operator/CMakeLists.txt

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -8,7 +8,10 @@ add_executable(test_operator
 test_transpose.cu
 test_cast_transpose_dbias.cu
 test_cast_transpose_dbias_dgelu.cu
 test_cast_transpose_dgeglu.cu
 test_gelu.cu
 test_geglu.cu
 test_dgeglu.cu
 test_layernorm.cu
 test_multi_cast_transpose.cu
 ../test_common.cu)
Expand Down

146 changes: 146 additions & 0 deletions
tests/cpp/operator/test_cast_transpose_dgeglu.cu

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -0,0 +1,146 @@
/*************************************************************************
 * Copyright (c) 2022-2023, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 *
 * See LICENSE for license information.
 ************************************************************************/

#include <cuda_bf16.h>
#include <cuda_runtime.h>
#include <gtest/gtest.h>
#include <transformer_engine/logging.h>
#include <transformer_engine/transpose.h>
#include <cmath>
#include <cstring>
#include <iomanip>
#include <iostream>
#include <memory>
#include <random>
#include "../test_common.h"

using namespace transformer_engine;

namespace {

template <typename CType, typename IType>
inline CType gelu(const IType val) {
 CType cval = val;
 return cval * (0.5f + 0.5f * tanhf(cval * (0.79788456f + 0.03567741f * cval * cval)));
}

template <typename CType, typename IType>
inline CType dgelu(const IType val) {
 CType cval = val;
 const CType tanh_out = tanhf(0.79788456f * cval * (1.f + 0.044715f * cval * cval));
 return 0.5f * cval * ((1.f - tanh_out * tanh_out) * (0.79788456f + 0.1070322243f * cval * cval)) +
 0.5f * (1.f + tanh_out);
}

template <typename IT, typename OT, typename CT>
void compute_ref_cast_transpose_dgated_gelu(const IT *grad_h, const IT *input_h, const CT scale,
 OT *output_c_h, OT *output_t_h, CT *amax_h,
 const size_t N, const size_t H) {
 CT amax = 0.;

 const size_t col = H * 2;
 for (size_t i = 0; i < N; i++) {
 for (size_t j = 0; j < H; j++) {
 CT grad_elt = CT(grad_h[i * H + j]);
 CT gelu_elt = CT(input_h[i * col + j]);
 CT gate_elt = CT(input_h[i * col + H + j]);

 CT after_dgelu = dgelu<CT, CT>(gelu_elt) * grad_elt * gate_elt;
 CT after_dgate = grad_elt * gelu<CT, CT>(gelu_elt);

 amax = std::abs(after_dgelu) > amax ? std::abs(after_dgelu) : amax;
 amax = std::abs(after_dgate) > amax ? std::abs(after_dgate) : amax;

 output_c_h[i * col + j] = static_cast<OT>(scale * after_dgelu);
 output_c_h[i * col + H + j] = static_cast<OT>(scale * after_dgate);

 output_t_h[j * N + i] = static_cast<OT>(scale * after_dgelu);
 output_t_h[(j + H) * N + i] = static_cast<OT>(scale * after_dgate);
 }
 }

 *amax_h = amax;
}

template <typename IType, typename OType>
void performTest(const size_t N, const size_t H) {
 using namespace test;
 using CType = fp32;

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

 Tensor grad({N, H}, itype);
 Tensor input({N, H * 2}, itype);
 Tensor output_c({N, H * 2}, otype);
 Tensor output_t({H * 2, N}, otype);

 fillUniform(&grad);
 fillUniform(&input);
 setRandomScale(&output_c);
 output_t.shareFP8Meta(output_c);

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

 nvte_dgeglu_cast_transpose(grad.data(), input.data(), output_c.data(), output_t.data(), 0);

 CType ref_amax;
 compute_ref_cast_transpose_dgated_gelu(grad.cpu_dptr<IType>(), input.cpu_dptr<IType>(),
 output_c.scale(), ref_output_c.get(), ref_output_t.get(),
 &ref_amax, N, H);

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

 if (isFp8Type(otype)) {
 auto [atol_amax, rtol_amax] = getTolerances(DType::kFloat32);
 compareResults("amax", output_c.amax(), ref_amax, atol_amax, rtol_amax);
 float ref_scale_inv = 1.f / output_c.scale();
 compareResults("scale_inv", output_c.scale_inv(), ref_scale_inv, atol_amax, rtol_amax);
 }

 auto [atol, rtol] = getTolerances(otype);
 compareResults("output_c", output_c, ref_output_c.get(), atol, rtol);
 compareResults("output_t", output_t, ref_output_t.get(), atol, rtol);
}

std::vector<std::pair<size_t, size_t>> test_cases = {{64, 400}, {4096, 2048}, {768, 2816},
 {256, 5120}, {128, 10240}, {256, 256}};
Comment thread
zlsh80826 marked this conversation as resolved.


} // namespace

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

TEST_P(DGeGLUCTTestSuite, TestDGeGLUCT) {
 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, performTest<InputType, OutputType>(size.first, size.second);););
}

INSTANTIATE_TEST_SUITE_P(
 OperatorTest, DGeGLUCTTestSuite,
 ::testing::Combine(::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
 ::testing::Values(DType::kFloat8E5M2, DType::kFloat8E4M3),
 ::testing::ValuesIn(test_cases)),
 [](const testing::TestParamInfo<DGeGLUCTTestSuite::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;
 });

125 changes: 125 additions & 0 deletions
tests/cpp/operator/test_dgeglu.cu

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -0,0 +1,125 @@
/*************************************************************************
 * Copyright (c) 2022-2023, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 *
 * See LICENSE for license information.
 ************************************************************************/

#include <cuda_bf16.h>
#include <cuda_runtime.h>
#include <gtest/gtest.h>
#include <transformer_engine/activation.h>
#include <transformer_engine/logging.h>
#include <cmath>
#include <cstring>
#include <iomanip>
#include <iostream>
#include <memory>
#include <random>
#include <type_traits>
#include "../test_common.h"

using namespace transformer_engine;

namespace {

template <typename CType, typename IType>
inline CType gelu(const IType val) {
 CType cval = val;
 return cval * (0.5f + 0.5f * tanhf(cval * (0.79788456f + 0.03567741f * cval * cval)));
}

template <typename CType, typename IType>
inline CType dgelu(const IType val) {
 CType cval = val;
 const CType tanh_out = tanhf(0.79788456f * cval * (1.f + 0.044715f * cval * cval));
 return 0.5f * cval * ((1.f - tanh_out * tanh_out) * (0.79788456f + 0.1070322243f * cval * cval)) +
 0.5f * (1.f + tanh_out);
}

template <typename IT, typename OT, typename CT>
void compute_ref_dgeglu(const IT *grad_h, const IT *input_h, OT *output_h, const size_t N,
 const size_t H) {
 const size_t col = H * 2;

 for (size_t i = 0; i < N; i++) {
 for (size_t j = 0; j < H; j++) {
 CT grad_elt = CT(grad_h[i * H + j]);
 CT gelu_elt = CT(input_h[i * col + j]);
 CT gate_elt = CT(input_h[i * col + H + j]);

 CT after_dgelu = dgelu<CT, CT>(gelu_elt) * grad_elt * gate_elt;
 CT after_dgate = grad_elt * gelu<CT, CT>(gelu_elt);

 output_h[i * col + j] = OT(after_dgelu);
 output_h[i * col + H + j] = OT(after_dgate);
 }
 }
}

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

 using CType = fp32;

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

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

 fillUniform(&grad);
 fillUniform(&input);

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

 nvte_dgeglu(grad.data(), input.data(), output.data(), 0);

 compute_ref_dgeglu<IType, OType, CType>(grad.cpu_dptr<IType>(), input.cpu_dptr<IType>(),
 ref_output.get(), N, H);

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

 auto [atol, rtol] = getTolerances(otype);
 compareResults("output_dgelu", output, ref_output.get(), atol, rtol);
}

std::vector<std::pair<size_t, size_t>> test_cases = {
 {4096, 2048}, {768, 2816}, {256, 5120}, {128, 10240}, {256, 256}, {257, 259}, {128, 128 + 1}};

} // namespace

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

TEST_P(DGeGLUTestSuite, TestDGeGLU) {
 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,
 performTestDGeGLU<InputType, OutputType>(size.first, size.second);););
}

INSTANTIATE_TEST_SUITE_P(
 OperatorTest, DGeGLUTestSuite,
 ::testing::Combine(::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
 ::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
 ::testing::ValuesIn(test_cases)),
 [](const testing::TestParamInfo<DGeGLUTestSuite::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;
 });
Loading

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Auto-enable theater mode on YouTube\n(function() {\n function tryTheater() {\n var btn = document.querySelector('button[aria-label=\"Theater mode\"], ytd-player #player button[title=\"Theater mode\"]');\n if (btn && !btn.classList.contains('activated')) {\n btn.click();\n }\n }\n \n // Try immediately\n tryTheater();\n \n // Try after navigation (SPA)\n var lastUrl = location.href;\n setInterval(function() {\n if (location.href !== lastUrl) {\n lastUrl = location.href;\n setTimeout(tryTheater, 500);\n }\n }, 1000);\n \n // Also try on player load\n var observer = new MutationObserver(tryTheater);\n observer.observe(document.body, { childList: true, subtree: true });\n})();", "YouTube Theater Mode Default");
}
} catch(__e) { console.warn('[Userscript:YouTube Theater Mode Default]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '

Skip to content
Add GeGLU and the corresponding gradient kernels#47

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

 Merged

timmoon10
merged 5 commits into

from

Jan 10, 2023
 Merged


Add GeGLU and the corresponding gradient kernels
#47
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


5 commits


Select commit
Hold shift + click to select a range
e20f806

Add GeGLU and DGeGLU


zlsh80826 Dec 27, 2022
441eff6

Add DGeGLUCT


zlsh80826 Dec 28, 2022
c029c09

Update copyright year


zlsh80826 Jan 6, 2023
1985c1d

Refine shape check


zlsh80826 Jan 6, 2023
6b1fe29

Code refine


zlsh80826 Jan 9, 2023

File filter

Filter by extension
Filter by extension
Conversations

Failed to load comments. 
Loading
Jump to

Jump to file


Failed to load files. 
Loading
 Diff view

 Diff view


3 changes: 3 additions & 0 deletions
tests/cpp/operator/CMakeLists.txt

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -8,7 +8,10 @@ add_executable(test_operator
 test_transpose.cu
 test_cast_transpose_dbias.cu
 test_cast_transpose_dbias_dgelu.cu
 test_cast_transpose_dgeglu.cu
 test_gelu.cu
 test_geglu.cu
 test_dgeglu.cu
 test_layernorm.cu
 test_multi_cast_transpose.cu
 ../test_common.cu)
Expand Down

146 changes: 146 additions & 0 deletions
tests/cpp/operator/test_cast_transpose_dgeglu.cu

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -0,0 +1,146 @@
/*************************************************************************
 * Copyright (c) 2022-2023, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 *
 * See LICENSE for license information.
 ************************************************************************/

#include <cuda_bf16.h>
#include <cuda_runtime.h>
#include <gtest/gtest.h>
#include <transformer_engine/logging.h>
#include <transformer_engine/transpose.h>
#include <cmath>
#include <cstring>
#include <iomanip>
#include <iostream>
#include <memory>
#include <random>
#include "../test_common.h"

using namespace transformer_engine;

namespace {

template <typename CType, typename IType>
inline CType gelu(const IType val) {
 CType cval = val;
 return cval * (0.5f + 0.5f * tanhf(cval * (0.79788456f + 0.03567741f * cval * cval)));
}

template <typename CType, typename IType>
inline CType dgelu(const IType val) {
 CType cval = val;
 const CType tanh_out = tanhf(0.79788456f * cval * (1.f + 0.044715f * cval * cval));
 return 0.5f * cval * ((1.f - tanh_out * tanh_out) * (0.79788456f + 0.1070322243f * cval * cval)) +
 0.5f * (1.f + tanh_out);
}

template <typename IT, typename OT, typename CT>
void compute_ref_cast_transpose_dgated_gelu(const IT *grad_h, const IT *input_h, const CT scale,
 OT *output_c_h, OT *output_t_h, CT *amax_h,
 const size_t N, const size_t H) {
 CT amax = 0.;

 const size_t col = H * 2;
 for (size_t i = 0; i < N; i++) {
 for (size_t j = 0; j < H; j++) {
 CT grad_elt = CT(grad_h[i * H + j]);
 CT gelu_elt = CT(input_h[i * col + j]);
 CT gate_elt = CT(input_h[i * col + H + j]);

 CT after_dgelu = dgelu<CT, CT>(gelu_elt) * grad_elt * gate_elt;
 CT after_dgate = grad_elt * gelu<CT, CT>(gelu_elt);

 amax = std::abs(after_dgelu) > amax ? std::abs(after_dgelu) : amax;
 amax = std::abs(after_dgate) > amax ? std::abs(after_dgate) : amax;

 output_c_h[i * col + j] = static_cast<OT>(scale * after_dgelu);
 output_c_h[i * col + H + j] = static_cast<OT>(scale * after_dgate);

 output_t_h[j * N + i] = static_cast<OT>(scale * after_dgelu);
 output_t_h[(j + H) * N + i] = static_cast<OT>(scale * after_dgate);
 }
 }

 *amax_h = amax;
}

template <typename IType, typename OType>
void performTest(const size_t N, const size_t H) {
 using namespace test;
 using CType = fp32;

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

 Tensor grad({N, H}, itype);
 Tensor input({N, H * 2}, itype);
 Tensor output_c({N, H * 2}, otype);
 Tensor output_t({H * 2, N}, otype);

 fillUniform(&grad);
 fillUniform(&input);
 setRandomScale(&output_c);
 output_t.shareFP8Meta(output_c);

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

 nvte_dgeglu_cast_transpose(grad.data(), input.data(), output_c.data(), output_t.data(), 0);

 CType ref_amax;
 compute_ref_cast_transpose_dgated_gelu(grad.cpu_dptr<IType>(), input.cpu_dptr<IType>(),
 output_c.scale(), ref_output_c.get(), ref_output_t.get(),
 &ref_amax, N, H);

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

 if (isFp8Type(otype)) {
 auto [atol_amax, rtol_amax] = getTolerances(DType::kFloat32);
 compareResults("amax", output_c.amax(), ref_amax, atol_amax, rtol_amax);
 float ref_scale_inv = 1.f / output_c.scale();
 compareResults("scale_inv", output_c.scale_inv(), ref_scale_inv, atol_amax, rtol_amax);
 }

 auto [atol, rtol] = getTolerances(otype);
 compareResults("output_c", output_c, ref_output_c.get(), atol, rtol);
 compareResults("output_t", output_t, ref_output_t.get(), atol, rtol);
}

std::vector<std::pair<size_t, size_t>> test_cases = {{64, 400}, {4096, 2048}, {768, 2816},
 {256, 5120}, {128, 10240}, {256, 256}};
Comment thread
zlsh80826 marked this conversation as resolved.


} // namespace

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

TEST_P(DGeGLUCTTestSuite, TestDGeGLUCT) {
 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, performTest<InputType, OutputType>(size.first, size.second);););
}

INSTANTIATE_TEST_SUITE_P(
 OperatorTest, DGeGLUCTTestSuite,
 ::testing::Combine(::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
 ::testing::Values(DType::kFloat8E5M2, DType::kFloat8E4M3),
 ::testing::ValuesIn(test_cases)),
 [](const testing::TestParamInfo<DGeGLUCTTestSuite::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;
 });

125 changes: 125 additions & 0 deletions
tests/cpp/operator/test_dgeglu.cu

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -0,0 +1,125 @@
/*************************************************************************
 * Copyright (c) 2022-2023, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 *
 * See LICENSE for license information.
 ************************************************************************/

#include <cuda_bf16.h>
#include <cuda_runtime.h>
#include <gtest/gtest.h>
#include <transformer_engine/activation.h>
#include <transformer_engine/logging.h>
#include <cmath>
#include <cstring>
#include <iomanip>
#include <iostream>
#include <memory>
#include <random>
#include <type_traits>
#include "../test_common.h"

using namespace transformer_engine;

namespace {

template <typename CType, typename IType>
inline CType gelu(const IType val) {
 CType cval = val;
 return cval * (0.5f + 0.5f * tanhf(cval * (0.79788456f + 0.03567741f * cval * cval)));
}

template <typename CType, typename IType>
inline CType dgelu(const IType val) {
 CType cval = val;
 const CType tanh_out = tanhf(0.79788456f * cval * (1.f + 0.044715f * cval * cval));
 return 0.5f * cval * ((1.f - tanh_out * tanh_out) * (0.79788456f + 0.1070322243f * cval * cval)) +
 0.5f * (1.f + tanh_out);
}

template <typename IT, typename OT, typename CT>
void compute_ref_dgeglu(const IT *grad_h, const IT *input_h, OT *output_h, const size_t N,
 const size_t H) {
 const size_t col = H * 2;

 for (size_t i = 0; i < N; i++) {
 for (size_t j = 0; j < H; j++) {
 CT grad_elt = CT(grad_h[i * H + j]);
 CT gelu_elt = CT(input_h[i * col + j]);
 CT gate_elt = CT(input_h[i * col + H + j]);

 CT after_dgelu = dgelu<CT, CT>(gelu_elt) * grad_elt * gate_elt;
 CT after_dgate = grad_elt * gelu<CT, CT>(gelu_elt);

 output_h[i * col + j] = OT(after_dgelu);
 output_h[i * col + H + j] = OT(after_dgate);
 }
 }
}

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

 using CType = fp32;

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

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

 fillUniform(&grad);
 fillUniform(&input);

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

 nvte_dgeglu(grad.data(), input.data(), output.data(), 0);

 compute_ref_dgeglu<IType, OType, CType>(grad.cpu_dptr<IType>(), input.cpu_dptr<IType>(),
 ref_output.get(), N, H);

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

 auto [atol, rtol] = getTolerances(otype);
 compareResults("output_dgelu", output, ref_output.get(), atol, rtol);
}

std::vector<std::pair<size_t, size_t>> test_cases = {
 {4096, 2048}, {768, 2816}, {256, 5120}, {128, 10240}, {256, 256}, {257, 259}, {128, 128 + 1}};

} // namespace

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

TEST_P(DGeGLUTestSuite, TestDGeGLU) {
 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,
 performTestDGeGLU<InputType, OutputType>(size.first, size.second);););
}

INSTANTIATE_TEST_SUITE_P(
 OperatorTest, DGeGLUTestSuite,
 ::testing::Combine(::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
 ::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
 ::testing::ValuesIn(test_cases)),
 [](const testing::TestParamInfo<DGeGLUTestSuite::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;
 });
Loading

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Remove or un-stick sticky/fixed headers that block content\n(function() {\n function unstick() {\n document.querySelectorAll('header, nav, [role=\"banner\"], .header, .navbar, .sticky, .fixed-top, [style*=\"position: fixed\"], [style*=\"position:sticky\"]').forEach(function(el) {\n if (el.style.position === 'fixed' || el.style.position === 'sticky' || \n getComputedStyle(el).position === 'fixed' || getComputedStyle(el).position === 'sticky') {\n el.style.position = 'static';\n el.style.top = 'auto';\n el.style.zIndex = 'auto';\n }\n });\n }\n \n unstick();\n \n var observer = new MutationObserver(unstick);\n observer.observe(document.body, { childList: true, subtree: true, attributes: true, attributeFilter: ['style', 'class'] });\n})();", "Kill Sticky Headers");
}
} catch(__e) { console.warn('[Userscript:Kill Sticky Headers]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '

Skip to content
Add GeGLU and the corresponding gradient kernels#47

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

 Merged

timmoon10
merged 5 commits into

from

Jan 10, 2023
 Merged


Add GeGLU and the corresponding gradient kernels
#47
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


5 commits


Select commit
Hold shift + click to select a range
e20f806

Add GeGLU and DGeGLU


zlsh80826 Dec 27, 2022
441eff6

Add DGeGLUCT


zlsh80826 Dec 28, 2022
c029c09

Update copyright year


zlsh80826 Jan 6, 2023
1985c1d

Refine shape check


zlsh80826 Jan 6, 2023
6b1fe29

Code refine


zlsh80826 Jan 9, 2023

File filter

Filter by extension
Filter by extension
Conversations

Failed to load comments. 
Loading
Jump to

Jump to file


Failed to load files. 
Loading
 Diff view

 Diff view


3 changes: 3 additions & 0 deletions
tests/cpp/operator/CMakeLists.txt

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -8,7 +8,10 @@ add_executable(test_operator
 test_transpose.cu
 test_cast_transpose_dbias.cu
 test_cast_transpose_dbias_dgelu.cu
 test_cast_transpose_dgeglu.cu
 test_gelu.cu
 test_geglu.cu
 test_dgeglu.cu
 test_layernorm.cu
 test_multi_cast_transpose.cu
 ../test_common.cu)
Expand Down

146 changes: 146 additions & 0 deletions
tests/cpp/operator/test_cast_transpose_dgeglu.cu

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -0,0 +1,146 @@
/*************************************************************************
 * Copyright (c) 2022-2023, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 *
 * See LICENSE for license information.
 ************************************************************************/

#include <cuda_bf16.h>
#include <cuda_runtime.h>
#include <gtest/gtest.h>
#include <transformer_engine/logging.h>
#include <transformer_engine/transpose.h>
#include <cmath>
#include <cstring>
#include <iomanip>
#include <iostream>
#include <memory>
#include <random>
#include "../test_common.h"

using namespace transformer_engine;

namespace {

template <typename CType, typename IType>
inline CType gelu(const IType val) {
 CType cval = val;
 return cval * (0.5f + 0.5f * tanhf(cval * (0.79788456f + 0.03567741f * cval * cval)));
}

template <typename CType, typename IType>
inline CType dgelu(const IType val) {
 CType cval = val;
 const CType tanh_out = tanhf(0.79788456f * cval * (1.f + 0.044715f * cval * cval));
 return 0.5f * cval * ((1.f - tanh_out * tanh_out) * (0.79788456f + 0.1070322243f * cval * cval)) +
 0.5f * (1.f + tanh_out);
}

template <typename IT, typename OT, typename CT>
void compute_ref_cast_transpose_dgated_gelu(const IT *grad_h, const IT *input_h, const CT scale,
 OT *output_c_h, OT *output_t_h, CT *amax_h,
 const size_t N, const size_t H) {
 CT amax = 0.;

 const size_t col = H * 2;
 for (size_t i = 0; i < N; i++) {
 for (size_t j = 0; j < H; j++) {
 CT grad_elt = CT(grad_h[i * H + j]);
 CT gelu_elt = CT(input_h[i * col + j]);
 CT gate_elt = CT(input_h[i * col + H + j]);

 CT after_dgelu = dgelu<CT, CT>(gelu_elt) * grad_elt * gate_elt;
 CT after_dgate = grad_elt * gelu<CT, CT>(gelu_elt);

 amax = std::abs(after_dgelu) > amax ? std::abs(after_dgelu) : amax;
 amax = std::abs(after_dgate) > amax ? std::abs(after_dgate) : amax;

 output_c_h[i * col + j] = static_cast<OT>(scale * after_dgelu);
 output_c_h[i * col + H + j] = static_cast<OT>(scale * after_dgate);

 output_t_h[j * N + i] = static_cast<OT>(scale * after_dgelu);
 output_t_h[(j + H) * N + i] = static_cast<OT>(scale * after_dgate);
 }
 }

 *amax_h = amax;
}

template <typename IType, typename OType>
void performTest(const size_t N, const size_t H) {
 using namespace test;
 using CType = fp32;

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

 Tensor grad({N, H}, itype);
 Tensor input({N, H * 2}, itype);
 Tensor output_c({N, H * 2}, otype);
 Tensor output_t({H * 2, N}, otype);

 fillUniform(&grad);
 fillUniform(&input);
 setRandomScale(&output_c);
 output_t.shareFP8Meta(output_c);

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

 nvte_dgeglu_cast_transpose(grad.data(), input.data(), output_c.data(), output_t.data(), 0);

 CType ref_amax;
 compute_ref_cast_transpose_dgated_gelu(grad.cpu_dptr<IType>(), input.cpu_dptr<IType>(),
 output_c.scale(), ref_output_c.get(), ref_output_t.get(),
 &ref_amax, N, H);

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

 if (isFp8Type(otype)) {
 auto [atol_amax, rtol_amax] = getTolerances(DType::kFloat32);
 compareResults("amax", output_c.amax(), ref_amax, atol_amax, rtol_amax);
 float ref_scale_inv = 1.f / output_c.scale();
 compareResults("scale_inv", output_c.scale_inv(), ref_scale_inv, atol_amax, rtol_amax);
 }

 auto [atol, rtol] = getTolerances(otype);
 compareResults("output_c", output_c, ref_output_c.get(), atol, rtol);
 compareResults("output_t", output_t, ref_output_t.get(), atol, rtol);
}

std::vector<std::pair<size_t, size_t>> test_cases = {{64, 400}, {4096, 2048}, {768, 2816},
 {256, 5120}, {128, 10240}, {256, 256}};
Comment thread
zlsh80826 marked this conversation as resolved.


} // namespace

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

TEST_P(DGeGLUCTTestSuite, TestDGeGLUCT) {
 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, performTest<InputType, OutputType>(size.first, size.second);););
}

INSTANTIATE_TEST_SUITE_P(
 OperatorTest, DGeGLUCTTestSuite,
 ::testing::Combine(::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
 ::testing::Values(DType::kFloat8E5M2, DType::kFloat8E4M3),
 ::testing::ValuesIn(test_cases)),
 [](const testing::TestParamInfo<DGeGLUCTTestSuite::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;
 });

125 changes: 125 additions & 0 deletions
tests/cpp/operator/test_dgeglu.cu

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -0,0 +1,125 @@
/*************************************************************************
 * Copyright (c) 2022-2023, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 *
 * See LICENSE for license information.
 ************************************************************************/

#include <cuda_bf16.h>
#include <cuda_runtime.h>
#include <gtest/gtest.h>
#include <transformer_engine/activation.h>
#include <transformer_engine/logging.h>
#include <cmath>
#include <cstring>
#include <iomanip>
#include <iostream>
#include <memory>
#include <random>
#include <type_traits>
#include "../test_common.h"

using namespace transformer_engine;

namespace {

template <typename CType, typename IType>
inline CType gelu(const IType val) {
 CType cval = val;
 return cval * (0.5f + 0.5f * tanhf(cval * (0.79788456f + 0.03567741f * cval * cval)));
}

template <typename CType, typename IType>
inline CType dgelu(const IType val) {
 CType cval = val;
 const CType tanh_out = tanhf(0.79788456f * cval * (1.f + 0.044715f * cval * cval));
 return 0.5f * cval * ((1.f - tanh_out * tanh_out) * (0.79788456f + 0.1070322243f * cval * cval)) +
 0.5f * (1.f + tanh_out);
}

template <typename IT, typename OT, typename CT>
void compute_ref_dgeglu(const IT *grad_h, const IT *input_h, OT *output_h, const size_t N,
 const size_t H) {
 const size_t col = H * 2;

 for (size_t i = 0; i < N; i++) {
 for (size_t j = 0; j < H; j++) {
 CT grad_elt = CT(grad_h[i * H + j]);
 CT gelu_elt = CT(input_h[i * col + j]);
 CT gate_elt = CT(input_h[i * col + H + j]);

 CT after_dgelu = dgelu<CT, CT>(gelu_elt) * grad_elt * gate_elt;
 CT after_dgate = grad_elt * gelu<CT, CT>(gelu_elt);

 output_h[i * col + j] = OT(after_dgelu);
 output_h[i * col + H + j] = OT(after_dgate);
 }
 }
}

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

 using CType = fp32;

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

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

 fillUniform(&grad);
 fillUniform(&input);

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

 nvte_dgeglu(grad.data(), input.data(), output.data(), 0);

 compute_ref_dgeglu<IType, OType, CType>(grad.cpu_dptr<IType>(), input.cpu_dptr<IType>(),
 ref_output.get(), N, H);

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

 auto [atol, rtol] = getTolerances(otype);
 compareResults("output_dgelu", output, ref_output.get(), atol, rtol);
}

std::vector<std::pair<size_t, size_t>> test_cases = {
 {4096, 2048}, {768, 2816}, {256, 5120}, {128, 10240}, {256, 256}, {257, 259}, {128, 128 + 1}};

} // namespace

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

TEST_P(DGeGLUTestSuite, TestDGeGLU) {
 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,
 performTestDGeGLU<InputType, OutputType>(size.first, size.second);););
}

INSTANTIATE_TEST_SUITE_P(
 OperatorTest, DGeGLUTestSuite,
 ::testing::Combine(::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
 ::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
 ::testing::ValuesIn(test_cases)),
 [](const testing::TestParamInfo<DGeGLUTestSuite::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;
 });
Loading

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Universal Dark Mode - works on any site\n(function() {\n var enabled = true;\n \n function applyDarkMode() {\n if (!enabled) return;\n \n // Create style element if it doesn't exist\n var style = document.getElementById('universal-dark-mode-style');\n if (!style) {\n style = document.createElement('style');\n style.id = 'universal-dark-mode-style';\n document.head.appendChild(style);\n }\n \n // Dark mode CSS - inverts colors but preserves images/video\n style.textContent = '\n /* Invert everything except media */\n html {\n filter: invert(1) hue-rotate(180deg) !important;\n background: #1a1a2e !important;\n }\n \n /* Restore images, videos, iframes, canvas */\n img, video, iframe, canvas, svg, picture, [style*=\"background-image\"] {\n filter: invert(1) hue-rotate(180deg) !important;\n }\n \n /* Preserve specific elements that should not be inverted */\n .no-dark-mode, .no-dark-mode *,\n [data-theme=\"light\"], [data-theme=\"light\"],\n .ace_editor, .ace_editor *,\n .CodeMirror, .CodeMirror *,\n .monaco-editor, .monaco-editor *,\n .markdown-body pre, .markdown-body pre *,\n .highlight, .highlight *,\n pre code, pre code * {\n filter: none !important;\n }\n \n /* Fix common UI elements */\n .modal, .popup, .dropdown-menu, .tooltip, .popover {\n filter: invert(1) hue-rotate(180deg) !important;\n background: #2d2d44 !important;\n border-color: #444 !important;\n }\n \n /* Scrollbars */\n ::-webkit-scrollbar { background: #1a1a2e !important; }\n ::-webkit-scrollbar-thumb { background: #444 !important; }\n ::-webkit-scrollbar-thumb:hover { background: #555 !important; }\n \n /* Selection */\n ::selection { background: #4ecdc4 !important; color: #1a1a2e !important; }\n ::-moz-selection { background: #4ecdc4 !important; color: #1a1a2e !important; }\n ';\n }\n \n function removeDarkMode() {\n var style = document.getElementById('universal-dark-mode-style');\n if (style) style.remove();\n }\n \n // Toggle with Alt+Shift+D\n document.addEventListener('keydown', function(e) {\n if (e.altKey && e.shiftKey && e.key === 'D') {\n e.preventDefault();\n enabled = !enabled;\n if (enabled) {\n applyDarkMode();\n console.log('[Universal Dark Mode] Enabled');\n } else {\n removeDarkMode();\n console.log('[Universal Dark Mode] Disabled');\n }\n }\n });\n \n // Apply on load\n applyDarkMode();\n \n // Re-apply on dynamic content\n var observer = new MutationObserver(function(mutations) {\n if (enabled && !document.getElementById('universal-dark-mode-style')) {\n applyDarkMode();\n }\n });\n observer.observe(document.head, { childList: true });\n \n console.log('[Universal Dark Mode] Loaded - Press Alt+Shift+D to toggle');\n})();", "Universal Dark Mode");
}
} catch(__e) { console.warn('[Userscript:Universal Dark Mode]', __e); }
})();
})();

Skip to content
Add GeGLU and the corresponding gradient kernels#47

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

 Merged

timmoon10
merged 5 commits into

from

Jan 10, 2023
 Merged


Add GeGLU and the corresponding gradient kernels
#47
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


5 commits


Select commit
Hold shift + click to select a range
e20f806

Add GeGLU and DGeGLU


zlsh80826 Dec 27, 2022
441eff6

Add DGeGLUCT


zlsh80826 Dec 28, 2022
c029c09

Update copyright year


zlsh80826 Jan 6, 2023
1985c1d

Refine shape check


zlsh80826 Jan 6, 2023
6b1fe29

Code refine


zlsh80826 Jan 9, 2023

File filter

Filter by extension
Filter by extension
Conversations

Failed to load comments. 
Loading
Jump to

Jump to file


Failed to load files. 
Loading
 Diff view

 Diff view


3 changes: 3 additions & 0 deletions
tests/cpp/operator/CMakeLists.txt

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -8,7 +8,10 @@ add_executable(test_operator
 test_transpose.cu
 test_cast_transpose_dbias.cu
 test_cast_transpose_dbias_dgelu.cu
 test_cast_transpose_dgeglu.cu
 test_gelu.cu
 test_geglu.cu
 test_dgeglu.cu
 test_layernorm.cu
 test_multi_cast_transpose.cu
 ../test_common.cu)
Expand Down

146 changes: 146 additions & 0 deletions
tests/cpp/operator/test_cast_transpose_dgeglu.cu

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -0,0 +1,146 @@
/*************************************************************************
 * Copyright (c) 2022-2023, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 *
 * See LICENSE for license information.
 ************************************************************************/

#include <cuda_bf16.h>
#include <cuda_runtime.h>
#include <gtest/gtest.h>
#include <transformer_engine/logging.h>
#include <transformer_engine/transpose.h>
#include <cmath>
#include <cstring>
#include <iomanip>
#include <iostream>
#include <memory>
#include <random>
#include "../test_common.h"

using namespace transformer_engine;

namespace {

template <typename CType, typename IType>
inline CType gelu(const IType val) {
 CType cval = val;
 return cval * (0.5f + 0.5f * tanhf(cval * (0.79788456f + 0.03567741f * cval * cval)));
}

template <typename CType, typename IType>
inline CType dgelu(const IType val) {
 CType cval = val;
 const CType tanh_out = tanhf(0.79788456f * cval * (1.f + 0.044715f * cval * cval));
 return 0.5f * cval * ((1.f - tanh_out * tanh_out) * (0.79788456f + 0.1070322243f * cval * cval)) +
 0.5f * (1.f + tanh_out);
}

template <typename IT, typename OT, typename CT>
void compute_ref_cast_transpose_dgated_gelu(const IT *grad_h, const IT *input_h, const CT scale,
 OT *output_c_h, OT *output_t_h, CT *amax_h,
 const size_t N, const size_t H) {
 CT amax = 0.;

 const size_t col = H * 2;
 for (size_t i = 0; i < N; i++) {
 for (size_t j = 0; j < H; j++) {
 CT grad_elt = CT(grad_h[i * H + j]);
 CT gelu_elt = CT(input_h[i * col + j]);
 CT gate_elt = CT(input_h[i * col + H + j]);

 CT after_dgelu = dgelu<CT, CT>(gelu_elt) * grad_elt * gate_elt;
 CT after_dgate = grad_elt * gelu<CT, CT>(gelu_elt);

 amax = std::abs(after_dgelu) > amax ? std::abs(after_dgelu) : amax;
 amax = std::abs(after_dgate) > amax ? std::abs(after_dgate) : amax;

 output_c_h[i * col + j] = static_cast<OT>(scale * after_dgelu);
 output_c_h[i * col + H + j] = static_cast<OT>(scale * after_dgate);

 output_t_h[j * N + i] = static_cast<OT>(scale * after_dgelu);
 output_t_h[(j + H) * N + i] = static_cast<OT>(scale * after_dgate);
 }
 }

 *amax_h = amax;
}

template <typename IType, typename OType>
void performTest(const size_t N, const size_t H) {
 using namespace test;
 using CType = fp32;

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

 Tensor grad({N, H}, itype);
 Tensor input({N, H * 2}, itype);
 Tensor output_c({N, H * 2}, otype);
 Tensor output_t({H * 2, N}, otype);

 fillUniform(&grad);
 fillUniform(&input);
 setRandomScale(&output_c);
 output_t.shareFP8Meta(output_c);

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

 nvte_dgeglu_cast_transpose(grad.data(), input.data(), output_c.data(), output_t.data(), 0);

 CType ref_amax;
 compute_ref_cast_transpose_dgated_gelu(grad.cpu_dptr<IType>(), input.cpu_dptr<IType>(),
 output_c.scale(), ref_output_c.get(), ref_output_t.get(),
 &ref_amax, N, H);

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

 if (isFp8Type(otype)) {
 auto [atol_amax, rtol_amax] = getTolerances(DType::kFloat32);
 compareResults("amax", output_c.amax(), ref_amax, atol_amax, rtol_amax);
 float ref_scale_inv = 1.f / output_c.scale();
 compareResults("scale_inv", output_c.scale_inv(), ref_scale_inv, atol_amax, rtol_amax);
 }

 auto [atol, rtol] = getTolerances(otype);
 compareResults("output_c", output_c, ref_output_c.get(), atol, rtol);
 compareResults("output_t", output_t, ref_output_t.get(), atol, rtol);
}

std::vector<std::pair<size_t, size_t>> test_cases = {{64, 400}, {4096, 2048}, {768, 2816},
 {256, 5120}, {128, 10240}, {256, 256}};
Comment thread
zlsh80826 marked this conversation as resolved.


} // namespace

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

TEST_P(DGeGLUCTTestSuite, TestDGeGLUCT) {
 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, performTest<InputType, OutputType>(size.first, size.second);););
}

INSTANTIATE_TEST_SUITE_P(
 OperatorTest, DGeGLUCTTestSuite,
 ::testing::Combine(::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
 ::testing::Values(DType::kFloat8E5M2, DType::kFloat8E4M3),
 ::testing::ValuesIn(test_cases)),
 [](const testing::TestParamInfo<DGeGLUCTTestSuite::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;
 });

125 changes: 125 additions & 0 deletions
tests/cpp/operator/test_dgeglu.cu

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -0,0 +1,125 @@
/*************************************************************************
 * Copyright (c) 2022-2023, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 *
 * See LICENSE for license information.
 ************************************************************************/

#include <cuda_bf16.h>
#include <cuda_runtime.h>
#include <gtest/gtest.h>
#include <transformer_engine/activation.h>
#include <transformer_engine/logging.h>
#include <cmath>
#include <cstring>
#include <iomanip>
#include <iostream>
#include <memory>
#include <random>
#include <type_traits>
#include "../test_common.h"

using namespace transformer_engine;

namespace {

template <typename CType, typename IType>
inline CType gelu(const IType val) {
 CType cval = val;
 return cval * (0.5f + 0.5f * tanhf(cval * (0.79788456f + 0.03567741f * cval * cval)));
}

template <typename CType, typename IType>
inline CType dgelu(const IType val) {
 CType cval = val;
 const CType tanh_out = tanhf(0.79788456f * cval * (1.f + 0.044715f * cval * cval));
 return 0.5f * cval * ((1.f - tanh_out * tanh_out) * (0.79788456f + 0.1070322243f * cval * cval)) +
 0.5f * (1.f + tanh_out);
}

template <typename IT, typename OT, typename CT>
void compute_ref_dgeglu(const IT *grad_h, const IT *input_h, OT *output_h, const size_t N,
 const size_t H) {
 const size_t col = H * 2;

 for (size_t i = 0; i < N; i++) {
 for (size_t j = 0; j < H; j++) {
 CT grad_elt = CT(grad_h[i * H + j]);
 CT gelu_elt = CT(input_h[i * col + j]);
 CT gate_elt = CT(input_h[i * col + H + j]);

 CT after_dgelu = dgelu<CT, CT>(gelu_elt) * grad_elt * gate_elt;
 CT after_dgate = grad_elt * gelu<CT, CT>(gelu_elt);

 output_h[i * col + j] = OT(after_dgelu);
 output_h[i * col + H + j] = OT(after_dgate);
 }
 }
}

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

 using CType = fp32;

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

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

 fillUniform(&grad);
 fillUniform(&input);

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

 nvte_dgeglu(grad.data(), input.data(), output.data(), 0);

 compute_ref_dgeglu<IType, OType, CType>(grad.cpu_dptr<IType>(), input.cpu_dptr<IType>(),
 ref_output.get(), N, H);

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

 auto [atol, rtol] = getTolerances(otype);
 compareResults("output_dgelu", output, ref_output.get(), atol, rtol);
}

std::vector<std::pair<size_t, size_t>> test_cases = {
 {4096, 2048}, {768, 2816}, {256, 5120}, {128, 10240}, {256, 256}, {257, 259}, {128, 128 + 1}};

} // namespace

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

TEST_P(DGeGLUTestSuite, TestDGeGLU) {
 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,
 performTestDGeGLU<InputType, OutputType>(size.first, size.second);););
}

INSTANTIATE_TEST_SUITE_P(
 OperatorTest, DGeGLUTestSuite,
 ::testing::Combine(::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
 ::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16),
 ::testing::ValuesIn(test_cases)),
 [](const testing::TestParamInfo<DGeGLUTestSuite::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;
 });
Loading