Page MenuHomePhabricator

cchen (Chi Chun Chen)
User

Projects

User does not belong to any projects.

User Details

User Since
Aug 21 2019, 11:47 AM (45 w, 3 d)

Recent Activity

Thu, Jul 2

cchen updated the diff for D82800: [OPENMP50] extend array section for stride (Parsing/Sema/AST).

Fix based on feedback

Thu, Jul 2, 2:36 PM · Restricted Project
cchen updated the diff for D82800: [OPENMP50] extend array section for stride (Parsing/Sema/AST).

Rebase

Thu, Jul 2, 12:57 PM · Restricted Project

Wed, Jul 1

cchen added reviewers for D82245: [libomptarget] Add support for target update non-contiguous: hfinkel, Meinersbur, kkwli0, Hahnfeld.
Wed, Jul 1, 2:36 PM · Restricted Project, Restricted Project

Tue, Jun 30

cchen added a comment to D79972: [OpenMP5.0] map item can be non-contiguous for target update.

@ABataev , I'm considering emitting an extra dimension for a non-contiguous descriptor to support stride in this patch (stride = 1 in array section is just a special case for computing stride, however, the formula computing stride do not change). Do you think I should do it in this patch?

Tue, Jun 30, 3:46 PM · Restricted Project
cchen added inline comments to D79972: [OpenMP5.0] map item can be non-contiguous for target update.
Tue, Jun 30, 10:18 AM · Restricted Project

Mon, Jun 29

cchen updated the diff for D82800: [OPENMP50] extend array section for stride (Parsing/Sema/AST).
  1. Remove a commit that is not on master branch
  2. Add a rule in Sema to avoid non-contiguous array section for map clause
  3. Add tests
Mon, Jun 29, 4:06 PM · Restricted Project
cchen added inline comments to D82800: [OPENMP50] extend array section for stride (Parsing/Sema/AST).
Mon, Jun 29, 1:36 PM · Restricted Project
cchen updated the diff for D82800: [OPENMP50] extend array section for stride (Parsing/Sema/AST).

Revert unnecessary changes made by bot

Mon, Jun 29, 1:01 PM · Restricted Project
cchen created D82800: [OPENMP50] extend array section for stride (Parsing/Sema/AST).
Mon, Jun 29, 12:28 PM · Restricted Project
cchen added a comment to D82245: [libomptarget] Add support for target update non-contiguous.

ping

Mon, Jun 29, 12:28 PM · Restricted Project, Restricted Project

Fri, Jun 26

cchen updated the diff for D79972: [OpenMP5.0] map item can be non-contiguous for target update.

Rebase and resolve conflictions

Fri, Jun 26, 10:55 AM · Restricted Project
cchen updated the diff for D79972: [OpenMP5.0] map item can be non-contiguous for target update.

Fix based on feedback

Fri, Jun 26, 8:44 AM · Restricted Project

Thu, Jun 25

cchen updated the diff for D79972: [OpenMP5.0] map item can be non-contiguous for target update.

Pass Info directly

Thu, Jun 25, 3:18 PM · Restricted Project
cchen updated the diff for D79972: [OpenMP5.0] map item can be non-contiguous for target update.

Fix based on feedback

Thu, Jun 25, 2:44 PM · Restricted Project
cchen added inline comments to D79972: [OpenMP5.0] map item can be non-contiguous for target update.
Thu, Jun 25, 1:39 PM · Restricted Project
cchen added inline comments to D79972: [OpenMP5.0] map item can be non-contiguous for target update.
Thu, Jun 25, 1:39 PM · Restricted Project
cchen updated the diff for D82245: [libomptarget] Add support for target update non-contiguous.

Fix non_contig type

Thu, Jun 25, 1:39 PM · Restricted Project, Restricted Project
cchen updated the diff for D79972: [OpenMP5.0] map item can be non-contiguous for target update.

Fix coding style

Thu, Jun 25, 1:05 PM · Restricted Project
cchen added a comment to D79972: [OpenMP5.0] map item can be non-contiguous for target update.

ping

Thu, Jun 25, 9:40 AM · Restricted Project

Wed, Jun 24

cchen updated the diff for D82245: [libomptarget] Add support for target update non-contiguous.

Update test so that it don't depends on compiler generated code

Wed, Jun 24, 2:40 PM · Restricted Project, Restricted Project
cchen updated the diff for D82245: [libomptarget] Add support for target update non-contiguous.

Revert llvm pre-merge lint for existing code

Wed, Jun 24, 1:34 PM · Restricted Project, Restricted Project
cchen updated the diff for D82245: [libomptarget] Add support for target update non-contiguous.

Update test

Wed, Jun 24, 1:01 PM · Restricted Project, Restricted Project

Mon, Jun 22

cchen added a comment to D79972: [OpenMP5.0] map item can be non-contiguous for target update.

How do you plan to support
#pragma omp target update to (arr[1:2][1:2][0:2], x, b[1:5][0:2])
Are you going to split this into 3 updates since your are using the arg fields.

Mon, Jun 22, 4:10 PM · Restricted Project
cchen updated the diff for D79972: [OpenMP5.0] map item can be non-contiguous for target update.

Updated test for clarification

Mon, Jun 22, 4:07 PM · Restricted Project
cchen added a comment to D79972: [OpenMP5.0] map item can be non-contiguous for target update.

How do you plan to support
#pragma omp target update to (arr[1:2][1:2][0:2], x, b[1:5][0:2])
Are you going to split this into 3 updates since your are using the arg fields.

Mon, Jun 22, 3:35 PM · Restricted Project
cchen added inline comments to D82245: [libomptarget] Add support for target update non-contiguous.
Mon, Jun 22, 1:26 PM · Restricted Project, Restricted Project
cchen updated the diff for D82245: [libomptarget] Add support for target update non-contiguous.

Add test

Mon, Jun 22, 1:26 PM · Restricted Project, Restricted Project

Fri, Jun 19

cchen created D82245: [libomptarget] Add support for target update non-contiguous.
Fri, Jun 19, 3:14 PM · Restricted Project, Restricted Project

Tue, Jun 16

cchen updated the diff for D79972: [OpenMP5.0] map item can be non-contiguous for target update.

Fix based on feedback

Tue, Jun 16, 3:23 PM · Restricted Project
cchen updated the diff for D79972: [OpenMP5.0] map item can be non-contiguous for target update.

Fix Int64Ty issue (The bitNum of APInt I used before is 32)

Tue, Jun 16, 1:44 PM · Restricted Project
cchen added inline comments to D79972: [OpenMP5.0] map item can be non-contiguous for target update.
Tue, Jun 16, 9:54 AM · Restricted Project

Mon, Jun 15

cchen added inline comments to D79972: [OpenMP5.0] map item can be non-contiguous for target update.
Mon, Jun 15, 3:29 PM · Restricted Project
cchen updated the diff for D79972: [OpenMP5.0] map item can be non-contiguous for target update.

Resolve issues

Mon, Jun 15, 3:29 PM · Restricted Project

Fri, Jun 12

cchen updated the diff for D79972: [OpenMP5.0] map item can be non-contiguous for target update.

Fix based on feedback

Fri, Jun 12, 3:58 PM · Restricted Project
cchen added inline comments to D79972: [OpenMP5.0] map item can be non-contiguous for target update.
Fri, Jun 12, 3:58 PM · Restricted Project

Wed, Jun 10

cchen updated the diff for D79972: [OpenMP5.0] map item can be non-contiguous for target update.

Fix based on feedback

Wed, Jun 10, 1:55 PM · Restricted Project

Tue, Jun 9

cchen added a comment to D79972: [OpenMP5.0] map item can be non-contiguous for target update.

Do you have a test for mapping of something like arr[0][:n], where the base is an array subscript and the remaining part is an array section?

Tue, Jun 9, 2:21 PM · Restricted Project
cchen updated the diff for D79972: [OpenMP5.0] map item can be non-contiguous for target update.

Fix based on feedback

Tue, Jun 9, 2:21 PM · Restricted Project

Mon, Jun 8

cchen added a comment to D79972: [OpenMP5.0] map item can be non-contiguous for target update.

ping

Mon, Jun 8, 10:37 PM · Restricted Project

Jun 2 2020

cchen added a comment to D79972: [OpenMP5.0] map item can be non-contiguous for target update.

Still: Did you think about implementing it in the compiler instead of the runtime?

I'm not sure I understand your question, which part of code are you asking?
The main work compiler needs to do is to send the {offset, count, stride} struct to runtime.

I mean did you think about calling __tgt_target_data_update function in a loop in the compiler-generated code instead of putting it into the runtime?

Oh, I would prefer to call tgt_target_data_update once in the compiler and I'm also doing it now.

I was not quite correct. What I mean, is to generate the array with the array section as VLA in the compiler, and fill it in the loop generated by the compiler for non-contiguous sections but not in the runtime?
Say, we have the code:

int arr[3][3]
...
 #pragma omp update to(arr[1:2][1:2]

In this case, we're going to transfer the next elements:

000
0xx
0xx

In the compiler-generated code we emit something like this:

void *bptr[<n>];
void *ptr[<n>];
int64 sizes[<n>];
int64 maptypes[<n>];
for (int i = 0; i < <n>; ++i) {
  bptr[i] = &arr[1+i][1];
  ptr[i] = &arr[1+i][1];
  sizes[i] = ...;'
  maptypes[i] = ...;
}
call void @__tgt_target_data_update(i64 -1, i32 <n>, bptr, ptr, sizes, maptypes);

With this solution, you won't need to modify the runtime and add a new mapping flag.

Jun 2 2020, 4:29 PM · Restricted Project
cchen updated the diff for D79972: [OpenMP5.0] map item can be non-contiguous for target update.

Fix based on feedback

Jun 2 2020, 3:56 PM · Restricted Project
cchen abandoned D80590: [WIP][OPENMP] Fix assertion error for using alignas with OpenMP directive.
Jun 2 2020, 2:16 PM · Restricted Project
cchen added a comment to D79972: [OpenMP5.0] map item can be non-contiguous for target update.

Still: Did you think about implementing it in the compiler instead of the runtime?

I'm not sure I understand your question, which part of code are you asking?
The main work compiler needs to do is to send the {offset, count, stride} struct to runtime.

I mean did you think about calling __tgt_target_data_update function in a loop in the compiler-generated code instead of putting it into the runtime?

Jun 2 2020, 12:37 PM · Restricted Project
cchen added a comment to D79972: [OpenMP5.0] map item can be non-contiguous for target update.

Still: Did you think about implementing it in the compiler instead of the runtime?

Jun 2 2020, 12:05 PM · Restricted Project

Jun 1 2020

cchen updated the diff for D79972: [OpenMP5.0] map item can be non-contiguous for target update.
  • Use PointerIntPair to pass non-contiguous information in AST
  • Error out in Sema if we don't have enough size information for cases involving pointers
  • Allows *arr[N][M] since we don't need size information for the last dimension
  • Add more test cases
Jun 1 2020, 2:05 PM · Restricted Project

May 28 2020

cchen added inline comments to D79972: [OpenMP5.0] map item can be non-contiguous for target update.
May 28 2020, 9:54 AM · Restricted Project
cchen updated the diff for D79972: [OpenMP5.0] map item can be non-contiguous for target update.

Fix based on feedback

May 28 2020, 9:50 AM · Restricted Project

May 27 2020

cchen added inline comments to D79972: [OpenMP5.0] map item can be non-contiguous for target update.
May 27 2020, 3:50 PM · Restricted Project
cchen added a comment to D79972: [OpenMP5.0] map item can be non-contiguous for target update.

Is my guess correct that for OpenMP >= 50 for target update directive we always emit possibly non-continuous runtime calls?

My intent is to emit possibly non-contiguous runtime calls only if the analysis in Sema set the IsNonContiguous flag to true.

But this analysis only checks for the directive and the version,nothing else.

May 27 2020, 3:49 PM · Restricted Project
cchen updated the summary of D79972: [OpenMP5.0] map item can be non-contiguous for target update.
May 27 2020, 3:16 PM · Restricted Project
cchen added a comment to D79972: [OpenMP5.0] map item can be non-contiguous for target update.

Is my guess correct that for OpenMP >= 50 for target update directive we always emit possibly non-continuous runtime calls?

May 27 2020, 3:16 PM · Restricted Project
cchen updated the summary of D79972: [OpenMP5.0] map item can be non-contiguous for target update.
May 27 2020, 2:46 PM · Restricted Project
cchen added inline comments to D79972: [OpenMP5.0] map item can be non-contiguous for target update.
May 27 2020, 2:43 PM · Restricted Project
cchen added a comment to D79972: [OpenMP5.0] map item can be non-contiguous for target update.

How are you going to pass this non-contiguous data in the runtime? Are you going to map it in a loop or convert this non-contiguous data into the contiguous and map it as a contiguous chunk of data? Your presentation provides interface only interface changes but has nothing about implementation in the runtime.

May 27 2020, 2:10 PM · Restricted Project

May 26 2020

cchen updated the diff for D80590: [WIP][OPENMP] Fix assertion error for using alignas with OpenMP directive.

Add test

May 26 2020, 4:23 PM · Restricted Project
cchen added a comment to D80590: [WIP][OPENMP] Fix assertion error for using alignas with OpenMP directive.

Haven't added test yet since I'm not sure in which file should I add the test.

Where is the directive?

May 26 2020, 4:23 PM · Restricted Project
cchen updated the summary of D80590: [WIP][OPENMP] Fix assertion error for using alignas with OpenMP directive.
May 26 2020, 4:23 PM · Restricted Project
cchen added a reviewer for D80590: [WIP][OPENMP] Fix assertion error for using alignas with OpenMP directive: ABataev.
May 26 2020, 3:17 PM · Restricted Project
cchen created D80590: [WIP][OPENMP] Fix assertion error for using alignas with OpenMP directive.
May 26 2020, 3:17 PM · Restricted Project
cchen added a comment to D80590: [WIP][OPENMP] Fix assertion error for using alignas with OpenMP directive.

Haven't added test yet since I'm not sure in which file should I add the test.

May 26 2020, 3:17 PM · Restricted Project

May 21 2020

cchen updated the diff for D79972: [OpenMP5.0] map item can be non-contiguous for target update.

Remove redundant code

May 21 2020, 11:53 AM · Restricted Project
cchen added a comment to D79972: [OpenMP5.0] map item can be non-contiguous for target update.

ping

May 21 2020, 11:53 AM · Restricted Project

May 15 2020

cchen updated the diff for D79972: [OpenMP5.0] map item can be non-contiguous for target update.

Rebase

May 15 2020, 9:45 AM · Restricted Project

May 14 2020

cchen created D79972: [OpenMP5.0] map item can be non-contiguous for target update.
May 14 2020, 3:47 PM · Restricted Project

Apr 6 2020

cchen added a comment to D77581: Add map-type check for target and target data directive.

@ABataev , can you land it for me when you have time? thanks

Apr 6 2020, 3:16 PM · Restricted Project
cchen updated the diff for D77581: Add map-type check for target and target data directive.

Reuse existing function for target/target data check

Apr 6 2020, 2:43 PM · Restricted Project
cchen added inline comments to D77581: Add map-type check for target and target data directive.
Apr 6 2020, 2:43 PM · Restricted Project
cchen updated the diff for D77581: Add map-type check for target and target data directive.

Do not use lint on test

Apr 6 2020, 2:10 PM · Restricted Project
cchen created D77581: Add map-type check for target and target data directive.
Apr 6 2020, 1:04 PM · Restricted Project

Feb 28 2020

cchen added a comment to D75077: [OpenMP5.0] Allow pointer arithmetic in motion/map clause.

@ABataev, can you land it for me when you have time?

Feb 28 2020, 11:30 AM · Restricted Project
cchen updated the diff for D75077: [OpenMP5.0] Allow pointer arithmetic in motion/map clause.

return RelevantExpr || Visit(LE);

Feb 28 2020, 10:23 AM · Restricted Project
cchen updated the diff for D75077: [OpenMP5.0] Allow pointer arithmetic in motion/map clause.

Add assert to make sure that we never overwrite RelevantExpr

Feb 28 2020, 10:05 AM · Restricted Project
cchen added inline comments to D75077: [OpenMP5.0] Allow pointer arithmetic in motion/map clause.
Feb 28 2020, 9:37 AM · Restricted Project

Feb 27 2020

cchen updated the diff for D75077: [OpenMP5.0] Allow pointer arithmetic in motion/map clause.

Add test

Feb 27 2020, 3:24 PM · Restricted Project
cchen added a comment to D75077: [OpenMP5.0] Allow pointer arithmetic in motion/map clause.

Add back the check of UO_Deref

Feb 27 2020, 1:46 PM · Restricted Project
cchen updated the diff for D75077: [OpenMP5.0] Allow pointer arithmetic in motion/map clause.

Add back the check of UO_Deref

Feb 27 2020, 1:46 PM · Restricted Project
cchen updated the diff for D75077: [OpenMP5.0] Allow pointer arithmetic in motion/map clause.

Fix based on feedback

Feb 27 2020, 1:32 PM · Restricted Project
cchen updated the diff for D75077: [OpenMP5.0] Allow pointer arithmetic in motion/map clause.

Remove redundant else

Feb 27 2020, 1:04 PM · Restricted Project
cchen updated the diff for D75077: [OpenMP5.0] Allow pointer arithmetic in motion/map clause.

Fix test and add check for unary operator.

Feb 27 2020, 1:04 PM · Restricted Project
cchen added inline comments to D75077: [OpenMP5.0] Allow pointer arithmetic in motion/map clause.
Feb 27 2020, 12:23 PM · Restricted Project
cchen added inline comments to D75077: [OpenMP5.0] Allow pointer arithmetic in motion/map clause.
Feb 27 2020, 12:23 PM · Restricted Project
cchen updated the diff for D75077: [OpenMP5.0] Allow pointer arithmetic in motion/map clause.

Fix based on feedback

Feb 27 2020, 10:15 AM · Restricted Project

Feb 26 2020

cchen added inline comments to D75077: [OpenMP5.0] Allow pointer arithmetic in motion/map clause.
Feb 26 2020, 11:59 AM · Restricted Project
cchen added inline comments to D75077: [OpenMP5.0] Allow pointer arithmetic in motion/map clause.
Feb 26 2020, 8:41 AM · Restricted Project

Feb 25 2020

cchen added inline comments to D75077: [OpenMP5.0] Allow pointer arithmetic in motion/map clause.
Feb 25 2020, 10:18 AM · Restricted Project
cchen added inline comments to D75077: [OpenMP5.0] Allow pointer arithmetic in motion/map clause.
Feb 25 2020, 10:18 AM · Restricted Project
cchen updated the diff for D75077: [OpenMP5.0] Allow pointer arithmetic in motion/map clause.

Add BinaryOperator in components and fix logic in MapBaseChecker.

Feb 25 2020, 10:09 AM · Restricted Project

Feb 24 2020

cchen added inline comments to D75077: [OpenMP5.0] Allow pointer arithmetic in motion/map clause.
Feb 24 2020, 2:20 PM · Restricted Project
cchen updated the diff for D75077: [OpenMP5.0] Allow pointer arithmetic in motion/map clause.

Remove the counting in MapBaseChecker, instead, comparing the type of root with
the type of LHS and RHS, and only visit the subtree having the same type as root.
This scheme pass all the test I have. I'll add BinOp into components in the next
revision. Thanks

Feb 24 2020, 2:20 PM · Restricted Project
cchen added inline comments to D75077: [OpenMP5.0] Allow pointer arithmetic in motion/map clause.
Feb 24 2020, 1:52 PM · Restricted Project
cchen updated the diff for D75077: [OpenMP5.0] Allow pointer arithmetic in motion/map clause.

Remove redundant code in test.

Feb 24 2020, 1:43 PM · Restricted Project
cchen created D75077: [OpenMP5.0] Allow pointer arithmetic in motion/map clause.
Feb 24 2020, 1:34 PM · Restricted Project

Feb 21 2020

cchen added a comment to D74970: [OpenMP] Refactor the analysis in checkMapClauseBaseExpression using StmtVisitor class..

@ABataev, can you land it for me when you have time? Thanks

Feb 21 2020, 2:39 PM · Restricted Project, Restricted Project
cchen updated the diff for D74970: [OpenMP] Refactor the analysis in checkMapClauseBaseExpression using StmtVisitor class..

Update test

Feb 21 2020, 2:20 PM · Restricted Project, Restricted Project
cchen updated the diff for D74970: [OpenMP] Refactor the analysis in checkMapClauseBaseExpression using StmtVisitor class..

Fix test and sorry for my carelessness

Feb 21 2020, 1:53 PM · Restricted Project, Restricted Project
cchen updated the diff for D74970: [OpenMP] Refactor the analysis in checkMapClauseBaseExpression using StmtVisitor class..

Add test for mapping function

Feb 21 2020, 1:25 PM · Restricted Project, Restricted Project
cchen added a comment to D74970: [OpenMP] Refactor the analysis in checkMapClauseBaseExpression using StmtVisitor class..

Did you try to run the tests? I would suggest adding a test, at least, where a function is mapped. We should see the error message here. Seems to me, we don't have this one.

We already have test for err_omp_invalid_map_this_expr, note_omp_invalid_subscript_on_this_ptr_map, etc.. in target_message.cpp line 44 without checking for value dependence. Do you mean that I should add a test for the check of value dependence? In that case, we don't print any messages.

No. But previously, if we tried to map a function, not a variable, it would be skipped silently. With this patch, this behavior will change. I suggest adding a test with mapping a function to check that error message is emitted. Something like:

int foo();
#pragma omp target map(foo) // <- must be an error here

And I just asked if you tried to run the tests with this patch. Did the result change or it is the same?

Feb 21 2020, 1:16 PM · Restricted Project, Restricted Project
cchen added a comment to D74970: [OpenMP] Refactor the analysis in checkMapClauseBaseExpression using StmtVisitor class..

Did you try to run the tests? I would suggest adding a test, at least, where a function is mapped. We should see the error message here. Seems to me, we don't have this one.

Feb 21 2020, 12:55 PM · Restricted Project, Restricted Project
cchen added inline comments to D74970: [OpenMP] Refactor the analysis in checkMapClauseBaseExpression using StmtVisitor class..
Feb 21 2020, 12:13 PM · Restricted Project, Restricted Project
cchen updated the diff for D74970: [OpenMP] Refactor the analysis in checkMapClauseBaseExpression using StmtVisitor class..

Fix by feedback

Feb 21 2020, 12:13 PM · Restricted Project, Restricted Project