Add passing of kernel schedule parameters.

    
    * gcc/builtin-types.def: Change signature of builtin.
    * gcc/fortran/types.def: Likewise.
    * gcc/oacc-builtins.def: Likewise.
    * gcc/oacc-low.c (expand_oacc_kernels): Change builtin call generation.
    * liboacc/liboacc-internal.h: Change signature.
    * liboacc/liboacc.c: Likewise.
    * liboacc/liboacc.h: Likewise.
    * liboacc/runtime.c: Pass schedule parameters to OpenCL runtime.


git-svn-id: https://gcc.gnu.org/svn/gcc/branches/openacc-1_0-branch@204214 138bc75d-0d04-0410-961f-82ee72b054a4

9 files changed, 55 insertions, 23 deletions

diff --git a/ChangeLog.ACC b/ChangeLog.ACC
index bd23daeb71f..0bf9d311088 100644
--- a/ChangeLog.ACC
+++ b/ChangeLog.ACC
@@ -1,4 +1,16 @@
 30-10-2013 Dmitry Bocharnikov <dmitry.b@samsung.com>
+    Add passing of kernel schedule parameters.

+    

+    * gcc/builtin-types.def: Change signature of builtin.

+    * gcc/fortran/types.def: Likewise.

+    * gcc/oacc-builtins.def: Likewise.

+    * gcc/oacc-low.c (expand_oacc_kernels): Change builtin call generation.

+    * liboacc/liboacc-internal.h: Change signature.

+    * liboacc/liboacc.c: Likewise.

+    * liboacc/liboacc.h: Likewise.

+    * liboacc/runtime.c: Pass schedule parameters to OpenCL runtime.

+
+30-10-2013 Dmitry Bocharnikov <dmitry.b@samsung.com>
     Fix formatting.
     
     * liboacc/liboacc-internal.h: Fix formatting.

diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def
index 71336f2a44a..1e555a30a3d 100644
--- a/gcc/builtin-types.def
+++ b/gcc/builtin-types.def
@@ -590,10 +590,11 @@ DEF_FUNCTION_TYPE_2 (BT_FN_VOIDPTR_CONST_STRING_INT, BT_PTR_VOID,
 DEF_FUNCTION_TYPE_2	(BT_FN_VOID_INT_VOIDPTR, BT_VOID, BT_INT, BT_PTR_VOID)
 DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VOIDPTR_UINT_VOIDPTR, BT_VOID, BT_PTR_VOID,
 			BT_UINT, BT_PTR_VOID)
-DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VOIDPTR_UINT_INT, BT_VOID, BT_PTR_VOID, BT_UINT, BT_INT)
-DEF_FUNCTION_TYPE_4 (BT_FN_VOID_VOIDPTR_UINT_VOIDPTR_UINT, BT_VOID, BT_PTR_VOID,
-			BT_UINT, BT_PTR_VOID, BT_UINT)
+DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VOIDPTR_UINT_INT, BT_VOID, BT_PTR_VOID, BT_UINT,
+     BT_INT)
 DEF_FUNCTION_TYPE_5 (BT_FN_VOIDPTR_VOIDPTR_UNIT_INT_VOIDPTR_UINT, BT_PTR_VOID,
 			BT_PTR_VOID, BT_UINT, BT_INT, BT_PTR_VOID, BT_UINT)
-DEF_FUNCTION_TYPE_5 (BT_FN_VOID_VOIDPTR_UNIT_INT_VOIDPTR_UINT, BT_VOID, BT_PTR_VOID,
-			BT_UINT, BT_INT, BT_PTR_VOID, BT_UINT)
+DEF_FUNCTION_TYPE_5 (BT_FN_VOID_VOIDPTR_UNIT_INT_VOIDPTR_UINT, BT_VOID,
+      BT_PTR_VOID, BT_UINT, BT_INT, BT_PTR_VOID, BT_UINT)
+DEF_FUNCTION_TYPE_6 (BT_FN_VOID_VOIDPTR_UINT_UINT_INT_VOIDPTR_UINT, BT_VOID,
+      BT_PTR_VOID, BT_UINT, BT_UINT, BT_INT, BT_PTR_VOID, BT_UINT)
diff --git a/gcc/fortran/types.def b/gcc/fortran/types.def
index 414260e92a6..75ddd06b516 100644
--- a/gcc/fortran/types.def
+++ b/gcc/fortran/types.def
@@ -225,10 +225,11 @@ DEF_FUNCTION_TYPE_2 (BT_FN_VOIDPTR_CONST_STRING_INT, BT_PTR_VOID,
 DEF_FUNCTION_TYPE_2	(BT_FN_VOID_INT_VOIDPTR, BT_VOID, BT_INT, BT_PTR_VOID)
 DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VOIDPTR_UINT_VOIDPTR, BT_VOID, BT_PTR_VOID,
 			BT_UINT, BT_PTR_VOID)
-DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VOIDPTR_UINT_INT, BT_VOID, BT_PTR_VOID, BT_UINT, BT_INT)
-DEF_FUNCTION_TYPE_4 (BT_FN_VOID_VOIDPTR_UINT_VOIDPTR_UINT, BT_VOID, BT_PTR_VOID,
-			BT_UINT, BT_PTR_VOID, BT_UINT)
+DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VOIDPTR_UINT_INT, BT_VOID, BT_PTR_VOID, BT_UINT,
+      BT_INT)
 DEF_FUNCTION_TYPE_5 (BT_FN_VOIDPTR_VOIDPTR_UNIT_INT_VOIDPTR_UINT, BT_PTR_VOID,
 			BT_PTR_VOID, BT_UINT, BT_INT, BT_PTR_VOID, BT_UINT)
-DEF_FUNCTION_TYPE_5 (BT_FN_VOID_VOIDPTR_UNIT_INT_VOIDPTR_UINT, BT_VOID, BT_PTR_VOID,
-			BT_UINT, BT_INT, BT_PTR_VOID, BT_UINT)
+DEF_FUNCTION_TYPE_5 (BT_FN_VOID_VOIDPTR_UNIT_INT_VOIDPTR_UINT, BT_VOID,
+    BT_PTR_VOID, BT_UINT, BT_INT, BT_PTR_VOID, BT_UINT)
+DEF_FUNCTION_TYPE_6 (BT_FN_VOID_VOIDPTR_UINT_UINT_INT_VOIDPTR_UINT, BT_VOID,
+      BT_PTR_VOID, BT_UINT, BT_UINT, BT_INT, BT_PTR_VOID, BT_UINT)
diff --git a/gcc/oacc-builtins.def b/gcc/oacc-builtins.def
index 97a25dbb893..6c06e68acca 100644
--- a/gcc/oacc-builtins.def
+++ b/gcc/oacc-builtins.def
@@ -31,19 +31,23 @@ DEF_OACC_BUILTIN (BUILT_IN_OACC_CHECK_CUR_DEV, "OACC_check_cur_dev",
                   BT_FN_VOID , ATTR_NOTHROW_LEAF_LIST)
 
 DEF_OACC_BUILTIN (BUILT_IN_OACC_GET_KERNEL, "OACC_get_kernel",
-                  BT_FN_VOIDPTR_CONST_STRING_CONST_STRING, ATTR_NOTHROW_LEAF_LIST)
+                  BT_FN_VOIDPTR_CONST_STRING_CONST_STRING,
+                  ATTR_NOTHROW_LEAF_LIST)
 
 DEF_OACC_BUILTIN (BUILT_IN_OACC_COPYIN, "OACC_copyin",
-                  BT_FN_VOIDPTR_VOIDPTR_UNIT_INT_VOIDPTR_UINT, ATTR_NOTHROW_LEAF_LIST)
+                  BT_FN_VOIDPTR_VOIDPTR_UNIT_INT_VOIDPTR_UINT,
+                  ATTR_NOTHROW_LEAF_LIST)
 
 DEF_OACC_BUILTIN (BUILT_IN_OACC_SET_KERNEL_ARG, "OACC_set_kernel_arg",
                   BT_FN_VOID_VOIDPTR_UINT_VOIDPTR, ATTR_NOTHROW_LEAF_LIST)
 
 DEF_OACC_BUILTIN (BUILT_IN_OACC_START_KERNEL, "OACC_start_kernel",
-                  BT_FN_VOID_VOIDPTR_UINT_VOIDPTR_UINT, ATTR_NOTHROW_LEAF_LIST)
+                  BT_FN_VOID_VOIDPTR_UINT_UINT_INT_VOIDPTR_UINT,
+                  ATTR_NOTHROW_LEAF_LIST)
 
 DEF_OACC_BUILTIN (BUILT_IN_OACC_COPYOUT, "OACC_copyout",
-                  BT_FN_VOID_VOIDPTR_UNIT_INT_VOIDPTR_UINT, ATTR_NOTHROW_LEAF_LIST)
+                  BT_FN_VOID_VOIDPTR_UNIT_INT_VOIDPTR_UINT,
+                  ATTR_NOTHROW_LEAF_LIST)
 
 DEF_OACC_BUILTIN (BUILT_IN_OACC_CREATE_EVENTS, "OACC_create_events",
                   BT_FN_VOIDPTR_CONST_STRING_INT, ATTR_NOTHROW_LEAF_LIST)
@@ -64,7 +68,8 @@ DEF_OACC_BUILTIN (BUILT_IN_OACC_CHECK_PRESENT, "OACC_check_present",
                   BT_FN_VOIDPTR_VOIDPTR, ATTR_NOTHROW_LEAF_LIST)
 
 DEF_OACC_BUILTIN (BUILT_IN_OACC_CREATE_ON_DEVICE, "OACC_create_on_device",
-                  BT_FN_VOIDPTR_VOIDPTR_UNIT_INT_VOIDPTR_UINT, ATTR_NOTHROW_LEAF_LIST)
+                  BT_FN_VOIDPTR_VOIDPTR_UNIT_INT_VOIDPTR_UINT,
+                  ATTR_NOTHROW_LEAF_LIST)
 
 DEF_OACC_BUILTIN (BUILT_IN_OACC_ADD_NAMED_ASYNC, "OACC_add_named_async",
                   BT_FN_VOID_INT_VOIDPTR, ATTR_NOTHROW_LEAF_LIST)
diff --git a/gcc/oacc-low.c b/gcc/oacc-low.c
index 14911973059..d27b99e9c8f 100644
--- a/gcc/oacc-low.c
+++ b/gcc/oacc-low.c
@@ -2732,13 +2732,17 @@ expand_oacc_kernels(gimple_stmt_iterator* gsi)
           }
       else
           {
-              worksize = build_int_cst(uint32_type_node, 1);
+              worksize = integer_one_node;
           }
 
       gen_add(gsi, build_call(locus,
-                    builtin_decl_explicit(BUILT_IN_OACC_START_KERNEL), 4,
-                    kernels[i]->kernel_handle, worksize, queue_handle,
-                    build_int_cst(uint32_type_node, 0)));
+                    builtin_decl_explicit(BUILT_IN_OACC_START_KERNEL), 6,
+                    kernels[i]->kernel_handle,
+                    worksize,               /* WORKITEMS */
+                    integer_zero_node,      /* OFFSET */
+                    integer_minus_one_node, /* GROUPSIZE */
+                    queue_handle,
+                    integer_zero_node));
       gen_add(gsi, build_call(locus,
                   builtin_decl_explicit(BUILT_IN_OACC_ADVANCE_EVENTS), 1,
                   queue_handle));
diff --git a/liboacc/liboacc-internal.h b/liboacc/liboacc-internal.h
index 860d5372603..fcca756b7c8 100755
--- a/liboacc/liboacc-internal.h
+++ b/liboacc/liboacc-internal.h
@@ -250,6 +250,8 @@ extern struct OACC_kernel_data* OACC_find_kernel(const char* prog_name,
                                                  const char* kern_name);
 extern void OACC_enqueue_kernel(struct OACC_kernel_data* kern,
                                 unsigned worksize,
+                                unsigned offset,
+                                int groupsize, 
                                 struct OACC_queue_data* queue,
                                 unsigned idx);
 extern void OACC_set_arg_buf(struct OACC_kernel_data* kern,
diff --git a/liboacc/liboacc.c b/liboacc/liboacc.c
index 9d8f060daa6..469fe0d8cd1 100755
--- a/liboacc/liboacc.c
+++ b/liboacc/liboacc.c
@@ -55,14 +55,15 @@ OACC_get_kernel(const char* prog_name, const char* kern_name)
 
 /* enqueue kernel KERN for WORKSIZE threads */
 void
-OACC_start_kernel(oacc_kernel kern, unsigned worksize, oacc_event ev,
+OACC_start_kernel(oacc_kernel kern, unsigned worksize,
+                  unsigned offset, int groupsize, oacc_event ev,
                   unsigned idx)
 {
   OACC_CHECK_KERN_PTR(kern)
   OACC_CHECK_INIT
   OACC_CHECK_DEV_INIT(OACC_curr_dev[OACC_curr_num])
   OACC_enqueue_kernel((struct OACC_kernel_data*)kern, worksize,
-  (struct OACC_queue_data*)ev, idx);
+         offset, groupsize, (struct OACC_queue_data*)ev, idx);
 }
 
 oacc_buffer
diff --git a/liboacc/liboacc.h b/liboacc/liboacc.h
index 4b6d198417d..0f373e87e4e 100755
--- a/liboacc/liboacc.h
+++ b/liboacc/liboacc.h
@@ -43,7 +43,8 @@ void OACC_check_cur_dev(void);
 /* create or get kernel from cache */
 oacc_kernel OACC_get_kernel(const char* prog_name, const char* kern_name);
 /* start a kernel with specified worksize */
-void OACC_start_kernel(oacc_kernel kernel, unsigned worksize, oacc_event ev,
+void OACC_start_kernel(oacc_kernel kernel, unsigned worksize,
+                       unsigned offset, int groupsize, oacc_event ev,
                        unsigned ev_idx);
 /* associate memory object with kernel argument */
 void OACC_set_kernel_arg(oacc_kernel kern, unsigned idx, oacc_buffer buf);
diff --git a/liboacc/runtime.c b/liboacc/runtime.c
index 52f8857024a..d3ccf6ea3d5 100755
--- a/liboacc/runtime.c
+++ b/liboacc/runtime.c
@@ -129,16 +129,21 @@ OACC_dev_fini(OACC_device_ptr pdev)
 /* enqueue kernel object */
 void
 OACC_enqueue_kernel(struct OACC_kernel_data* kern, unsigned worksize,
+                    unsigned offset, int groupsize,
                     struct OACC_queue_data* queue, unsigned idx)
 {
   cl_int err;
   cl_uint nev;
   cl_event *ev_wait = NULL, *ev_set = NULL;
+  unsigned *ploc_size = NULL;
+
+  if(groupsize > 0)
+    ploc_size = (unsigned *)&groupsize;
 
   OACC_EQ_EVENTS(queue, nev, ev_wait, ev_set, idx)
 
   err = clEnqueueNDRangeKernel(OACC_curr_dev[OACC_curr_num]->queue, kern->kern,
-                  1 ,NULL, &worksize, NULL, nev, ev_wait, ev_set);
+                  1, &offset, &worksize, ploc_size, nev, ev_wait, ev_set);
   if(err < 0)
     {
       OACC_fatal("Can't enqueue kernel: %d\n", err);