[SYCL] Save user specified names in lambda object

[elizabethandrews]

This patch retains user specified names for lambda captures in
lambda object. As a result, the name of openCL kernel arguments
generated for SYCL kernel specified as a lamdba, now includes the
user names in kernel argument name (matches current behavior for
SYCL kernel specified as a functor object).

Signed-off-by: Elizabeth Andrews elizabeth.andrews@intel.com

[elizabethandrews]

I am also working on a patch to upstream this support to LLVM project. I'm hoping to put up that patch for review this week. I will share the link here once I put it up. I'm not sure if they will accept the change since there is no strong justification for this other than "it improves SYCL".

Since there is some urgency for this patch for DPCPP, I am hoping to get the ball rolling here in the meantime.

[elizabethandrews]

I don't have the environment to test CUDA fails. The failures however seem to be due to IR changes because of name change. I've fixed what I saw in failing tests, but could not verify fix locally. Hopefully I've caught all the issues.

[alexbatashev]

Runtime LGTM

[premanandrao]

LGTM

[keryell]

Cool feature for the debug!

[elizabethandrews]

@bader I don't believe the failures are a result of this patch. This is the 3rd run with the same code (first run had missed a few CUDA tests which I now fixed by changing tests, second run had one flaky test which longer fails) and this is third run which has new fails which did not show up in previous runs

[keryell]

Just rethinking last night about it, what happens it you have a non qualified name which appears twice, so you have a name conflict?
For example you capture a::c and b::c, with c are USM pointers defined in 2 different namespaces or 2 accessors defined in 2 different nested structs inside the submit around the lambda?

[elizabethandrews]

Just rethinking last night about it, what happens it you have a non qualified name which appears twice, so you have a name conflict? For example you capture a::c and b::c, with c are USM pointers defined in 2 different namespaces or 2 accessors defined in 2 different nested structs inside the submit around the lambda?

I believe you will have the same name. This isn't new behavior though. Accessors defined in nested structs already had name retained in kernel lambda object. For example -

Code:

struct AccessorNested {
  int A, B;
  sycl::accessor<char, 1, sycl::access::mode::read> AccField;
};

struct AccessorOuter {
  sycl::accessor<char, 1, sycl::access::mode::read> AccField;
  AccessorNested obj;
  int C;
};

int main() {
  AccessorOuter Object;
  myQueue.submit([&](sycl::handler &h) {
    h.single_task<class kernel>(
        [=] {
          Object.use();
        });
  });
  return 0;
}

AST for kernel parameters -

|-ParmVarDecl 0xa127378 <<invalid sloc>> <invalid sloc> used _arg_AccField '__global char *'
  | |-SYCLAccessorReadonlyAttr 0xa1273e0 <<invalid sloc>> Implicit
  | `-SYCLAccessorPtrAttr 0xa127438 <<invalid sloc>> Implicit
  |-ParmVarDecl 0xa127470 <<invalid sloc>> <invalid sloc> used _arg_AccField 'sycl::range<1>'
  |-ParmVarDecl 0xa1274e8 <<invalid sloc>> <invalid sloc> used _arg_AccField 'sycl::range<1>'
  |-ParmVarDecl 0xa127560 <<invalid sloc>> <invalid sloc> used _arg_AccField 'sycl::id<1>'
  |-ParmVarDecl 0xa127dc0 <<invalid sloc>> <invalid sloc> used _arg_A 'int'
  |-ParmVarDecl 0xa127e78 <<invalid sloc>> <invalid sloc> used _arg_B 'int'
  |-ParmVarDecl 0xa127f30 <<invalid sloc>> <invalid sloc> used _arg_AccField '__global char *'
  | |-SYCLAccessorReadonlyAttr 0xa127f98 <<invalid sloc>> Implicit
  | `-SYCLAccessorPtrAttr 0xa127ff0 <<invalid sloc>> Implicit
  |-ParmVarDecl 0xa128028 <<invalid sloc>> <invalid sloc> used _arg_AccField 'sycl::range<1>'
  |-ParmVarDecl 0xa1280a0 <<invalid sloc>> <invalid sloc> used _arg_AccField 'sycl::range<1>'
  |-ParmVarDecl 0xa128118 <<invalid sloc>> <invalid sloc> used _arg_AccField 'sycl::id<1>'
  |-ParmVarDecl 0xa128420 <<invalid sloc>> <invalid sloc> used _arg_C 'int'

This is existing behavior though. Variable names nested in some class were available prior to this patch as well. Only names of direct captures were missing - since clang didn't save these in lambda object. This patch just added that. There might be value in using qualified names instead.

[againull]

@elizabethandrews I am not sure I fully understand this PR but I have a question.

Value names are discarded by default for release build of the Clang (no assertions), please see:
https://clang.llvm.org/docs/UsersManual.html#controlling-value-names-in-llvm-ir

And to disable this behaviour the following option can be used:
-fno-discard-value-names

If we want to preserve value names for DPCPP then shouldn't we enforce -fno-discard-value-names in driver if -fsycl is used?

[elizabethandrews]

@elizabethandrews I am not sure I fully understand this PR but I have a question.

Value names are discarded by default for release build of the Clang (no assertions), please see: https://clang.llvm.org/docs/UsersManual.html#controlling-value-names-in-llvm-ir

And to disable this behaviour the following option can be used: -fno-discard-value-names

If we want to preserve value names for DPCPP then shouldn't we enforce -fno-discard-value-names in driver if -fsycl is used?

@againull this PR wasn't about preserving names for llvm::values. This PR preserves names when generating an openCL kernel (from SYCL kernel). We were not saving user-specified names in some cases resulting in generic names in LLVM IR for openCL kernel arguments, irrespective of options used.

[keryell]

Would it make sense to save a fully qualified mangled name then instead of the short one?

[elizabethandrews]

Would it make sense to save a fully qualified mangled name then instead of the short one?

From a functional perspective I don't think it makes a difference since same names don't seem to bother backend. Name generated in IR will be unique. For debugging purposes, if we have names which repeat within a kernel maybe fully qualified names helps? I'm not sure if it will make the IR verbose though.