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

[etherzhhb]

Summary

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

include/llvm/IR/Intrinsics.td

def int_experimental_gc_statepoint : Intrinsic<[llvm_token_ty],
                                 llvm_i32_ty, llvm_vararg_ty],
                                 [Throws]>;
 
 def int_experimental_gc_result   : Intrinsic<[llvm_any_ty], [llvm_token_ty],
                                              [IntrReadMem]>;
 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 --------------------------===//
 //
 def int_flt_rounds : Intrinsic<[llvm_i32_ty]>,
                      GCCBuiltin<"__builtin_flt_rounds">;
 def int_trap : Intrinsic<[], [], [IntrNoReturn]>,
                GCCBuiltin<"__builtin_trap">;
 def int_debugtrap : Intrinsic<[]>,