On Tue, 20 Oct 2015, Alexander Monakov wrote:
> On Tue, 20 Oct 2015, Bernd Schmidt wrote:
>
> > On 10/20/2015 08:34 PM, Alexander Monakov wrote:
> > > Due to special treatment of types, emitting variables of type _Bool in
> > > global scope is impossible: extern references are emitted with .u8, but
> > > definitions use .u64. This patch fixes the issue by treating boolean type
> > > as
> > > integer types.
> > >
> > > * config/nvptx/nvptx.c (init_output_initializer): Also accept
> > > BOOLEAN_TYPE.
> >
> > Interesting, what was the testcase? I didn't stumble over this one. In any
> > case, I think this patch is ok for trunk.
>
> libgomp has 'bool gomp_cancel_var' in global scope, and since it is not
> compiled with -ffunction-sections, GOMP_parallel pulls in
> GOMP_cancel (same TU, parallel.c), which references the variable. Anything
> with "#pragma omp parallel" would fail to link.
Hi Bernd,
There's another issue with handling of types in init_output_initializer.
Test libgomp.c++/examples-4/declare_target-2.C fails with "libgomp: Can't map
target variables (size mismatch)".
Variables varX and varY of struct/class type have size 4, but their definition
in PTX code has size 8 because init_output_initializer changes 'type' to
ptr_type_node and emits them in size-8 chunks. Could you please explain why
the chunking is in a pointer type rather than bytewise? It seems the current
scheme is problematic for all types with sizes non divisible by 8, except
those mapped to PTX types already.
With the trivial patch changing ptr_type_node to char_type_node (pasted below)
I get the desired behavior with no libgomp testsuite regressions.
While looking at relevant code I found what appears to be a bad assert in
nvptx.c:
gcc_assert (size = decl_chunk_size);
and also above that:
gcc_assert (vec_pos = XVECLEN (pat, 0));
Those should have been equality checks rather than assignments, correct?
Apart from those, I found another assignment-in-assertion in GCC in tree-eh.c:
gcc_assert (ri = (int)ri);
I'll send a separate mail about that shortly.
Thanks.
Alexander
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -1893,7 +1893,7 @@ init_output_initializer (FILE *file, const char *name,
const_tree type,
&& TREE_CODE (type) != REAL_TYPE)
|| sz < 0
|| sz > HOST_BITS_PER_WIDE_INT)
- type = ptr_type_node;
+ type = char_type_node;
decl_chunk_size = int_size_in_bytes (type);
decl_chunk_mode = int_mode_for_mode (TYPE_MODE (type));
decl_offset = 0;
