diff mbox

[gomp4] force vector size to 32

Message ID 55D7907B.8040002@acm.org
State New
Headers show

Commit Message

Nathan Sidwell Aug. 21, 2015, 8:56 p.m. UTC 
I've committed this patch to force the vector_length to constant 32.  This helps 
in 2  places

1) Cesar's reduction patch set made this assumption, as we can't emit a loop for 
vector reductions, as he explained.

2) Routines can now optimize on this axis knowing it is 32.

The only case lost is where one selects a single element from each gang, which 
is a bizarre way of operating the GPGPU, so no real loss.

Also fixed a testcase that had got a duplicated body.

nathan
diff mbox

Patch

2015-08-21  Nathan Sidwell  <nathan@codesourcery.com>

	* config/nvptx/nvptx.c (nvptx_validate_dims): Set vector length to
	32 always.

	testsuite/
	* c-c++-common/goacc/routine-5.c: Remove duplication.

Index: gcc/config/nvptx/nvptx.c
===================================================================
--- gcc/config/nvptx/nvptx.c	(revision 227082)
+++ gcc/config/nvptx/nvptx.c	(working copy)
@@ -3570,29 +3570,18 @@  nvptx_validate_dims (tree decl, int dims
 {
   bool changed = false;
 
-  if (fn_level >= 0)
-    /* This is a routine.  All dimensions are dynamic and controlled
-       by the  calling function.  Because we permit a 1vx1wxNg
-       geometry, we can't take the opportunity to fix the vector
-       dimension inside a routine.  Perhaps we should?  */
-    return false;
-  
-  /* If the worker size is not 1, the vector size must be 32.  If
-     the vector size is not 1, it must be 32.  */
-  if ((dims[GOMP_DIM_WORKER] > 1 || dims[GOMP_DIM_WORKER] == 0)
-      || (dims[GOMP_DIM_VECTOR] > 1 || dims[GOMP_DIM_VECTOR] == 0))
+  /* The vector size must be 32, unless this is a SEQ routine.  */
+  if (fn_level <= GOMP_DIM_VECTOR
+      && dims[GOMP_DIM_VECTOR] != PTX_VECTOR_LENGTH)
     {
-      if (dims[GOMP_DIM_VECTOR] != PTX_VECTOR_LENGTH)
-	{
-	  if (dims[GOMP_DIM_VECTOR] >= 0)
-	    warning_at (DECL_SOURCE_LOCATION (decl), 0,
-			dims[GOMP_DIM_VECTOR]
-			? "using vector_length (%d), ignoring %d"
-			: "using vector_length (%d), ignoring runtime setting",
-			PTX_VECTOR_LENGTH, dims[GOMP_DIM_VECTOR]);
-	  dims[GOMP_DIM_VECTOR] = PTX_VECTOR_LENGTH;
-	  changed = true;
-	}
+      if (dims[GOMP_DIM_VECTOR] >= 0 && fn_level < 0)
+	warning_at (DECL_SOURCE_LOCATION (decl), 0,
+		    dims[GOMP_DIM_VECTOR]
+		    ? "using vector_length (%d), ignoring %d"
+		    : "using vector_length (%d), ignoring runtime setting",
+		    PTX_VECTOR_LENGTH, dims[GOMP_DIM_VECTOR]);
+      dims[GOMP_DIM_VECTOR] = PTX_VECTOR_LENGTH;
+      changed = true;
     }
 
   /* Check the num workers is not too large.  */
Index: gcc/testsuite/c-c++-common/goacc/routine-5.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/routine-5.c	(revision 227082)
+++ gcc/testsuite/c-c++-common/goacc/routine-5.c	(working copy)
@@ -17,24 +17,3 @@  void Foo ()
 
 #pragma acc routine  // { dg-error "not followed by function" }
 int i;
-
-/* Test invalid use of clauses with routine.  */
-/* { dg-do compile } */
-/* { dg-additional-options "-W */
-
-void Bar ();
-
-void Foo ()
-{
-  Bar ();
-}
-
-#pragma acc routine (Bar) // { dg-error "must be applied before use" }
-
-#pragma acc routine (Foo) gang // { dg-error "must be applied before definition" }
-
-#pragma acc routine (Baz) // { dg-error "not been declared" }
-
-#pragma acc routine  // { dg-error "not followed by function" }
-int i;
-