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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
