Skip to content

[Clang][AMDGPU] Add __builtin_amdgcn_cvt_off_f32_i4 #133741

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

[Clang][AMDGPU] Add __builtin_amdgcn_cvt_off_f32_i4 #133741

Changes from 1 commit
Commits
Show all changes
10 commits
Select commit Hold shift + click to select a range
41af387
[Clang][AMDGPU] Add __builtin_amdgcn_cvt_off_f32_i4
jmmartinez Mar 27, 2025
dc78e8a
[Review] Turn builtin input from unsigned to signed
jmmartinez Apr 1, 2025
1cf0a9c
[Review] Added signed OpenCL CodeGen test
jmmartinez Apr 1, 2025
c10a091
[Review] Add line in the Release notes
jmmartinez Apr 1, 2025
73efe54
[Review] Move pattern into instruction declaration
jmmartinez Apr 1, 2025
89520ee
[Review] Handle undef & poison and update tests
jmmartinez Apr 1, 2025
eea660c
CI4BitAsInt -> CArg4BitAsInt
jmmartinez Apr 1, 2025
4b7a154
[Review] Forgot the constexpr case
jmmartinez Apr 1, 2025
57f299e
[Review] Tabulate the operation result to avoid operations that could…
jmmartinez Apr 1, 2025
e87dd39
[Review] static constexpr and mask
jmmartinez Apr 2, 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
10 changes: 7 additions & 3 deletions llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -743,9 +743,13 @@ GCNTTIImpl::instCombineIntrinsic(InstCombiner &IC, IntrinsicInst &II) const {
if (!CArg)
break;

int CArg4BitAsInt = CArg->getValue().trunc(4).getSExtValue();
float ResVal = 0.0625 * CArg4BitAsInt;
Constant *Res = ConstantFP::get(II.getType(), ResVal);
// Tabulated 0.0625 * (sext (CArg & 0xf)).
constexpr size_t ResValsSize = 16;
const float ResVals[ResValsSize] = {
0.0, 0.0625, 0.125, 0.1875, 0.25, 0.3125, 0.375, 0.4375,
-0.5, -0.4375, -0.375, -0.3125, -0.25, -0.1875, -0.125, -0.0625};
Constant *Res =
ConstantFP::get(Ty, ResVals[CArg->getZExtValue() % ResValsSize]);
return IC.replaceInstUsesWith(II, Res);
}
case Intrinsic::amdgcn_ubfe:
Expand Down
Loading