Skip to content

[NVPTX] Cleanup and document nvvm.fabs intrinsics, adding f16 support #135644

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
merged 7 commits into from
Apr 17, 2025
Merged

[NVPTX] Cleanup and document nvvm.fabs intrinsics, adding f16 support #135644

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
fd11c2b
[NVPTX] Cleaup and document nvvm.fabs intrinsics, adding f16 support
AlexMaclean Apr 11, 2025
7d87fb5
fixup
AlexMaclean Apr 15, 2025
e0d28ec
address comments
AlexMaclean Apr 15, 2025
1bfd50d
fixup clang
AlexMaclean Apr 15, 2025
51f32c9
fixup
AlexMaclean Apr 15, 2025
391bfad
add f16 builtins
AlexMaclean Apr 15, 2025
fc08be3
fixup test
AlexMaclean Apr 16, 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
4 changes: 4 additions & 0 deletions clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1034,6 +1034,10 @@ Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID,
case NVPTX::BI__nvvm_fmin_xorsign_abs_f16x2:
return MakeHalfType(Intrinsic::nvvm_fmin_xorsign_abs_f16x2, BuiltinID, E,
*this);
case NVPTX::BI__nvvm_abs_bf16:
case NVPTX::BI__nvvm_abs_bf16x2:
return Builder.CreateUnaryIntrinsic(Intrinsic::nvvm_fabs,
EmitScalarExpr(E->getArg(0)));
case NVPTX::BI__nvvm_ldg_h:
case NVPTX::BI__nvvm_ldg_h2:
return MakeHalfType(Intrinsic::not_intrinsic, BuiltinID, E, *this);
Expand Down
4 changes: 2 additions & 2 deletions clang/test/CodeGen/builtins-nvptx.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1038,9 +1038,9 @@ __device__ void nvvm_cvt_sm89() {
__device__ void nvvm_abs_neg_bf16_bf16x2_sm80() {
#if __CUDA_ARCH__ >= 800

// CHECK_PTX70_SM80: call bfloat @llvm.nvvm.abs.bf16(bfloat 0xR3DCD)
// CHECK_PTX70_SM80: call bfloat @llvm.nvvm.fabs.bf16(bfloat 0xR3DCD)
__nvvm_abs_bf16(BF16);
// CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.abs.bf16x2(<2 x bfloat> splat (bfloat 0xR3DCD))
// CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.fabs.v2bf16(<2 x bfloat> splat (bfloat 0xR3DCD))
__nvvm_abs_bf16x2(BF16X2);

// CHECK_PTX70_SM80: call bfloat @llvm.nvvm.neg.bf16(bfloat 0xR3DCD)
Expand Down
Loading