[SYCL][ESIMD] Support SYCL accessors in ESIMD mode in the FE. #1755

kbobrovs · 2020-05-26T03:41:33Z

REVIEW ONLY THE LAST COMMIT

FE detects whether an accessor kernel parameter is an accessor
passed to a ESIMD kernel), and marks the corrsponding integration
header parameter descriptor accordingly.
ESIMD accessors are handled differently at runtime (image1d_buffer_t is created).
Accompanying runtime changes are coming later.

Signed-off-by: Konstantin S Bobrovsky [email protected]

Options name: -fsycl-explicit-simd. This option is temporary until ESIMD and normal SYCL kernels can co-exist in the same source. Signed-off-by: Konstantin S Bobrovsky <[email protected]>

FE detects whether an accessor kernel parameter is an accessor passed to a ESIMD kernel), and marks the corrsponding integration header parameter descriptor accordingly. Signed-off-by: Konstantin S Bobrovsky <[email protected]>

Fznamznon · 2020-05-26T10:33:15Z

clang/lib/Sema/SemaSYCL.cpp

+  // Cooperate with the SYCL API, which defines this method iff the accessor is
+  // ESIMD accessor.
+  assert(Util::isSyclAccessorType(AccTy) && "sycl::accessor expected");
+  return getMethodByName(AccTy, "ESIMDBufferAccessorMarker") != nullptr;


I don't get why we have to distinguish ESIMD accessors from default ones if user cannot specify that accessor is ESIMD accessor and it is not possible to get a program with ESIMD and default accessors at the same time.

There will be corresponding runtime patch, which makes it clear why, + provides more detailed comments.

In short, for accessors in ESIMD kernels memory object creation is different - it is not clCreateBuffer

Okay, in this case there could be difference in process of setting kernel arguments, that is why you need to change descriptor in the integration header.
But I still don't get this searching for internal method. Why we don't threat all accessors as ESIMD ones when ESIMD extension is enabled?

Because some day we want to mix SIMD and normal kernels.

I'm a bit concerned about having additional handling for ESIMD accessors. I think, we were moving to unification of kernel arguments handling. This stuff just breaks this direction. Is it possible to elaborate what is the difference in RT handling between SIMD accessors and default ones?

I explained above:

In short, for accessors in ESIMD kernels memory object creation is different - it is not clCreateBuffer

that's the only difference.

CUDA and L0 back-ends doesn't use clCreateBuffer to allocate the memory and still they don't require special accessors. I don't follow why ESIMD kernels memory object creation requires compiler changes.

CUDA and L0 back-ends doesn't use clCreateBuffer to allocate the memory and still they don't require special accessors. I don't follow why ESIMD kernels memory object creation requires compiler changes.

This is a good example, which can help clarify.

CUDA and L0 are separate back-ends, as you correctly mentioned. What I need here is to distinguish usual accessors and ESIMD accessors within the same backend - be it L0 or OpenCL. I see that L0 is finally integrated into open-source, and I'll need to adjust the changes, but their sense will be the same - different sequence of PI calls will be needed to create memory object for the ESIMD accessor.

I could tag a kernel with "I am SIMD kernel" info, then diverge accessor handling based on that. But my estimation was that that would be much more intrusive and unsafe change to the runtime. I can add a TODO naming this alternative.

I explained my concerns above. We were moving to unification of kernel arguments generation, so we can improve detection and handling of special SYCL types (i.e. accessors, sampler and streams) by using an attribute. This approach was discussed during upstreaming WG meeting and we agreed that it is better than what we have now. Adding new accessor type doesn't seem like the change compatible with this direction, so I would prefer tagging kernels. Although I don't have enough expertise in RT to evaluate how intrusive and unsafe change to the runtime. @bader, WDYT?

erichkeane

The splitting of the isSyclType and matchQualifiedTypeName is odd, and likely unnecessary. Can you explain your motivation for doing that?

erichkeane · 2020-05-26T15:48:21Z

clang/lib/CodeGen/CGSYCLRuntime.cpp

@@ -20,6 +20,8 @@
 using namespace clang;
 using namespace CodeGen;

+namespace {


I get the purpose here, but this seems to be an unrelated change. Please submit it as a separate aptch.

The problem of absent anonymous namespace (and exported symbol names clash) is triggered by this patch. Separate patch would make it hard to understand the real purpose.

Besides, all patches that I've already submitted are ~5% of the overall line count of the feature. I simply can't split all tiny unrelated changes into separate patches given the time I have.

The schedule of your feature doesn't alter our code review guidelines.

Still this.

I'm not asking to change code review guidelines, I'm providing context. Guidelines are quite flexible, and they are just guidelines. There maybe other factors like business needs.

In this particular case, you probably missed the first part of my reasoning:

The problem of absent anonymous namespace (and exported symbol names clash) is triggered by this patch. Separate patch would make it hard to understand the real purpose.

what do you think?

I did not miss that.

erichkeane · 2020-05-26T15:48:30Z

clang/lib/Sema/SemaSYCL.cpp

@@ -54,6 +54,8 @@ enum KernelInvocationKind {
 const static std::string InitMethodName = "__init";
 const static std::string FinalizeMethodName = "__finalize";

+namespace {
+


Same as above

replied above

erichkeane · 2020-05-26T15:49:09Z

clang/lib/Sema/SemaSYCL.cpp

@@ -62,6 +64,7 @@ class Util {
  /// Checks whether given clang type is a full specialization of the SYCL
  /// accessor class.
  static bool isSyclAccessorType(const QualType &Ty);
+  static bool isSyclAccessorType(const CXXRecordDecl *RecTy);


Why are these overloads necessary? Is it not possible to just make the uses do RecordDecl->getType()?

This is another interface for convenience. Plus this avoids unnecessary dereferences, dyn_cast and branch:

bool Util::isSyclAccessorType(const QualType &Ty) { if (const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl()) return isSyclAccessorType(RecTy); return false; }

I think I'd prefer just having people write ->getType instead of this. You end up repeating all the record-decl detection code 1/2 dozen times.

This is matter of taste. If you insist I can change - ?

I think your change reduces readability, and is otherwise an unrelated change. @Fznamznon , thoughts?

I don't have strong preference regarding this, but for me it seems that there become a bit more code than it was before, so I don't mind reverting it.

@Fznamznon, do you mind leaving it as is?

erichkeane · 2020-05-26T15:50:17Z

clang/lib/Sema/SemaSYCL.cpp

@@ -2017,7 +2041,13 @@ SYCLIntegrationHeader::SYCLIntegrationHeader(DiagnosticsEngine &_Diag,
 // -----------------------------------------------------------------------------

 bool Util::isSyclAccessorType(const QualType &Ty) {
-  return isSyclType(Ty, "accessor", true /*Tmpl*/);
+  if (const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl())


This seems odd here. Why can't isSyclType handle this?

Signed-off-by: Konstantin S Bobrovsky <[email protected]>

kbobrovs · 2020-05-26T19:49:14Z

The splitting of the isSyclType and matchQualifiedTypeName is odd, and likely unnecessary. Can you explain your motivation for doing that?

between isSyclType and matchQualifiedTypeName there is a lot of boilerplate code, whose duplication is avoided by this splitting, thus improving maintainability.

erichkeane

I'll leave it to @Fznamznon to comment on the split of this but I'm concerned that we have two types of accessors that are so close to eachother. I cannot fathom why we cannot unify the handling of these two types better, so I'll leave it to her to close with.

erichkeane · 2020-05-27T00:17:47Z

clang/lib/CodeGen/CGSYCLRuntime.cpp

@@ -20,6 +20,8 @@
 using namespace clang;
 using namespace CodeGen;

+namespace {


Still this.

erichkeane · 2020-05-27T00:18:33Z

clang/lib/Sema/SemaSYCL.cpp

@@ -62,6 +64,7 @@ class Util {
  /// Checks whether given clang type is a full specialization of the SYCL
  /// accessor class.
  static bool isSyclAccessorType(const QualType &Ty);
+  static bool isSyclAccessorType(const CXXRecordDecl *RecTy);


I think I'd prefer just having people write ->getType instead of this. You end up repeating all the record-decl detection code 1/2 dozen times.

bader · 2020-05-28T15:46:10Z

@kbobrovs, could you elaborate in the description what problem this patch is trying to resolve?
Current description list what changes were done, but it seems that we are missing why these changes are needed.

kbobrovs · 2020-05-28T22:05:49Z

@kbobrovs, could you elaborate in the description what problem this patch is trying to resolve?
Current description list what changes were done, but it seems that we are missing why these changes are needed.

I provided a short description in the intro for this PR:

ESIMD accessors are handled differently at runtime (image1d_buffer_t is created).

This is the essence. Runtime creates different memory object so that explicit SIMD APIs work on SYCL accessors. The FE part if enabling the upcoming changes in runtime where all details should be clear.

erichkeane · 2020-06-02T15:26:59Z

clang/lib/Sema/SemaSYCL.cpp

 }

 bool Util::matchQualifiedTypeName(const QualType &Ty,
                                  ArrayRef<Util::DeclContextDesc> Scopes) {
+  if (const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl())


Instead of doing this 3 line pattern everywhere, make this whole body be:

return matchQualifiedTypeName(Ty->getAsCXXRecordDecl(), Scopes);

Then, in matchQualifiedType name, leave the
if (!RecTy) return false;

part.

kbobrovs · 2020-06-06T00:55:49Z

Will implement temporary runtime-only solution instead.

kbobrovs added 2 commits May 25, 2020 20:28

[SYCL][ESIMD] Add driver option to enable SYCL Explicit SIMD extension.

c258a7b

Options name: -fsycl-explicit-simd. This option is temporary until ESIMD and normal SYCL kernels can co-exist in the same source. Signed-off-by: Konstantin S Bobrovsky <[email protected]>

kbobrovs requested review from AGindinson, erichkeane, Fznamznon and mdtoguchi as code owners May 26, 2020 03:41

Fznamznon reviewed May 26, 2020

View reviewed changes

erichkeane reviewed May 26, 2020

View reviewed changes

[SQUASH] Address review comments.

0300ad8

Signed-off-by: Konstantin S Bobrovsky <[email protected]>

kbobrovs requested review from erichkeane and Fznamznon May 27, 2020 00:15

erichkeane reviewed May 27, 2020

View reviewed changes

kbobrovs requested a review from erichkeane May 27, 2020 00:43

kbobrovs requested a review from bader May 30, 2020 04:18

erichkeane reviewed Jun 2, 2020

View reviewed changes

kbobrovs closed this Jun 6, 2020

kbobrovs deleted the esimd-acc branch July 30, 2020 12:30

[SYCL][ESIMD] Support SYCL accessors in ESIMD mode in the FE. #1755

[SYCL][ESIMD] Support SYCL accessors in ESIMD mode in the FE. #1755

Uh oh!

Conversation

kbobrovs commented May 26, 2020

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

erichkeane left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

kbobrovs commented May 26, 2020

Uh oh!

erichkeane left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

bader commented May 28, 2020

Uh oh!

kbobrovs commented May 28, 2020

Uh oh!

Choose a reason for hiding this comment

Uh oh!

kbobrovs commented Jun 6, 2020

Uh oh!

Uh oh!