Skip to content
New issue

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

Sign up for GitHub

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

Already on GitHub? Sign in to your account

[pre-commit.ci] pre-commit suggestions #979

Draft
wants to merge 2 commits into
base: main
Choose a base branch
from
Draft

[pre-commit.ci] pre-commit suggestions #979

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
b682881
[pre-commit.ci] pre-commit suggestions
pre-commit-ci[bot] Jan 7, 2025
f0138bf
[pre-commit.ci] auto fixes from pre-commit.com hooks
pre-commit-ci[bot] Jan 7, 2025
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
6 changes: 3 additions & 3 deletions .pre-commit-config.yaml
100755 → 100644
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,7 +14,7 @@ ci:

repos:
- repo: https://github.com/pre-commit/pre-commit-hooks
rev: v4.6.0
rev: v5.0.0
hooks:
- id: check-merge-conflict
- id: check-added-large-files
Expand All @@ -24,15 +24,15 @@ repos:
files: .*.(c|cc|cxx|cpp|cu|cuh|h|hpp|py)$

- repo: https://github.com/psf/black
rev: 24.4.2
rev: 24.10.0
hooks:
- id: black
name: Format python code
args: [--line-length=100, --preview, --enable-unstable-feature=string_processing]
types: [python]

- repo: https://github.com/pre-commit/mirrors-clang-format
rev: v18.1.6
rev: v19.1.6
hooks:
- id: clang-format
entry: clang-format -i
Expand Down
72 changes: 54 additions & 18 deletions transformer_engine/common/common.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -120,35 +120,51 @@ struct TypeInfo {
using namespace transformer_engine; \
case DType::kByte: { \
using type = unsigned char; \
{ __VA_ARGS__ } \
{ \
__VA_ARGS__ \
} \
} break; \
case DType::kInt32: { \
using type = int32_t; \
{ __VA_ARGS__ } \
{ \
__VA_ARGS__ \
} \
} break; \
case DType::kInt64: { \
using type = int64_t; \
{ __VA_ARGS__ } \
{ \
__VA_ARGS__ \
} \
} break; \
case DType::kFloat32: { \
using type = float; \
{ __VA_ARGS__ } \
{ \
__VA_ARGS__ \
} \
} break; \
case DType::kFloat16: { \
using type = fp16; \
{ __VA_ARGS__ } \
{ \
__VA_ARGS__ \
} \
} break; \
case DType::kBFloat16: { \
using type = bf16; \
{ __VA_ARGS__ } \
{ \
__VA_ARGS__ \
} \
} break; \
case DType::kFloat8E4M3: { \
using type = fp8e4m3; \
{ __VA_ARGS__ } \
{ \
__VA_ARGS__ \
} \
} break; \
case DType::kFloat8E5M2: { \
using type = fp8e5m2; \
{ __VA_ARGS__ } \
{ \
__VA_ARGS__ \
} \
} break; \
default: \
NVTE_ERROR("Invalid type."); \
Expand All @@ -159,23 +175,33 @@ struct TypeInfo {
using namespace transformer_engine; \
case DType::kFloat32: { \
using type = float; \
{ __VA_ARGS__ } \
{ \
__VA_ARGS__ \
} \
} break; \
case DType::kFloat16: { \
using type = fp16; \
{ __VA_ARGS__ } \
{ \
__VA_ARGS__ \
} \
} break; \
case DType::kBFloat16: { \
using type = bf16; \
{ __VA_ARGS__ } \
{ \
__VA_ARGS__ \
} \
} break; \
case DType::kFloat8E5M2: { \
using type = fp8e5m2; \
{ __VA_ARGS__ } \
{ \
__VA_ARGS__ \
} \
} break; \
case DType::kFloat8E4M3: { \
using type = fp8e4m3; \
{ __VA_ARGS__ } \
{ \
__VA_ARGS__ \
} \
} break; \
default: \
NVTE_ERROR("Invalid type."); \
Expand All @@ -186,11 +212,15 @@ struct TypeInfo {
using namespace transformer_engine; \
case DType::kFloat8E5M2: { \
using type = fp8e5m2; \
{ __VA_ARGS__ } \
{ \
__VA_ARGS__ \
} \
} break; \
case DType::kFloat8E4M3: { \
using type = fp8e4m3; \
{ __VA_ARGS__ } \
{ \
__VA_ARGS__ \
} \
} break; \
default: \
NVTE_ERROR("Invalid type."); \
Expand All @@ -201,15 +231,21 @@ struct TypeInfo {
using namespace transformer_engine; \
case DType::kFloat32: { \
using type = float; \
{ __VA_ARGS__ } \
{ \
__VA_ARGS__ \
} \
} break; \
case DType::kFloat16: { \
using type = fp16; \
{ __VA_ARGS__ } \
{ \
__VA_ARGS__ \
} \
} break; \
case DType::kBFloat16: { \
using type = bf16; \
{ __VA_ARGS__ } \
{ \
__VA_ARGS__ \
} \
} break; \
case DType::kFloat8E5M2: \
case DType::kFloat8E4M3: { \
Expand Down
4 changes: 2 additions & 2 deletions transformer_engine/pytorch/csrc/type_shim.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -292,7 +292,7 @@ reduce_block_into_lanes(T *x, T val, int lanes = 1,
final = x[tid] + x[tid + 32];
else
final = val;
// __SYNCWARP();
// __SYNCWARP();

#pragma unroll
for (int i = 16; i >= lanes; i >>= 1) final = final + __shfl_down_sync(0xffffffff, final, i);
Expand Down Expand Up @@ -333,7 +333,7 @@ reduce_block_into_lanes_max_op(T *x, T val, int lanes = 1,
final = fmaxf(fabsf(x[tid]), fabsf(x[tid + 32]));
else
final = val;
// __SYNCWARP();
// __SYNCWARP();

#pragma unroll
for (int i = 16; i >= lanes; i >>= 1)
Expand Down
Loading