Skip to content

[AMDGPU] Treat printf as builtin for OpenCL #72554

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

Closed
wants to merge 1 commit into from
Closed

[AMDGPU] Treat printf as builtin for OpenCL #72554

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
8 changes: 8 additions & 0 deletions clang/include/clang/Basic/BuiltinsAMDGPU.def
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,6 +21,10 @@
#if defined(BUILTIN) && !defined(TARGET_BUILTIN)
# define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE) BUILTIN(ID, TYPE, ATTRS)
#endif

#if defined(BUILTIN) && !defined(LANGBUILTIN)
#define LANGBUILTIN(ID, TYPE, ATTRS, BUILTIN_LANG) BUILTIN(ID, TYPE, ATTRS)
#endif
//===----------------------------------------------------------------------===//
// SI+ only builtins.
//===----------------------------------------------------------------------===//
Expand Down Expand Up @@ -406,5 +410,9 @@ TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_fp8_f32, "iffiIb", "nc", "fp8-insts")
TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_bf8_f32, "ifiiIi", "nc", "fp8-insts")
TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_fp8_f32, "ifiiIi", "nc", "fp8-insts")

// OpenCL
LANGBUILTIN(printf, "icC*4.", "fp:0:", ALL_OCL_LANGUAGES)
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Why do you need to define a new target builtin, just to hack it to the generic lang builtin later? Just handle the existing printf builtin?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is specifically to recognize the OpenCL version of printf (where fmt string arg is a pointer to const address space) as a builtin. The hack to generic builtin is just a option that I had as I did not want to add a new case to builtin expansion code (since the API used by both OpenCL and HIP are same ), however Im okay with adding a new case too if you feel it makes more sense.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I still don't see why this is necessary. A target-defined language-specific builtin is a whole new beast. What is missing in the current parsing of OpenCL printf?

Copy link
Contributor Author

@vikramRH vikramRH Nov 24, 2023 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@ssahasra , I still feel this is the way to move here since I dont see a way to access the printf option at IR level (i.e during optimization pipeline) and thus decide version of printf to use. It has to be at clang CodeGen. I ask other reviewers too if they feel there are major concerns with adding such a builtin variant (i.e AMDGPU and OCL specific). I might have to look for alternative approaches if so.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

If you're handling the builtin in clang directly, you can go off the original Builtin::BIprintf. I don't see what the alias AMDGPU::BIprintf is doing.

Copy link
Contributor Author

@vikramRH vikramRH Dec 1, 2023 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

OpenCL spec says printf format string should be in constant address space. This makes the OCL printf signature target specific and hence we would need a target specific builtinID to recognize this. Im not sure I understand how we can go ahead with generic "BIPrinf" here ?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The address space makes the type language dependent, it does not make it target dependent

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think what @vikramRH is saying is that the magic number "4" for OpenCL address space "__constant" is specific to AMDGPU.


#undef BUILTIN
#undef TARGET_BUILTIN
#undef LANGBUILTIN
7 changes: 7 additions & 0 deletions clang/lib/AST/Decl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -49,6 +49,7 @@
#include "clang/Basic/SourceLocation.h"
#include "clang/Basic/SourceManager.h"
#include "clang/Basic/Specifiers.h"
#include "clang/Basic/TargetBuiltins.h"
#include "clang/Basic/TargetCXXABI.h"
#include "clang/Basic/TargetInfo.h"
#include "clang/Basic/Visibility.h"
Expand Down Expand Up @@ -3598,6 +3599,12 @@ unsigned FunctionDecl::getBuiltinID(bool ConsiderWrapperFunctions) const {
if (!ConsiderWrapperFunctions && getStorageClass() == SC_Static)
return 0;

// AMDGCN implementation supports printf as a builtin
// for OpenCL
if (Context.getTargetInfo().getTriple().isAMDGCN() &&
Context.getLangOpts().OpenCL && BuiltinID == AMDGPU::BIprintf)
return BuiltinID;

// OpenCL v1.2 s6.9.f - The library functions defined in
// the C99 standard headers are not available.
if (Context.getLangOpts().OpenCL &&
Expand Down
3 changes: 3 additions & 0 deletions clang/lib/Basic/Targets/AMDGPU.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -91,6 +91,9 @@ static constexpr Builtin::Info BuiltinInfo[] = {
{#ID, TYPE, ATTRS, nullptr, HeaderDesc::NO_HEADER, ALL_LANGUAGES},
#define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE) \
{#ID, TYPE, ATTRS, FEATURE, HeaderDesc::NO_HEADER, ALL_LANGUAGES},
#define LANGBUILTIN(ID, TYPE, ATTRS, LANG) \
{ #ID, TYPE, ATTRS, nullptr, HeaderDesc::NO_HEADER, LANG } \
,
#include "clang/Basic/BuiltinsAMDGPU.def"
};

Expand Down
5 changes: 5 additions & 0 deletions clang/lib/CodeGen/CGBuiltin.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2458,6 +2458,11 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
&getTarget().getLongDoubleFormat() == &llvm::APFloat::IEEEquad())
BuiltinID = mutateLongDoubleBuiltin(BuiltinID);

// Mutate the printf builtin ID so that we use the same CodeGen path for
// HIP and OpenCL with AMDGPU targets.
if (getTarget().getTriple().isAMDGCN() && BuiltinID == AMDGPU::BIprintf)
BuiltinID = Builtin::BIprintf;

// If the builtin has been declared explicitly with an assembler label,
// disable the specialized emitting below. Ideally we should communicate the
// rename in IR, or at least avoid generating the intrinsic calls that are
Expand Down