Fixed conversion error

*    ChangeLog.ACC
*    gcc/oacc-builtins.def
*    gcc/fortran/types.def
*    gcc/oacc-low.c
*    gcc/builtin-types.def



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

5 files changed, 45 insertions, 11 deletions

diff --git a/ChangeLog.ACC b/ChangeLog.ACC
index 9599e50f7ba..4fe242bdd3d 100644
--- a/ChangeLog.ACC
+++ b/ChangeLog.ACC
@@ -1,3 +1,11 @@
+09-10-2013 Ilmir Usmanov <i.usmanov@samsung.com>
+    Fix return and param types of get_global_id() builtin 
+    
+    * gcc/builtin-types.def: Add new function type size_t(unsigned int)
+    * gcc/fortran/types.def: Likewise
+    * gcc/oacc-builtins.def: Change type of get_global_id() function
+    * gcc/oacc-low.c: Add convert expression to ssa form
+
 08-10-2013 Dmitry Bocharnikov <dmitry.b@samsung.com>
     Add check for irreducible loops
     
diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def
index 28cd9d32ff8..71336f2a44a 100644
--- a/gcc/builtin-types.def
+++ b/gcc/builtin-types.def
@@ -230,6 +230,7 @@ DEF_FUNCTION_TYPE_1 (BT_FN_ULONGLONG_ULONGLONG, BT_ULONGLONG, BT_ULONGLONG)
 DEF_FUNCTION_TYPE_1 (BT_FN_UINT16_UINT16, BT_UINT16, BT_UINT16)
 DEF_FUNCTION_TYPE_1 (BT_FN_UINT32_UINT32, BT_UINT32, BT_UINT32)
 DEF_FUNCTION_TYPE_1 (BT_FN_UINT64_UINT64, BT_UINT64, BT_UINT64)
+DEF_FUNCTION_TYPE_1 (BT_FN_SIZE_UINT, BT_SIZE, BT_UINT)
 
 DEF_POINTER_TYPE (BT_PTR_FN_VOID_PTR, BT_FN_VOID_PTR)
 
diff --git a/gcc/fortran/types.def b/gcc/fortran/types.def
index c32134a14da..414260e92a6 100644
--- a/gcc/fortran/types.def
+++ b/gcc/fortran/types.def
@@ -89,6 +89,7 @@ DEF_FUNCTION_TYPE_1 (BT_FN_VOID_VPTR, BT_VOID, BT_VOLATILE_PTR)
 DEF_FUNCTION_TYPE_1 (BT_FN_UINT_UINT, BT_UINT, BT_UINT)
 DEF_FUNCTION_TYPE_1 (BT_FN_PTR_PTR, BT_PTR, BT_PTR)
 DEF_FUNCTION_TYPE_1 (BT_FN_VOID_INT, BT_VOID, BT_INT)
+DEF_FUNCTION_TYPE_1 (BT_FN_SIZE_UINT, BT_SIZE, BT_UINT)
 
 
 DEF_POINTER_TYPE (BT_PTR_FN_VOID_PTR, BT_FN_VOID_PTR)
diff --git a/gcc/oacc-builtins.def b/gcc/oacc-builtins.def
index 7573cdc0efd..97a25dbb893 100644
--- a/gcc/oacc-builtins.def
+++ b/gcc/oacc-builtins.def
@@ -25,7 +25,7 @@
    See builtins.def for details.  */
 
 DEF_OACC_BUILTIN (BUILT_IN_OACC_GET_GLOBAL_ID, "get_global_id",
-                  BT_FN_INT_INT, ATTR_CONST_NOTHROW_LEAF_LIST)
+                  BT_FN_SIZE_UINT, ATTR_CONST_NOTHROW_LEAF_LIST)
 
 DEF_OACC_BUILTIN (BUILT_IN_OACC_CHECK_CUR_DEV, "OACC_check_cur_dev",
                   BT_FN_VOID , ATTR_NOTHROW_LEAF_LIST)
diff --git a/gcc/oacc-low.c b/gcc/oacc-low.c
index 122baa2eb5c..027acb4ee46 100644
--- a/gcc/oacc-low.c
+++ b/gcc/oacc-low.c
@@ -1263,7 +1263,6 @@ dump_fn_body(FILE* dump_file, const char* title)
   fflush(dump_file);
 }
 
-
 static void
 parallelize_loop(struct loop* l)
 {
@@ -1436,21 +1435,45 @@ parallelize_loop(struct loop* l)
           }
           gimple stmt = gsi_stmt(gsi);
           if(gimple_code(stmt) != GIMPLE_PHI) {
+        	  location_t location = gimple_location(stmt);
               tree lhs = get_gimple_def_var(stmt);
-              gimple call_stmt = build_call(gimple_location(stmt), builtin_decl_explicit(BUILT_IN_OACC_GET_GLOBAL_ID), 1, build_int_cst(TREE_TYPE(lhs), 0));
+              tree builtin_decl = builtin_decl_explicit(BUILT_IN_OACC_GET_GLOBAL_ID);
+              tree builtin_return_type = TREE_TYPE (TREE_TYPE (builtin_decl));
+              //FIXME: add more conversion magic???
+              gcc_assert(TREE_CODE(builtin_return_type) == TREE_CODE(TREE_TYPE(lhs)));
+
+              int builtin_unsigned = TYPE_UNSIGNED(builtin_return_type);
+              int lhs_unsigned = TYPE_UNSIGNED(TREE_TYPE(lhs));
 
-              tree tmp_var = create_tmp_reg(TREE_TYPE(lhs), "_oacc_tmp");
+              /* _oacc_tmp = __builtin_get_global_id (0); */
+              gimple call_stmt = build_call(location, builtin_decl, 1, build_int_cst(builtin_return_type, 0));
+              tree tmp_var = create_tmp_reg(builtin_return_type, "_oacc_tmp");
               tmp_var = make_ssa_name_fn(cfun, tmp_var, call_stmt);
               set_gimple_def_var(call_stmt, tmp_var);
-
-              gimple add_stmt = build_assign(PLUS_EXPR, tmp_var, lhs);
-              tree new_var = make_ssa_name_fn(cfun, SSA_NAME_VAR(lhs), add_stmt);
-              set_gimple_def_var(add_stmt, new_var);
-
               gen_add(&gsi, call_stmt);
-              gen_add(&gsi, add_stmt);
-
 
+              if (builtin_unsigned != lhs_unsigned)
+                {
+            	  /* _ = (int) _oacc_tmp; */
+            	  gimple convert_stmt = build_type_cast (TREE_TYPE(lhs), tmp_var);
+            	  tree convert_var = make_ssa_name_fn(cfun, SSA_NAME_VAR(lhs), convert_stmt);
+				  set_gimple_def_var(convert_stmt, convert_var);
+            	  gen_add (&gsi, convert_stmt);
+
+            	  /* i = _ + i; */
+            	  gimple add_stmt = build_assign(PLUS_EXPR, convert_var, lhs);
+				  tree new_var = make_ssa_name_fn(cfun, SSA_NAME_VAR(lhs), add_stmt);
+				  set_gimple_def_var(add_stmt, new_var);
+	              gen_add(&gsi, add_stmt);
+                }
+              else
+                {
+            	  /* i = _oacc_tmp + i; */
+            	  gimple add_stmt = build_assign(PLUS_EXPR, tmp_var, lhs);
+				  tree new_var = make_ssa_name_fn(cfun, SSA_NAME_VAR(lhs), add_stmt);
+				  set_gimple_def_var(add_stmt, new_var);
+	              gen_add(&gsi, add_stmt);
+                }
           }
       }
   }
@@ -1505,6 +1528,7 @@ expand_oacc_kernels(gimple_stmt_iterator* gsi)
           gimple_stmt_iterator gsi;
           for(gsi = gsi_start_bb(bb); !gsi_end_p(gsi); gsi_next(&gsi)) {
               gimple stmt = gsi_stmt(gsi);
+
               gimple_set_modified (stmt, true);
               update_stmt_operands(stmt);
           }