[PATCH] D17019: [OpenMP] Code generation for teams - kernel launching

[Alexey Bataev via cfe-commits]

ABataev added inline comments.

================
Comment at: lib/CodeGen/CGOpenMPRuntime.cpp:3799-3806
@@ -3780,4 +3798,10 @@
   OffloadEntriesInfoManager.registerTargetRegionEntryInfo(
       DeviceID, FileID, ParentName, Line, OutlinedFn, OutlinedFnID);
 }
 
+/// \brief Emit the num_teams clause of an enclosed teams directive at the
+/// target region scope. If there is no teams directive associated with the
+/// target directive, or if there is no num_teams clause associated with the
+/// enclosed teams directive, return nullptr.
+static llvm::Value *
+emitNumTeamsClauseForTargetDirective(CodeGenFunction &CGF,
----------------
sfantao wrote:
> ABataev wrote:
> > sfantao wrote:
> > > ABataev wrote:
> > > > I don't understand why global var is not captured in target region. If it is not implemented yet, it must be implemented. If it is not captured, it must be captured in Sema. We should not introduce some function/other objects to find a workaround for 'not implemented' features.
> > > Sorry, I was not clear in my comment. It is not that globals are not captured in target regions - they are, we already have Sema doing that. 
> > > 
> > > My point is that exactly because we capture globals in target regions the magic that `OMPCapturedExprDecl` introduces does not work for that specific case. So, if we have something like:
> > > 
> > > ```
> > > int Gbl;
> > > 
> > > foo() {
> > >   #pragma omp target
> > >   #pragma omp teams num_teams(Gbl)
> > >   {}
> > > }
> > > ```
> > > when the DeclRefExpr for Gbl used in num_teams is emitted in the scope that encloses '#pragma omp target', it will crash because Gbl is not a local and is marked as refer to enclosing capture. 
> > > 
> > > If I got it right, a solution based on `OMPCapturedExprDecl` basically makes local declarations whose initializers are the expression we are interested in. In the cases that  `OMPCapturedExprDecl` is currently employed we don't have globals being captured and that is why it works fine.
> > > 
> > > It is likely I am missing something here. Let me know if you need me to provide more details.
> > > 
> > > Thanks!
> > It should not crash, because if it is captured, we must use captured version of this variable, passed in arguments to outlined function
> I am afraid I may not be understanding what you want me to do. Going back to my example:
> 
> ```
> int Gbl;
> 
> foo() {
>   // a) I need to emit num_teams(Gbl) here. DeclRefExpr(Gbl) emission won't work because it is marked "refer to enclosing capture". 
>   #pragma omp target
>   // b) If I emit it here, that's fine because I already have the arguments of the outlined function, but that is not what I need.
>   #pragma omp teams num_teams(Gbl)
>   {}
> }
> ```
> Can you please elaborate on how `OMPCapturedExprDecl` would help me implement a). Sorry for the trouble.
> 
> Thanks!
Ok, why you don't want to emit it in b), but in a)?


http://reviews.llvm.org/D17019