[PATCH] D69938: [OpenCL] Use __generic addr space when generating internal representation of lambda

[Anastasia Stulova via Phabricator via cfe-commits]

Anastasia added a comment.

In D69938#1747735 <https://reviews.llvm.org/D69938#1747735>, @rjmccall wrote:

> > I was already planning to add this. I could look into it next and maybe just a add FIXME in the test for now.
>
> Sure.
>
> > Btw global lambda objects are in `__global` address space in OpenCL.
>
> Just based on being written outside of a function body or in a static-local initializer?  I guess it's supportable; you can reasonably allocate global memory for those temporaries.


If I compile this kernel:

  auto glambda = [](auto a)  { return 1; };
  
  __kernel void foo() {
    glambda(1);
  }

I will get the following IR:

  %class.anon = type { i8 }
  
  @glambda = internal addrspace(1) global %class.anon undef, align 1
  
  ; Function Attrs: convergent noinline nounwind optnone
  define spir_kernel void @foo() #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !3 !kernel_arg_type !3 !kernel_arg_base_type !3 !kernel_arg_type_qual !3 {
  entry:
    %call = call spir_func i32 @"_ZNU3AS43$_0clIiEEDaT_"(%class.anon addrspace(4)* addrspacecast (%class.anon addrspace(1)* @glambda to %class.anon addrspace(4)*), i32 1) #2
    ret void
  }
  
  ; Function Attrs: convergent noinline nounwind optnone
  define internal spir_func i32 @"_ZNU3AS43$_0clIiEEDaT_"(%class.anon addrspace(4)* %this, i32 %a) #1 align 2 {
  entry:
    %this.addr = alloca %class.anon addrspace(4)*, align 4
    %a.addr = alloca i32, align 4
    store %class.anon addrspace(4)* %this, %class.anon addrspace(4)** %this.addr, align 4
    store i32 %a, i32* %a.addr, align 4
    %this1 = load %class.anon addrspace(4)*, %class.anon addrspace(4)** %this.addr, align 4
    ret i32 1
  }

and the AST is:

  |-VarDecl 0x55e8d2789500 <test.cl:1:1, col:39> col:6 used glambda '__global (lambda at test.cl:1:16)':'__global (lambda at test.cl:1:16)' cinit
  | `-ImplicitCastExpr 0x55e8d278a1f0 <col:16, col:39> '__global (lambda at test.cl:1:16)':'__global (lambda at test.cl:1:16)' <NoOp>
  |   `-LambdaExpr 0x55e8d2789f90 <col:16, col:39> '(lambda at test.cl:1:16)'
  |     |-CXXRecordDecl 0x55e8d2789808 <col:16> col:16 implicit class definition
  |     | |-DefinitionData generic lambda pass_in_registers empty standard_layout trivially_copyable literal can_const_default_init
  |     | | |-DefaultConstructor defaulted_is_constexpr
  |     | | |-CopyConstructor simple trivial has_const_param needs_implicit implicit_has_const_param
  |     | | |-MoveConstructor exists simple trivial needs_implicit
  |     | | |-CopyAssignment trivial has_const_param needs_implicit implicit_has_const_param
  |     | | |-MoveAssignment
  |     | | `-Destructor simple irrelevant trivial
  |     | |-FunctionTemplateDecl 0x55e8d2789aa8 <<invalid sloc>, col:39> col:16 operator()
  |     | | |-TemplateTypeParmDecl 0x55e8d27895d8 <col:19> col:19 implicit class depth 0 index 0
  |     | | |-CXXMethodDecl 0x55e8d27899f8 <col:25, col:39> col:16 constexpr operator() 'auto (auto) const __generic' inline
  |     | | | |-ParmVarDecl 0x55e8d27896a8 <col:19, col:24> col:24 a 'auto'
  |     | | | `-CompoundStmt 0x55e8d2789b48 <col:27, col:39>
  |     | | |   `-ReturnStmt 0x55e8d2789b38 <col:29, col:36>
  |     | | |     `-IntegerLiteral 0x55e8d2789b18 <col:36> 'int' 1
  |     | | `-CXXMethodDecl 0x55e8d27b9750 <col:25, col:39> col:16 used constexpr operator() 'int (int) const __generic' inline
  |     | |   |-TemplateArgument type 'int'
  |     | |   |-ParmVarDecl 0x55e8d27b9608 <col:19, col:24> col:24 a 'int':'int'
  |     | |   `-CompoundStmt 0x55e8d27b9978 <col:27, col:39>
  |     | |     `-ReturnStmt 0x55e8d27b9968 <col:29, col:36>
  |     | |       `-IntegerLiteral 0x55e8d2789b18 <col:36> 'int' 1
  |     | |-FunctionTemplateDecl 0x55e8d2789e18 <<invalid sloc>, col:39> col:16 implicit operator auto (*)(type-parameter-0-0)
  |     | | |-TemplateTypeParmDecl 0x55e8d27895d8 <col:19> col:19 implicit class depth 0 index 0
  |     | | `-CXXConversionDecl 0x55e8d2789d70 <col:16, col:39> col:16 implicit constexpr operator auto (*)(type-parameter-0-0) 'auto (*() const noexcept)(auto)' inline
  |     | |-FunctionTemplateDecl 0x55e8d2789f20 <<invalid sloc>, col:39> col:16 implicit __invoke
  |     | | |-TemplateTypeParmDecl 0x55e8d27895d8 <col:19> col:19 implicit class depth 0 index 0
  |     | | `-CXXMethodDecl 0x55e8d2789e78 <col:16, col:39> col:16 implicit __invoke 'auto (auto)' static inline
  |     | |   `-ParmVarDecl 0x55e8d2789d08 <col:19, col:24> col:24 a 'auto'
  |     | `-CXXDestructorDecl 0x55e8d2789fc0 <col:16> col:16 implicit referenced ~ 'void () __generic noexcept' inline default trivial
  |     `-CompoundStmt 0x55e8d2789b48 <col:27, col:39>
  |       `-ReturnStmt 0x55e8d2789b38 <col:29, col:36>
  |         `-IntegerLiteral 0x55e8d2789b18 <col:36> 'int' 1
  |-CXXRecordDecl 0x55e8d2789808 <col:16> col:16 implicit class definition
  | |-DefinitionData generic lambda pass_in_registers empty standard_layout trivially_copyable literal can_const_default_init
  | | |-DefaultConstructor defaulted_is_constexpr
  | | |-CopyConstructor simple trivial has_const_param needs_implicit implicit_has_const_param
  | | |-MoveConstructor exists simple trivial needs_implicit
  | | |-CopyAssignment trivial has_const_param needs_implicit implicit_has_const_param
  | | |-MoveAssignment
  | | `-Destructor simple irrelevant trivial
  | |-FunctionTemplateDecl 0x55e8d2789aa8 <<invalid sloc>, col:39> col:16 operator()
  | | |-TemplateTypeParmDecl 0x55e8d27895d8 <col:19> col:19 implicit class depth 0 index 0
  | | |-CXXMethodDecl 0x55e8d27899f8 <col:25, col:39> col:16 constexpr operator() 'auto (auto) const __generic' inline
  | | | |-ParmVarDecl 0x55e8d27896a8 <col:19, col:24> col:24 a 'auto'
  | | | `-CompoundStmt 0x55e8d2789b48 <col:27, col:39>
  | | |   `-ReturnStmt 0x55e8d2789b38 <col:29, col:36>
  | | |     `-IntegerLiteral 0x55e8d2789b18 <col:36> 'int' 1
  | | `-CXXMethodDecl 0x55e8d27b9750 <col:25, col:39> col:16 used constexpr operator() 'int (int) const __generic' inline
  | |   |-TemplateArgument type 'int'
  | |   |-ParmVarDecl 0x55e8d27b9608 <col:19, col:24> col:24 a 'int':'int'
  | |   `-CompoundStmt 0x55e8d27b9978 <col:27, col:39>
  | |     `-ReturnStmt 0x55e8d27b9968 <col:29, col:36>
  | |       `-IntegerLiteral 0x55e8d2789b18 <col:36> 'int' 1
  | |-FunctionTemplateDecl 0x55e8d2789e18 <<invalid sloc>, col:39> col:16 implicit operator auto (*)(type-parameter-0-0)
  | | |-TemplateTypeParmDecl 0x55e8d27895d8 <col:19> col:19 implicit class depth 0 index 0
  | | `-CXXConversionDecl 0x55e8d2789d70 <col:16, col:39> col:16 implicit constexpr operator auto (*)(type-parameter-0-0) 'auto (*() const noexcept)(auto)' inline
  | |-FunctionTemplateDecl 0x55e8d2789f20 <<invalid sloc>, col:39> col:16 implicit __invoke
  | | |-TemplateTypeParmDecl 0x55e8d27895d8 <col:19> col:19 implicit class depth 0 index 0
  | | `-CXXMethodDecl 0x55e8d2789e78 <col:16, col:39> col:16 implicit __invoke 'auto (auto)' static inline
  | |   `-ParmVarDecl 0x55e8d2789d08 <col:19, col:24> col:24 a 'auto'
  | `-CXXDestructorDecl 0x55e8d2789fc0 <col:16> col:16 implicit referenced ~ 'void () __generic noexcept' inline default trivial
  `-FunctionDecl 0x55e8d278a258 <line:2:1, line:11:1> line:2:15 foo 'void ()'
    |-CompoundStmt 0x55e8d27b9a98 <col:21, line:11:1>
    | `-CXXOperatorCallExpr 0x55e8d27b9a60 <line:4:3, col:12> 'int':'int'
    |   |-ImplicitCastExpr 0x55e8d27b9a08 <col:10, col:12> 'int (*)(int) const __generic' <FunctionToPointerDecay>
    |   | `-DeclRefExpr 0x55e8d27b9990 <col:10, col:12> 'int (int) const __generic' lvalue CXXMethod 0x55e8d27b9750 'operator()' 'int (int) const __generic'
    |   |-ImplicitCastExpr 0x55e8d27b9a48 <col:3> 'const __generic (lambda at test.cl:1:16)' lvalue <AddressSpaceConversion>
    |   | `-DeclRefExpr 0x55e8d278a338 <col:3> '__global (lambda at test.cl:1:16)':'__global (lambda at test.cl:1:16)' lvalue Var 0x55e8d2789500 'glambda' '__global (lambda at test.cl:1:16)':'__global (lambda at test.cl:1:16)'
    |   `-IntegerLiteral 0x55e8d278a358 <col:11> 'int' 1
    `-OpenCLKernelAttr 0x55e8d278a2f8 <line:2:1>

It seems we don't actually have temporaries.


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D69938/new/

https://reviews.llvm.org/D69938