[PATCH] D28595: [RFC] Introduce directive.scope.entry and directive.scope.exit for IR-level Region Annotation
Hongbin Zheng via Phabricator via llvm-commits
llvm-commits at lists.llvm.org
Wed Jan 11 22:01:26 PST 2017
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>
More information about the llvm-commits
mailing list