diff mbox

GOMP_target: alignment (was: [gomp4] #pragma omp target* fixes)

Message ID 87y53jgu5m.fsf@schwinge.name
State New
Headers show

Commit Message

Thomas Schwinge Dec. 17, 2013, 7:21 p.m. UTC 
Hi!

On Mon, 16 Dec 2013 17:58:26 +0100, Jakub Jelinek <jakub@redhat.com> wrote:
> I'd indeed prefer if you just used one
> array, it can be say just uchar array of twice the width, with even chars
> for alignment and odd for kinds (or vice versa), compared to two arrays
> it is tiny bit cheaper at the caller side IMHO.

Like this, for gomp-4_0-branch?  Is hard-coding a shift by eight bits OK,
or am I to fiddle with CHAR_TYPE_SIZE, and the like?

While conceptually nicer, using some build_* machinery to actually build
a datatype and initializer for an array of a »struct { char kind; char
alignment; }« would be more difficult, for unclear benefit, so I didn't
look into that.

commit 46002ec0e69e2fbc1f14d2549a5cbb93849c1da1
Author: Thomas Schwinge <thomas@codesourcery.com>
Date:   Tue Dec 17 13:44:46 2013 +0100

    Prepare OpenACC memory mapping interface for additional mapping kinds.
    
    	gcc/
    	* omp-low.c (lower_oacc_parallel): Switch kinds array to unsigned
    	short, and shift alignment description to begin at bit 8.
    	libgomp/
    	* libgomp_g.h (GOACC_parallel): Switch kinds array to unsigned
    	short.
    	* oacc-parallel.c (GOACC_parallel): Likewise, and catch
    	unsupported kinds.



Grüße,
 Thomas

Comments

Jakub Jelinek Dec. 17, 2013, 7:35 p.m. UTC | #1 
On Tue, Dec 17, 2013 at 08:21:57PM +0100, Thomas Schwinge wrote:
> On Mon, 16 Dec 2013 17:58:26 +0100, Jakub Jelinek <jakub@redhat.com> wrote:
> > I'd indeed prefer if you just used one
> > array, it can be say just uchar array of twice the width, with even chars
> > for alignment and odd for kinds (or vice versa), compared to two arrays
> > it is tiny bit cheaper at the caller side IMHO.
> 
> Like this, for gomp-4_0-branch?  Is hard-coding a shift by eight bits OK,
> or am I to fiddle with CHAR_TYPE_SIZE, and the like?

I think shift by 8 bits is fine, I believe 8 bits is the minimum a char can
have and it is better if you have a constant you can use on both sides.

	Jakub
diff mbox

Patch

diff --git gcc/omp-low.c gcc/omp-low.c
index e0f7d1d..eb755c3 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -8775,7 +8775,7 @@  lower_oacc_parallel (gimple_stmt_iterator *gsi_p, omp_context *ctx)
       TREE_ADDRESSABLE (TREE_VEC_ELT (t, 1)) = 1;
       TREE_STATIC (TREE_VEC_ELT (t, 1)) = 1;
       TREE_VEC_ELT (t, 2)
-	= create_tmp_var (build_array_type_nelts (unsigned_char_type_node,
+	= create_tmp_var (build_array_type_nelts (short_unsigned_type_node,
 						  map_cnt),
 			  ".omp_data_kinds");
       DECL_NAMELESS (TREE_VEC_ELT (t, 2)) = 1;
@@ -8884,7 +8884,7 @@  lower_oacc_parallel (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	    if (TREE_CODE (s) != INTEGER_CST)
 	      TREE_STATIC (TREE_VEC_ELT (t, 1)) = 0;
 
-	    unsigned char tkind = 0;
+	    unsigned short tkind = 0;
 	    switch (OMP_CLAUSE_CODE (c))
 	      {
 	      case OMP_CLAUSE_MAP:
@@ -8903,9 +8903,9 @@  lower_oacc_parallel (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	    if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign)
 	      talign = DECL_ALIGN_UNIT (ovar);
 	    talign = ceil_log2 (talign);
-	    tkind |= talign << 3;
+	    tkind |= talign << 8;
 	    CONSTRUCTOR_APPEND_ELT (vkind, purpose,
-				    build_int_cst (unsigned_char_type_node,
+				    build_int_cst (short_unsigned_type_node,
 						   tkind));
 	    if (nc && nc != c)
 	      c = nc;
diff --git libgomp/libgomp_g.h libgomp/libgomp_g.h
index 394f3a8..34d26f6 100644
--- libgomp/libgomp_g.h
+++ libgomp/libgomp_g.h
@@ -217,6 +217,6 @@  extern void GOMP_teams (unsigned int, unsigned int);
 /* oacc-parallel.c */
 
 extern void GOACC_parallel (int, void (*) (void *), const void *,
-			    size_t, void **, size_t *, unsigned char *);
+			    size_t, void **, size_t *, unsigned short *);
 
 #endif /* LIBGOMP_G_H */
diff --git libgomp/oacc-parallel.c libgomp/oacc-parallel.c
index 730b83b..bf7b74c 100644
--- libgomp/oacc-parallel.c
+++ libgomp/oacc-parallel.c
@@ -25,12 +25,29 @@ 
 
 /* This file handles the OpenACC parallel construct.  */
 
+#include "libgomp.h"
 #include "libgomp_g.h"
 
 void
 GOACC_parallel (int device, void (*fn) (void *), const void *openmp_target,
 		size_t mapnum, void **hostaddrs, size_t *sizes,
-		unsigned char *kinds)
+		unsigned short *kinds)
 {
-  GOMP_target (device, fn, openmp_target, mapnum, hostaddrs, sizes, kinds);
+  unsigned char kinds_[mapnum];
+  size_t i;
+
+  /* TODO.  Eventually, we'll be interpreting all mapping kinds according to
+     the OpenACC semantics; for now we're re-using what is implemented for
+     OpenMP.  */
+  for (i = 0; i < mapnum; ++i)
+    {
+      unsigned char kind = kinds[i];
+      unsigned char align = kinds[i] >> 8;
+      if (kind > 4)
+	gomp_fatal ("memory mapping kind %x for %zd is not yet supported",
+		    kind, i);
+
+      kinds_[i] = kind | align << 3;
+    }
+  GOMP_target (device, fn, openmp_target, mapnum, hostaddrs, sizes, kinds_);
 }