Skip to content

[SYCL] Implement anonymous namespace spec-id functionality. #3576

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.

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 5 commits into from
Apr 22, 2021

Conversation

erichkeane
Copy link
Contributor

In order to get the specialization_id metadata to be usable in the
integration footer, we need to insert a 'shim' function at each
anonymous namespace so that we can do a lookup at the 'right' location,
then look it up again later by the type.

This implements all that functionality based on my latest understanding
of the design.

In order to get the specialization_id metadata to be usable in the
integration footer, we need to insert a 'shim' function at each
anonymous namespace so that we can do a lookup at the 'right' location,
then look it up again later by the type.

This implements all that functionality based on my latest understanding
of the design.
Copy link
Contributor

@premanandrao premanandrao left a comment

Choose a reason for hiding this comment

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

Thanks @erichkeane. Looks mostly okay to me; some nits:

// Emit the list of shims required for a DeclContext, calls itself recursively.
static std::string EmitSpecIdShims(raw_ostream &O, unsigned &ShimCounter,
const DeclContext *DC,
std::string NameForLastShim) {
Copy link
Contributor

Choose a reason for hiding this comment

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

Can NameForLastShim be a reference?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

The intent was for it to be a 'local variable' here so that it could modify it without affecting the parent. That said, it seems that it only does tail-call, so I believe it can. We just need to make sure with the change that we never allow anything that modifies this variable to be below the recursion.

constexpr sycl::specialization_id same_name{7};
// CHECK-NEXT: namespace {
// CHECK-NEXT: namespace __sycl_detail {
// CHECK-NEXT: static constexpr decltype(same_name) &__spec_id_shim_[[SHIM_ID:[0-9]+]]() {
Copy link
Contributor

Choose a reason for hiding this comment

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

Is it worth it to test that this SHIM_ID and others below are a different value than the previous ones?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

I'd thought about that, but I'm not aware of a good way to make them all be wildcards, yet still make sure they are different. Unfortunately there doesn't seem to be any sort of mechanism in CHECK to do something like that.

Do you have a suggestion?

Copy link
Contributor

Choose a reason for hiding this comment

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

Sorry, no, I don't have a suggestion. It was just a thought.

Copy link
Contributor

@AaronBallman AaronBallman left a comment

Choose a reason for hiding this comment

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

LGTM

@bader bader requested a review from premanandrao April 22, 2021 05:59
constexpr sycl::specialization_id same_name{7};
// CHECK-NEXT: namespace {
// CHECK-NEXT: namespace __sycl_detail {
// CHECK-NEXT: static constexpr decltype(same_name) &__spec_id_shim_[[SHIM_ID:[0-9]+]]() {
Copy link
Contributor

Choose a reason for hiding this comment

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

Sorry, no, I don't have a suggestion. It was just a thought.

@bader bader merged commit d482ac3 into intel:sycl Apr 22, 2021
Copy link
Contributor

@AlexeySachkov AlexeySachkov left a comment

Choose a reason for hiding this comment

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

Sorry for such late review, but there seems to be a bug in the patch

// CHECK-NEXT: namespace detail {
// CHECK-NEXT: template<>
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
// CHECK-NEXT: return ::__sycl_detail::__spec_id_shim_[[SHIM_ID]]();
Copy link
Contributor

Choose a reason for hiding this comment

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

It doesn't seem like this code will compile: we return a pointer to a specialization_id variable from a function which is expected to return const char *. symbolic_ID should be return there

OS << "template<>\n";
OS << "inline const char *get_spec_constant_symbolic_ID<" << TopShim
<< ">() {\n";
OS << " return " << TopShim << ";\n";
Copy link
Contributor

Choose a reason for hiding this comment

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

I guess here we should have emitSpecIDName instead of TopShim, see comment below in test

Copy link
Contributor Author

Choose a reason for hiding this comment

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

its not the spec-id name, because of course we cannot name that.

Sorry for such late review, but there seems to be a bug in the patch

#3669

erichkeane pushed a commit to erichkeane/llvm that referenced this pull request Apr 30, 2021
Looks like I got carried away with the implementation and we didn't
catch that I was returning the wrong things from the 'end call function'
when generating anonymous function shims.  This was brought up here:
intel#3576 (review)

This should fix that problem.
romanovvlad pushed a commit that referenced this pull request May 4, 2021
…3669)

Looks like I got carried away with the implementation and we didn't
catch that I was returning the wrong things from the 'end call function'
when generating anonymous function shims.  This was brought up here:
#3576 (review)

This should fix that problem.
VD = VD->getCanonicalDecl();
if (VD->isInAnonymousNamespace()) {
std::string TopShim = EmitSpecIdShims(OS, ShimCounter, VD);
OS << "namespace sycl {\n";
Copy link
Contributor

Choose a reason for hiding this comment

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

Shouldn't we insert __SYCL_INLINE_NAMESPACE(cl) like it is done for integration header on line 4330?
@erichkeane

Copy link
Contributor Author

Choose a reason for hiding this comment

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

I don't know, that wasn't a part of the spec I was given.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

See: #3752. @pglowney for your attention as well.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

6 participants