[PATCH] D28595: [RFC] Introduce directive.scope.entry and directive.scope.exit for IR-level Region Annotation

[Hongbin Zheng via Phabricator via llvm-commits]

etherzhhb created this revision.
etherzhhb added a reviewer: xtian.
etherzhhb added a subscriber: llvm-commits.
etherzhhb set the repository for this revision to rL LLVM.

This is a draft of the directive.scope.entry/directive.scope.exit intrinsics I mentioned in the "[llvm-dev] [RFC] IR-level Region Annotations" thread[1]. Currently this patch do not have a concrete use case, it just demonstrate the general idea:

1. directive.scope.entry defines a token and  directive.scope.exit consumes the token defined by directive.scope.entry.
2. directive.scope.entry mark the entry of the region and directive.scope.exit mark the exit of the token in the CFG
3. Extra information of the region is annotated as operand bundle of the directive.scope.entry intrinsic.

For example, from:

  #pragma prefetch x:1:20 y:0:10
  for (i=0; i<2*N; i++) { xp[i*m + j] = -1; yp[i*n +j] = -2; }

we can generate the region annotation as:

  %0 = tail call token @llvm.directive.scope.entry() [ "prefetch"(i32 *xp, i64 1, i64 20, i32 *yp, i64 0, i64 10) ]
  ... loop ...
  tail call void @llvm.directive.scope.exit(token %0)

Or

  #pragma omp parallel for
  for (i=0; i<2*N; i++) { xp[i] =0; }

we can generate:

  %0 = tail call token @llvm.directive.scope.entry() [ "omp.parallel.for"() ]
  ... loop ...
  tail call void @llvm.directive.scope.exit(token %0)

Based on these intrinsics, we can introduce wrapper classes to represent different regions:

  class OmpParallelForDirective: public IntrinsicInst {
  public:
    static inline bool classof(const IntrinsicInst *I) {
      return I->getIntrinsicID() == Intrinsic::directive_scope_entry && hasOperandBundle("omp.parallel.for"); // later we can replace this by an enum
    }
    ...
    
    Instruction *getRegionEntry() const;
    Instruction *getRegionExit() const;
  }
  
  class PrefetchDirective: public IntrinsicInst {
  public:
    static inline bool classof(const IntrinsicInst *I) {
      return I->getIntrinsicID() == Intrinsic::directive_scope_entry && hasOperandBundle("prefetch"); 
    }
  
    ArrayRef<Value*> getValues() const; // return { xp, yp } in the previous example
  
    Instruction *getRegionEntry() const;
    Instruction *getRegionExit() const;
  }

[1]https://groups.google.com/d/msg/llvm-dev/uSlv-hZJm_Y/KRSWE9pgEAAJ


Repository:
  rL LLVM

https://reviews.llvm.org/D28595

Files:
  include/llvm/IR/Intrinsics.td


Index: include/llvm/IR/Intrinsics.td
===================================================================
--- include/llvm/IR/Intrinsics.td
+++ include/llvm/IR/Intrinsics.td
@@ -596,6 +596,10 @@
 def int_experimental_gc_relocate : Intrinsic<[llvm_any_ty],
                                 [llvm_token_ty, llvm_i32_ty, llvm_i32_ty],
                                 [IntrReadMem]>;
+//===------------------------ Directive Intrinsics ------------------------===//
+//
+def int_directive_scope_entry : Intrinsic<[llvm_token_ty], [], []>;
+def int_directive_scope_exit  : Intrinsic<[], [llvm_token_ty], []>;
 
 //===-------------------------- Other Intrinsics --------------------------===//
 //


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D28595.84072.patch
Type: text/x-patch
Size: 692 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20170112/9293ad07/attachment.bin>