On 07/29/2014 02:07 AM, Thomas Schwinge wrote:
> On Thu, 20 Mar 2014 15:42:48 +0100, I wrote:
>> On Tue, 18 Mar 2014 14:50:44 +0100, I wrote:
>>> On Tue, 18 Mar 2014 16:37:24 +0400, Ilmir Usmanov <i.usma...@samsung.com>
>>> wrote:
>>>> This patch introduces support of OpenACC loop directive (and combined
>>>> directives) in C front-end up to GENERIC. Currently no clause is allowed.
>>>
>>> Thanks! I had worked on a simpler patch, not yet dealing with combined
>>> clauses. Also, I have some work for the GIMPLE level, namely building on
>>> GIMPLE_OMP_FOR, adding a new GF_OMP_FOR_KIND_OACC_LOOP. I'll post this
>>> soon.
>>
>> Here are the patches, committed in r208702..4 to gomp-4_0-branch.
>
> Cesar, I hope I'm not confusing things here, but I remember that you once
> pointed out that in Fortran OpenACC, we have to explicitly specify the
> loop iteration variable in a private clause, whereas the OpenACC
> specification says this needs to happen automatically (predetermined data
> attribute). I see gcc/fortran/openmp.c:gfc_resolve_do_iterator add this
> private clause, which I assume we're using also for OpenACC loops.
> Looking at -fdump-tree-all output for the following two test cases
> complied in OpenMP and OpenACC mode, I see that indeed an explicit
> private clause is added for Fortran code, but not for C. Why is that
> required? (I have not yet spent any time on figuring this out myself.)
>
> int
> main(void)
> {
> int i;
>
> #pragma acc parallel
> #pragma acc loop
> #pragma omp parallel for
> for (i = 0; i < 10; ++i)
> ;
>
> return 0;
> }
>
> program test
> implicit none
> integer :: i
>
> !$acc parallel
> !$acc loop
> !$omp parallel do
> DO i = 1, 10
> ENDDO
> !$omp end parallel do
> !$acc end parallel
> end
>
> In light of this, please also review whether the following
> gimplify_omp_for changes of mine (when I originally added the OACC_LOOP
> support) are correct. Evidently, this code still has TODO markers in it;
> this was well before you added preliminary support for the private
> clause. That is, should we use GOVD_PRIVATE also for OACC_LOOP?
OMP has both a private and a shared clause. And those neither of those
clauses can share variables, i.e. an index variable cannot be both
private and shared. By default, index variables are private in OMP, so
the gimplification pass makes those variables explicitly private so that
later passes can check for errors.
Because the fortran frontend needs to do a little more error handling
early on, it makes the induction variables private.
> Why does this nevertheless currently work for C for loops without any
> private clause for the loop iteration variable? Maybe because in C, the
> loop iteration variable ends up in a register and that is not shared
> between threads? I do see a private clause being added during
> gimplification for the OpenMP C test case, but not for OpenACC -- which I
> assume is precisely due to the code I'm quoting below.
OMP lowering pass makes the index variables local to the parallel
function/kernel. So each thread effectively has its private copy of the
index variable.
> Once this has been resolved, please also remove the explicit private
> clauses from the existing test cases (Fortran and also C, as applicable)
> to make sure that the predermined data attributes are tested there.
Considering that ACC lacks a shared clause, it should be harmless to
revert that patch below. Do you want to revert it with your
private/firstprivate patch?
Cesar
>> --- gcc/gimplify.c
>> +++ gcc/gimplify.c
>> @@ -6683,14 +6683,36 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>> gimple_seq for_body, for_pre_body;
>> int i;
>> bool simd;
>> + enum gimplify_omp_var_data govd_private;
>> + enum omp_region_type ort;
>> bitmap has_decl_expr = NULL;
>>
>> orig_for_stmt = for_stmt = *expr_p;
>>
>> - simd = TREE_CODE (for_stmt) == OMP_SIMD
>> - || TREE_CODE (for_stmt) == CILK_SIMD;
>> - gimplify_scan_omp_clauses (&OMP_FOR_CLAUSES (for_stmt), pre_p,
>> - simd ? ORT_SIMD : ORT_WORKSHARE);
>> + switch (TREE_CODE (for_stmt))
>> + {
>> + case OMP_FOR:
>> + case OMP_DISTRIBUTE:
>> + simd = false;
>> + govd_private = GOVD_PRIVATE;
>> + ort = ORT_WORKSHARE;
>> + break;
>> + case OACC_LOOP:
>> + simd = false;
>> + govd_private = /* TODO */ GOVD_LOCAL;
>> + ort = /* TODO */ ORT_WORKSHARE;
>> + break;
>> + case OMP_SIMD:
>> + case CILK_SIMD:
>> + simd = true;
>> + govd_private = GOVD_PRIVATE;
>> + ort = ORT_SIMD;
>> + break;
>> + default:
>> + gcc_unreachable ();
>> + }
>> +
>> + gimplify_scan_omp_clauses (&OMP_FOR_CLAUSES (for_stmt), pre_p, ort);
>>
>> /* Handle OMP_FOR_INIT. */
>> for_pre_body = NULL;
>> @@ -6722,6 +6744,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>>
>> if (OMP_FOR_INIT (for_stmt) == NULL_TREE)
>> {
>> + gcc_assert (TREE_CODE (for_stmt) != OACC_LOOP);
>> for_stmt = walk_tree (&OMP_FOR_BODY (for_stmt), find_combined_omp_for,
>> NULL, NULL);
>> gcc_assert (for_stmt != NULL_TREE);
>> @@ -6742,7 +6765,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>> gcc_assert (INTEGRAL_TYPE_P (TREE_TYPE (decl))
>> || POINTER_TYPE_P (TREE_TYPE (decl)));
>>
>> - /* Make sure the iteration variable is private. */
>> + /* Make sure the iteration variable is some kind of private. */
>> tree c = NULL_TREE;
>> if (orig_for_stmt != for_stmt)
>> /* Do this only on innermost construct for combined ones. */;
>> @@ -6768,6 +6791,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>> }
>> else
>> {
>> + gcc_assert (govd_private == GOVD_PRIVATE);
>> bool lastprivate
>> = (!has_decl_expr
>> || !bitmap_bit_p (has_decl_expr, DECL_UID (decl)));
>> @@ -6785,7 +6809,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>> else if (omp_is_private (gimplify_omp_ctxp, decl, simd))
>> omp_notice_variable (gimplify_omp_ctxp, decl, true);
>> else
>> - omp_add_variable (gimplify_omp_ctxp, decl, GOVD_PRIVATE | GOVD_SEEN);
>> + omp_add_variable (gimplify_omp_ctxp, decl, govd_private | GOVD_SEEN);
>>
>> /* If DECL is not a gimple register, create a temporary variable to
>> act
>> as an iteration counter. This is valid, since DECL cannot be
>> @@ -6799,7 +6823,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>>
>> gimplify_seq_add_stmt (&for_body, gimple_build_assign (decl, var));
>>
>> - omp_add_variable (gimplify_omp_ctxp, var, GOVD_PRIVATE | GOVD_SEEN);
>> + omp_add_variable (gimplify_omp_ctxp, var, govd_private | GOVD_SEEN);
>> }
>> else
>> var = decl;
>> @@ -6936,7 +6960,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>> t = TREE_VEC_ELT (OMP_FOR_INIT (for_stmt), i);
>> decl = TREE_OPERAND (t, 0);
>> var = create_tmp_var (TREE_TYPE (decl), get_name (decl));
>> - omp_add_variable (gimplify_omp_ctxp, var, GOVD_PRIVATE | GOVD_SEEN);
>> + omp_add_variable (gimplify_omp_ctxp, var, govd_private | GOVD_SEEN);
>> TREE_OPERAND (t, 0) = var;
>> t = TREE_VEC_ELT (OMP_FOR_INCR (for_stmt), i);
>> TREE_OPERAND (t, 1) = copy_node (TREE_OPERAND (t, 1));
>> @@ -6952,6 +6976,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>> case OMP_SIMD: kind = GF_OMP_FOR_KIND_SIMD; break;
>> case CILK_SIMD: kind = GF_OMP_FOR_KIND_CILKSIMD; break;
>> case OMP_DISTRIBUTE: kind = GF_OMP_FOR_KIND_DISTRIBUTE; break;
>> + case OACC_LOOP: kind = GF_OMP_FOR_KIND_OACC_LOOP; break;
>> default:
>> gcc_unreachable ();
>> }
>
>
> Grüße,
> Thomas
>
