This patch has some new execution tests, verifying loop partitioning is behaving
as expected.
There are more execution tests on the gomp4 branch, but many of them use
reductions. We'll merge those once reductions are merged.
nathan
2015-10-20 Nathan Sidwell <nat...@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/loop-g-1.c: New.
* testsuite/libgomp.oacc-c-c++-common/loop-w-1.c: New.
* testsuite/libgomp.oacc-c-c++-common/loop-g-1.s: New.
* testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c: New.
* testsuite/libgomp.oacc-c-c++-common/loop-g-2.c: New.
* testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c: New.
* testsuite/libgomp.oacc-c-c++-common/loop-v-1.c: New.
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c (revision 0)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c (working copy)
@@ -0,0 +1,57 @@
+/* { dg-do run } */
+/* { dg-additional-options "-O2" */
+
+#include <stdio.h>
+
+#define N (32*32*32+17)
+int main ()
+{
+ int ary[N];
+ int ix;
+ int exit = 0;
+ int ondev = 0;
+
+ for (ix = 0; ix < N;ix++)
+ ary[ix] = -1;
+
+#pragma acc parallel num_gangs(32) vector_length(32) copy(ary) copy(ondev)
+ {
+#pragma acc loop gang
+ for (unsigned ix = 0; ix < N; ix++)
+ {
+ if (__builtin_acc_on_device (5))
+ {
+ int g = 0, w = 0, v = 0;
+
+ __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g));
+ __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w));
+ __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v));
+ ary[ix] = (g << 16) | (w << 8) | v;
+ ondev = 1;
+ }
+ else
+ ary[ix] = ix;
+ }
+ }
+
+ for (ix = 0; ix < N; ix++)
+ {
+ int expected = ix;
+ if(ondev)
+ {
+ int g = ix / ((N + 31) / 32);
+ int w = 0;
+ int v = 0;
+
+ expected = (g << 16) | (w << 8) | v;
+ }
+
+ if (ary[ix] != expected)
+ {
+ exit = 1;
+ printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected);
+ }
+ }
+
+ return exit;
+}
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.s
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.s (revision 0)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.s (working copy)
@@ -0,0 +1,386 @@
+ .file "loop-g-1.c"
+ .section .gnu.offload_lto_.inline.f031cb8759bb7418,"e",@progbits
+ .string "x\234ce\200"
+ .string "i\006\004`d`P\220g``\262zp\205\231\201\205\201\t,\306\310\004\022\007"
+ .ascii "-\377\002_"
+ .text
+ .section .gnu.offload_lto_main._omp_fn.0.f031cb8759bb7418,"e",@progbits
+ .ascii "x\234\215W\373STG\026\276\347\366\2357\f\304\020b\310h0\331\321"
+ .ascii "$\204\031\311V\355\243\222\335\252\255\375i\377\201T~\263\310"
+ .ascii "@\b\273\300X0\032\363\023\027\034^E6\006bH\234\302\205\215(\020"
+ .ascii "\214H\214\213\016.\316\202\n\270\3403\260\242\361\301#\nb\214"
+ .ascii "\004\\$\b\3549\335\227a\200Y\365V\335\236\333\347~\337\351\257"
+ .ascii "O\237>}G'\211K\266J\222\003$)\233:f\td\to`\037\237\301_&+\222"
+ .ascii "\244\312Fl\024E6\320\263\236\236u!v=\350\020i\340\255Q6\021O"
+ .ascii "U\030\2756\001\272\225@\312\225\231\242\323\033\214&\310eL\033"
+ .ascii "U\321~?d;\031\264\025\335<*\277R\253\342u\243\264~U\236Z\260"
+ .ascii "\272\022\201\025\032\030\366\300+\210\362\265L\236\326\225\262"
+ .ascii "J\006`\304n\363\335\206\346H\352\312\020\220\260\377`\376r\233"
+ .ascii "\201\025\342\230\330\351\274t\266\n6\354U?\316W\253;U(F\030\203"
+ .ascii "}\300\320%\354\005\033\324\200M\252d\212\004\265\260\201\360"
+ .ascii "\347\275\007+\231\360\255`w\274\3077d\025\276\343\261\373\311"
+ .ascii "x\211\314\312\205\347\233\245\027K\242\354{\204\326\256\350\362"
+ .ascii "\374\202NU\251d:\370Rx\257G\357u\350\235.\204\227\r"
+ .string "\365\324D\bx\365\324\250RV\027\360\306#Zz\f\364v\311\037\362\312Ux\002\344\334\023#\017<1\262e\234=Jh\343\235\243\323<\362z\036\221\212\202\266\263\206R\036a\352n\357\231\373S\025\354\316S\243+\231\201\342\013_\023\225\3028\263w4\322\231H\274CP\005\2372\315_\351\341\341V\206\024\222\006\337\200\023\2324\202\267cG\263\314>\023\201o\236\275x\fP`\200\002_i-S\003\336uD\370\307\322\270\303at\2559\016*3>F\331\350\376\313\n\t3\302Q.\f\004\277\255\243\267V'\204\231\300\217\302\216h\370\017?\272\327\250\210t9\241R2V6\365\316\311\"a\252\275<\035g\206\232\017\033\330n\241\275\245\355\314\016\235\320^\375\275Q-\367r\217\255B\371?Q\3711R\336\262\250|\260i\356\323H\275\002\307\301\016\315x\037\304\273\021\357\003x\177\205\367~\274\033\300\316wG"
+ .string "p\035\244Xd\235\034\367W"
+ .ascii "\373\273\030\263~\276\244\"b!Q\033V\225i\211j\206\177\255LTh"
+ .ascii "\003;R\246\206\317\317\232\331W\202_\323<\277S\257%\304\211j"
+ .ascii "\320\310\355a\310'8Y\235\2716\027\305\374\202\\\374\267\276\021"
+ .ascii "\266\202|2\f\371\024'O\\\237h\265\2626A\036\031\236\270$\257"
+ .ascii " w\204!wr\262\377\352\221[Q\254[\220\217\237\275\360\302\nnW"
+ .ascii "\030\356i\316\275\032\270^``}\202\273\353\312\314\004W\235\257"
+ .ascii "\006\212w\306\323\006\260\300\277\303p\2739\267qv\270O\317\006"
+ .ascii "\004\267\244a\327\035\203\340\336(\335\037[\236\257\362\261-"
+ .ascii "\320\023\206\177F[\322\006\021\366\376\257k,\202Z]{\001\312\353"
+ .ascii "\002\237\023\325(\205V+\316;\313\361\355u\375\275z6,\306\355"
+ .ascii "\n\324\345Gj\023\036(\001U\354p8\027f\324\363\234\335\177\356"
+ .ascii "`"
+ .string "\251\225\335\024\354\351\2077\2765k\354}\213\354\013a\330\0279\273pO\373t$\033\021l\357\241\003G\264ZX=\270\310\3766\f\273we\002/\233\371\324\300\243f\336'F\357\231\274\246h\370\231\351 ^\277\022\377\037\034C\333<K\027\247\372P\277L\265#\376\377\255\315%\2607\226C@\235\237\337\256\252\215\277(g\t\200G\205Y\246\306BM\0045\221\324X\251\211\242&\032\033\313S|aW\341\243\356i\036\355\030\036\265g\370\354c\271\345YnY\315-\317\0210\216\232\347\251\241\003I\267\006\250.\255\305Gc<6\372\027\311\370\022u\3554\314zj6\360a$\251\210\327\317b\206\0076\033\203\302\342\366O\336Z\254\034\274,,)%8/\t\372\227\031\311\021\\&#\\\241V\202\327l\220`\243#\360\341\356f\023U4\035\304\300w`\213\223\306\240\240\370\217\032\f\001\337\234\350(4\022\300\002"
+ .string "W1\202\211\034\343\rb\320U\2344\n\371\036\311\307,\241\325\330\f\022\\\303p\217\301.\265\311Hh\037\213@\374\2534\356\341\242\356&\263@\001\\\0072\265\236\372\342ea)Se\270\201C\t\373|\375]\277^\274\310\205\0014\277j[\263\304i$:M\340\350\005V\020F\201c<p\276 \334\312\347ur~*\020%&.\303 \204\243\373\030\320@\213\314(\021\270pH\353\032i\030\274g\030\006\301\312\201\222t\037"
+ .ascii "#\"\243e\275\221t\340\267\020\3038\375d@K\264\017\327s\024\266"
+ .ascii "W\350\261\363T\350ak\301\250\r\001z\214\246(\327\035\0338\031"
+ .ascii "\025\272\334W\226\256\254\004\303|U\277'y\253\350\360\373\371"
+ .ascii "\370\2045wV\305/\253\210\n:\340\340&\344\362\002R\343\177\030"
+ .ascii "!\346\253\203[`\253\304\260\241\342+\3551\032\031\247!T\3138"
+ .ascii "\005v\355R9\3631\273\216\204\353\270p\241\325\216\302}\202\312"
+ .ascii "'k\327&\2130\2756\277\242\306\2326Y8}\021'G\271\341\357\212\023"
+ .ascii "\206\227PI\355\271\300>s\360\200\366\2615\301\363\271\242\370"
+ .ascii "\257vq>\217\004\317g\030\005'\271\250\350z\231\\P5\3518]\034"
+ .ascii "\241\035\327\3501xZ\267\364\336\337!'5\3622\365g\232\347mtB\002"
+ .ascii "\310\313\030$\241\034\256\323 tN\216\034y3d\213\340\256\245="
+ .ascii "K;\226\207\0239wH\361\323\024\027\202w&\nCL\320\360\2460<C\206"
+ .ascii "\333\220\367\305X\277\201,\370)\214]\232\316G\025U\312\272<\271"
+ .ascii "J\375\034\313\220\2254\316\016\026\f\231\024\372\366\376\001"
+ .ascii "\326q\325\247\006\017\371\365h\201\273\240\370X,\276\371\021"
+ .ascii "\024\356\356\307\256\240\273\002\371q\356\356-w7N\356\236\305"
+ .ascii "7?i\356\306a\301\233\357\261\336&\226{\233$o\253\361\315}\364"
+ .ascii "\206\271p\257\360\352-\263\230\376s\250\032\241\027\356u\357"
+ .ascii "\266\210\374\212\206\377\202M\240\246\r\002\024\207ZB\277\355"
+ .ascii "u`\204)\001zPX\257\201\236\347\316\342\204\261\337$\2146n_-\262"
+ .ascii "\310&lk\2370\213\036,f\321\364B\026\275M.Vd\321\332\260Y\224"
+ .ascii "MB\177\016\311\242\031\312\"\276\372\376\321\326\337\0131\361"
+ .ascii "\217\024\323}\252\366K&>\206\037\006?\206aV\250\2314jb\232\275"
+ .ascii "\307&\025\355\323\030=\212/c\230C\002\0266\236\264XA\330\024"
+ .ascii "\315da#\256\307\352D\025(\270\023\327\363\352\312\301\246\220"
+ .ascii "\r\273\201\252\t\375\365\351\333\327\244\210\002C\345u\236\027"
+ .ascii "\346\202\342\366\317\372R\026\213}\360\377\221\016\024Pe:#V\034"
+ .ascii "\b\301\302i\301\302\231'\2079\020\220\030'\375@\3770\351*\242"
+ .ascii "\005\300\347|p\346\270\262\223=\256\367\234Y9\351)\357\247fd"
+ .ascii "8\335\233S\263\222].\247'{K\326_\0223S\263\323R\2359\331.g\232"
+ .ascii "\313\225\230\231\234\236\225\221\236\225\352\314H\177'\315\235"
+ .ascii "\271\331\351I\315\361\344lI\367\004-\0167r\023]\t\tN\207c\231"
+ .ascii "\215\254\211.wf\246;\313\231\341voNLKLr\270$\213\003!\233R\222"
+ .ascii "=\311\233\322\225\364m\216\215:wVJ\352V\226\234\375\201\305\341"
+ .ascii "z\017El\312r;~\023\241=g&os\374\332\340\310\361\244nv\374\312"
+ .ascii "\344p\277\373nN\252\307\361[\223\343\035\367\226\254\024G\322"
+ .ascii "F\363\202-)IN\337\006i\360"
+ .ascii ">l\215\315toul\371\345\353\361\366\215\257\331\355.Orz\212c\333"
+ .ascii "\033\022\373]\266\024\263\344\225\007_|\360F\030#\242\377\007"
+ .ascii "!Kr\335"
+ .text
+ .section .gnu.offload_lto_.symbol_nodes.f031cb8759bb7418,"e",@progbits
+ .string "x\234ce``\320\003b\006&\236z\006\206\t\347\030\030\200\324\212\205\013\0170300\362\3263\202\205\030\030\032\032\024\030\030\230\031\030\031\216\264\277\231\317\301"
+ .string ""
+ .ascii "\004N\0139"
+ .text
+ .section .gnu.offload_lto_.refs.f031cb8759bb7418,"e",@progbits
+ .string "x\234ce```\004b\006"
+ .string ""
+ .string ";"
+ .ascii "\007"
+ .text
+ .section .gnu.offload_lto_.offload_table.f031cb8759bb7418,"e",@progbits
+ .string "x\234ce```\006bF\006\006"
+ .string ""
+ .string "Z"
+ .ascii "\n"
+ .text
+ .section .gnu.offload_lto_.decls.f031cb8759bb7418,"e",@progbits
+ .ascii "x\234\215T\337O\333U\024\377\236\336oa\226\2262@C\f\017d!\031"
+ .ascii "\311\322v\350\037\240\017>\360\270\355\3057I\375\322\261F\370"
+ .ascii "\226\264_4{\362\322\221XA\035L\030Jp\351\346F\221!k\351X\367"
+ .ascii "\013\2500\030l\300\306&\242\213 \272\200/\023\331d\262lq\365"
+ .ascii "\334{\373\205\002\242\336\344\334~\317\271\237\3639\347\334{"
+ .ascii "N\215\222X?\247I\322\f\376~\206\222\300e\300\337\003\322\372"
+ .ascii "bz\030%\"m]\314nM\342\263QrQ\206Q\366\243\354C)A\031B\351G\351"
+ .ascii "C\031E\331\2152\2012\216r/\311?\200\022\377\217XM\004\306\347"
+ .ascii "\342\2473\346\340G@\025\330\302}#\b\222E\360\005\034\245\327"
+ .ascii " \255A7;I lI\360\026\312T\302u\216d|\236\213\b\260\201\017\203"
+ .ascii "nJ\006\326\230ak\246\234\b\364L6\227$ID\277\003\332\031x\f?\301"
+ .ascii ",p\255g52b&\030\t\277\037?j9m%;\371\367\330\350\261\363f\331"
+ .ascii " A;\020\341\367p\242\365~f=\201}&(@\365\344\374\303f\006\220"
+ .ascii "\240\003\3629\340\346\344\314\220\221Y$\370\ndfy\336\031\355"
+ .ascii "5\350\030\350\024\306\225\313\253\263F\262\233GI4v\f\247\261"
+ .ascii "(]@\340,\310\020\022\230\356\256ga\023w\334o\206\267\270\351"
+ .ascii "\367\225\346\017w\344\035\247\270\226kh\026\300\201\227\340\325"
+ .ascii "|\330\213\261\341\034\344\341\336\215\276<\221\341\253\023\027"
+ .ascii "\255\273\2024\036\244\363\261O\215\237\237\210\323#\224\326\310"
+ .ascii "\f\007gx=\3605\006\343\237\360%\314\341\265\205a\027z\266/\236"
+ .ascii "\250I\253#$\232&A\204[Z\257\254\3340b\325\331&\330\201j"
+ .string "l\351l\314\202j\256\t\342,T\355\365\2431\003\271\004\274\232?f\273\226\254d@(w\226\307\277\310\020\216Y\250^X\034\0331\tG\n\250?\273\037\353M'\023\002{t\241\277\305\"\260\r\324\200\206\245\310jM\232)\300\212\r\372\351\316\232bv\024\305;\354aw\r\227\301\204\273#\037.\241\212\360\217>Y>'\013\377k\224\245\331\026\231~n\020\321\202\265<\321\370J\364\003\310n`\214\270\025\310\354\031\257\"\341\025N\330\007\331:-\007O\266\005\300\\G\251_\017\217\200\224\370\375`\026\315$7O\306#2\024\026\346\024\235b\350\371\330\020\t\371i\340H:\224\230\341<>\3527\220\263\346\210\031\243s1n\022G\007\027\032\200\3762\017\214\017.\376#8\312\342I\215\224\006h\330O_\317\223 \206\347\027\330\371 \0241\007\004\r@!\304\241\2207x;\253\220\277:\215gq\302^ )\377"
+ .ascii "\277u\207Z2r\222\204\301\377A\210>\017\236\034{\323R\237r\033"
+ .ascii "C)yJp\r,I\340\355\216!j\325\221\376m\2210\f\026DO\325\206\333"
+ .ascii "\210x59e\276r\371|\221\256xS\263\025L\371\200\336\300;\241\301"
+ .ascii "O_(\221\340:\032F8\335(o\203\033\250\217\241\336\212%\321\345"
+ .ascii "D\342\375D\"\235\037q\003m<\365F\240\361en`\333\206\341\302\036"
+ .ascii "\330f\270\360D\037\256q>\\7\365\341\252\3778\324\277\2617\020"
+ .ascii "\220R\343\204\350\r\366\313F1\261\032\267\212\032Y_\327>8^g\331"
+ .ascii "\276\257o\351}=)\372z\360Q8d\026\336\204ME\355\302|\206\270\241"
+ .ascii "<T'Cg\232LdJ\214\320\364\223\276\247fr\227+\362X\357\323\036"
+ .ascii "\302\357\256H\277\274\345\223\224\340c\213\242\340;|\352oy\260"
+ .ascii "i\f>\243\337\036E2fn\027c\322 KL\205\273\311\016\202;\354C\334"
+ .ascii "\363k\034\310ns\365\327\2772\327\007\213\023\300\024\262\336"
+ .ascii "\346%}\317'\353\226>Y-=\177\006eS2\243\300\277f\204\350\305p"
+ .ascii "t\320\222\312\375\303\266\334p\017\225\314J\247[\265\227z*\253"
+ .ascii "J\017\252\366\275/\342G\201\346\364\226\273\264\002\227\252y"
+ .ascii "\017Wy\334\252f\3618\025\245\340`\265\252hn\217\352\007\207O"
+ .ascii "\361:5\345\220C\365\271\313\336sUT8<U.\0251\016\315[\255\276"
+ .ascii "c\253t!\201\303\347U\034\345\212bc\021*\334\252\313Q\341~\273"
+ .ascii "\034\371\035\232\313\247\371\252\335\332\232\305\316\370m\312"
+ .ascii "\236=\016\273}\223\215Ym\212\247\262\322\243:*<\236*[\271\255"
+ .ascii "\330\256HF\217Z\346z\327bgy\22795g\251f\177\2058\275\207A\221"
+ .ascii "\225CN\357\337(\266/\210"
+ .text
+ .section .gnu.offload_lto_.symtab.f031cb8759bb7418,"e",@progbits
+ .text
+ .section .gnu.offload_lto_.opts,"e",@progbits
+ .string "'-fexceptions' '-fmath-errno' '-fsigned-zeros' '-ftrapping-math' '-fno-trapv' '-fno-strict-overflow' '-fno-openmp' '-foffload-abi=lp64' '-fopenacc'"
+ .text
+ .section .gnu.offload_lto_.mode_table.f031cb8759bb7418,"e",@progbits
+ .string "x\234ce\200"
+ .string "e \026"
+ .string "\342\376#\r\035\r{:\004&\2664-h8\322\0210\251\245\345@\303\211\216\t\314\223[:\032\032\317t\\`f`\016\364d`\016\006b\027 \016\361d"
+ .string ""
+ .ascii "\225\020\024\253"
+ .text
+ .section .rodata
+.LC0:
+ .string "ary[%d]=%x expected %x\n"
+ .text
+ .globl main
+ .type main, @function
+main:
+.LFB11:
+ .cfi_startproc
+ pushq %rbp
+ .cfi_def_cfa_offset 16
+ .cfi_offset 6, -16
+ movq %rsp, %rbp
+ .cfi_def_cfa_register 6
+ subq $131216, %rsp
+ movl $0, -8(%rbp)
+ movl $0, -131204(%rbp)
+ movl $0, -4(%rbp)
+.L3:
+ cmpl $32784, -4(%rbp)
+ jg .L2
+ movl -4(%rbp), %eax
+ cltq
+ movl $-1, -131200(%rbp,%rax,4)
+ addl $1, -4(%rbp)
+ jmp .L3
+.L2:
+ leaq -131204(%rbp), %rax
+ movq %rax, -48(%rbp)
+ leaq -131200(%rbp), %rax
+ movq %rax, -40(%rbp)
+ leaq -48(%rbp), %rax
+ subq $8, %rsp
+ pushq $0
+ movl $_ZZ4mainE17.omp_data_kinds.5, %r9d
+ movl $_ZZ4mainE17.omp_data_sizes.4, %r8d
+ movq %rax, %rcx
+ movl $2, %edx
+ movl $main._omp_fn.0, %esi
+ movl $-1, %edi
+ movl $0, %eax
+ call GOACC_parallel_keyed
+ addq $16, %rsp
+ movl $0, -4(%rbp)
+.L7:
+ cmpl $32784, -4(%rbp)
+ jg .L4
+ movl -4(%rbp), %eax
+ movl %eax, -12(%rbp)
+ movl -131204(%rbp), %eax
+ testl %eax, %eax
+ je .L5
+ movl -4(%rbp), %eax
+ movslq %eax, %rdx
+ imulq $2145388543, %rdx, %rdx
+ shrq $32, %rdx
+ sarl $9, %edx
+ sarl $31, %eax
+ subl %eax, %edx
+ movl %edx, %eax
+ movl %eax, -16(%rbp)
+ movl $0, -20(%rbp)
+ movl $0, -24(%rbp)
+ movl -16(%rbp), %eax
+ sall $16, %eax
+ movl %eax, %edx
+ movl -20(%rbp), %eax
+ sall $8, %eax
+ orl %edx, %eax
+ orl -24(%rbp), %eax
+ movl %eax, -12(%rbp)
+.L5:
+ movl -4(%rbp), %eax
+ cltq
+ movl -131200(%rbp,%rax,4), %eax
+ cmpl -12(%rbp), %eax
+ je .L6
+ movl $1, -8(%rbp)
+ movl -4(%rbp), %eax
+ cltq
+ movl -131200(%rbp,%rax,4), %edx
+ movl -12(%rbp), %ecx
+ movl -4(%rbp), %eax
+ movl %eax, %esi
+ movl $.LC0, %edi
+ movl $0, %eax
+ call printf
+.L6:
+ addl $1, -4(%rbp)
+ jmp .L7
+.L4:
+ movl -8(%rbp), %eax
+ leave
+ .cfi_def_cfa 7, 8
+ ret
+ .cfi_endproc
+.LFE11:
+ .size main, .-main
+ .type main._omp_fn.0, @function
+main._omp_fn.0:
+.LFB12:
+ .cfi_startproc
+ pushq %rbp
+ .cfi_def_cfa_offset 16
+ .cfi_offset 6, -16
+ movq %rsp, %rbp
+ .cfi_def_cfa_register 6
+ pushq %r15
+ pushq %r14
+ pushq %r13
+ pushq %r12
+ pushq %rbx
+ subq $40, %rsp
+ .cfi_offset 15, -24
+ .cfi_offset 14, -32
+ .cfi_offset 13, -40
+ .cfi_offset 12, -48
+ .cfi_offset 3, -56
+ movq %rdi, -72(%rbp)
+ movl $0, %r12d
+ movl $1, %r14d
+ movl $1, %eax
+ movl %eax, %r15d
+.L15:
+ movl $0, %eax
+ movl %eax, %ebx
+ movl $32785, %r13d
+ cmpl %r13d, %ebx
+ jge .L10
+.L13:
+ movl %ebx, %eax
+ movl %eax, -52(%rbp)
+ movl $5, %edi
+ call acc_on_device
+ testl %eax, %eax
+ jne .L11
+ jmp .L16
+.L14:
+ addl %r15d, %ebx
+ cmpl %r13d, %ebx
+ jl .L13
+ jmp .L10
+.L16:
+ movl -52(%rbp), %ecx
+ movq -72(%rbp), %rax
+ movq 8(%rax), %rax
+ movl -52(%rbp), %edx
+ movl %ecx, (%rax,%rdx,4)
+ jmp .L14
+.L11:
+ movl $0, -56(%rbp)
+ movl $0, -60(%rbp)
+ movl $0, -64(%rbp)
+#APP
+# 26 "/scratch/nsidwell/openacc/trunk-merge/src/gcc-mainline/libgomp/testsuite/libgomp.oacc-c++/../libgomp.oacc-c-c++-common/loop-g-1.c" 1
+ mov.u32 %eax,%ctaid.x;
+# 0 "" 2
+#NO_APP
+ movl %eax, -56(%rbp)
+#APP
+# 27 "/scratch/nsidwell/openacc/trunk-merge/src/gcc-mainline/libgomp/testsuite/libgomp.oacc-c++/../libgomp.oacc-c-c++-common/loop-g-1.c" 1
+ mov.u32 %eax,%tid.y;
+# 0 "" 2
+#NO_APP
+ movl %eax, -60(%rbp)
+#APP
+# 28 "/scratch/nsidwell/openacc/trunk-merge/src/gcc-mainline/libgomp/testsuite/libgomp.oacc-c++/../libgomp.oacc-c-c++-common/loop-g-1.c" 1
+ mov.u32 %eax,%tid.x;
+# 0 "" 2
+#NO_APP
+ movl %eax, -64(%rbp)
+ movl -56(%rbp), %eax
+ sall $16, %eax
+ movl %eax, %edx
+ movl -60(%rbp), %eax
+ sall $8, %eax
+ orl %edx, %eax
+ orl -64(%rbp), %eax
+ movl %eax, %ecx
+ movq -72(%rbp), %rax
+ movq 8(%rax), %rax
+ movl -52(%rbp), %edx
+ movl %ecx, (%rax,%rdx,4)
+ movq -72(%rbp), %rax
+ movq (%rax), %rax
+ movl $1, (%rax)
+ jmp .L14
+.L10:
+ addl $1, %r12d
+ cmpl %r14d, %r12d
+ jl .L15
+ movl $32785, -52(%rbp)
+ addq $40, %rsp
+ popq %rbx
+ popq %r12
+ popq %r13
+ popq %r14
+ popq %r15
+ popq %rbp
+ .cfi_def_cfa 7, 8
+ ret
+ .cfi_endproc
+.LFE12:
+ .size main._omp_fn.0, .-main._omp_fn.0
+ .data
+ .align 16
+ .type _ZZ4mainE17.omp_data_sizes.4, @object
+ .size _ZZ4mainE17.omp_data_sizes.4, 16
+_ZZ4mainE17.omp_data_sizes.4:
+ .quad 4
+ .quad 131140
+ .align 2
+ .type _ZZ4mainE17.omp_data_kinds.5, @object
+ .size _ZZ4mainE17.omp_data_kinds.5, 4
+_ZZ4mainE17.omp_data_kinds.5:
+ .value 643
+ .value 643
+ .section .gnu.offload_vars,"aw",@progbits
+ .align 8
+ .type .offload_var_table, @object
+ .size .offload_var_table, 0
+.offload_var_table:
+ .section .gnu.offload_funcs,"aw",@progbits
+ .align 8
+ .type .offload_func_table, @object
+ .size .offload_func_table, 8
+.offload_func_table:
+ .quad main._omp_fn.0
+ .comm __gnu_lto_v1,1,1
+ .ident "GCC: (Sourcery CodeBench (OpenACC/PTX) Lite 2016.05-999999) 6.0.0 20151019 (experimental)"
+ .section .note.GNU-stack,"",@progbits
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c (revision 0)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c (working copy)
@@ -0,0 +1,57 @@
+/* { dg-do run } */
+/* { dg-additional-options "-O2" */
+
+#include <stdio.h>
+
+#define N (32*32*32+17)
+int main ()
+{
+ int ary[N];
+ int ix;
+ int exit = 0;
+ int ondev = 0;
+
+ for (ix = 0; ix < N;ix++)
+ ary[ix] = -1;
+
+#pragma acc parallel num_gangs(32) vector_length(32) copy(ary) copy(ondev)
+ {
+#pragma acc loop gang (static:1)
+ for (unsigned ix = 0; ix < N; ix++)
+ {
+ if (__builtin_acc_on_device (5))
+ {
+ int g = 0, w = 0, v = 0;
+
+ __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g));
+ __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w));
+ __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v));
+ ary[ix] = (g << 16) | (w << 8) | v;
+ ondev = 1;
+ }
+ else
+ ary[ix] = ix;
+ }
+ }
+
+ for (ix = 0; ix < N; ix++)
+ {
+ int expected = ix;
+ if(ondev)
+ {
+ int g = ix % 32;
+ int w = 0;
+ int v = 0;
+
+ expected = (g << 16) | (w << 8) | v;
+ }
+
+ if (ary[ix] != expected)
+ {
+ exit = 1;
+ printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected);
+ }
+ }
+
+ return exit;
+}
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c (revision 0)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c (working copy)
@@ -0,0 +1,59 @@
+/* { dg-do run } */
+/* { dg-additional-options "-O2" */
+
+#include <stdio.h>
+
+#define N (32*32*32+17)
+int main ()
+{
+ int ary[N];
+ int ix;
+ int exit = 0;
+ int ondev = 0;
+
+ for (ix = 0; ix < N;ix++)
+ ary[ix] = -1;
+
+#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) copy(ary) copy(ondev)
+ {
+#pragma acc loop gang worker vector
+ for (unsigned ix = 0; ix < N; ix++)
+ {
+ if (__builtin_acc_on_device (5))
+ {
+ int g = 0, w = 0, v = 0;
+
+ __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g));
+ __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w));
+ __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v));
+ ary[ix] = (g << 16) | (w << 8) | v;
+ ondev = 1;
+ }
+ else
+ ary[ix] = ix;
+ }
+ }
+
+ for (ix = 0; ix < N; ix++)
+ {
+ int expected = ix;
+ if(ondev)
+ {
+ int chunk_size = (N + 32*32*32 - 1) / (32*32*32);
+
+ int g = ix / (chunk_size * 32 * 32);
+ int w = ix / 32 % 32;
+ int v = ix % 32;
+
+ expected = (g << 16) | (w << 8) | v;
+ }
+
+ if (ary[ix] != expected)
+ {
+ exit = 1;
+ printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected);
+ }
+ }
+
+ return exit;
+}
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c (revision 0)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c (working copy)
@@ -0,0 +1,57 @@
+/* { dg-do run } */
+/* { dg-additional-options "-O2" */
+
+#include <stdio.h>
+
+#define N (32*32*32+17)
+int main ()
+{
+ int ary[N];
+ int ix;
+ int exit = 0;
+ int ondev = 0;
+
+ for (ix = 0; ix < N;ix++)
+ ary[ix] = -1;
+
+#pragma acc parallel vector_length(32) copy(ary) copy(ondev)
+ {
+#pragma acc loop vector
+ for (unsigned ix = 0; ix < N; ix++)
+ {
+ if (__builtin_acc_on_device (5))
+ {
+ int g = 0, w = 0, v = 0;
+
+ __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g));
+ __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w));
+ __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v));
+ ary[ix] = (g << 16) | (w << 8) | v;
+ ondev = 1;
+ }
+ else
+ ary[ix] = ix;
+ }
+ }
+
+ for (ix = 0; ix < N; ix++)
+ {
+ int expected = ix;
+ if(ondev)
+ {
+ int g = 0;
+ int w = 0;
+ int v = ix % 32;
+
+ expected = (g << 16) | (w << 8) | v;
+ }
+
+ if (ary[ix] != expected)
+ {
+ exit = 1;
+ printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected);
+ }
+ }
+
+ return exit;
+}
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c (revision 0)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c (working copy)
@@ -0,0 +1,57 @@
+/* { dg-do run } */
+/* { dg-additional-options "-O2" */
+
+#include <stdio.h>
+
+#define N (32*32*32+17)
+int main ()
+{
+ int ary[N];
+ int ix;
+ int exit = 0;
+ int ondev = 0;
+
+ for (ix = 0; ix < N;ix++)
+ ary[ix] = -1;
+
+#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev)
+ {
+#pragma acc loop worker
+ for (unsigned ix = 0; ix < N; ix++)
+ {
+ if (__builtin_acc_on_device (5))
+ {
+ int g = 0, w = 0, v = 0;
+
+ __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g));
+ __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w));
+ __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v));
+ ary[ix] = (g << 16) | (w << 8) | v;
+ ondev = 1;
+ }
+ else
+ ary[ix] = ix;
+ }
+ }
+
+ for (ix = 0; ix < N; ix++)
+ {
+ int expected = ix;
+ if(ondev)
+ {
+ int g = 0;
+ int w = ix % 32;
+ int v = 0;
+
+ expected = (g << 16) | (w << 8) | v;
+ }
+
+ if (ary[ix] != expected)
+ {
+ exit = 1;
+ printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected);
+ }
+ }
+
+ return exit;
+}
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c (revision 0)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c (working copy)
@@ -0,0 +1,57 @@
+/* { dg-do run } */
+/* { dg-additional-options "-O2" */
+
+#include <stdio.h>
+
+#define N (32*32*32+17)
+int main ()
+{
+ int ary[N];
+ int ix;
+ int exit = 0;
+ int ondev = 0;
+
+ for (ix = 0; ix < N;ix++)
+ ary[ix] = -1;
+
+#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev)
+ {
+#pragma acc loop worker vector
+ for (unsigned ix = 0; ix < N; ix++)
+ {
+ if (__builtin_acc_on_device (5))
+ {
+ int g = 0, w = 0, v = 0;
+
+ __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g));
+ __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w));
+ __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v));
+ ary[ix] = (g << 16) | (w << 8) | v;
+ ondev = 1;
+ }
+ else
+ ary[ix] = ix;
+ }
+ }
+
+ for (ix = 0; ix < N; ix++)
+ {
+ int expected = ix;
+ if(ondev)
+ {
+ int g = 0;
+ int w = (ix / 32) % 32;
+ int v = ix % 32;
+
+ expected = (g << 16) | (w << 8) | v;
+ }
+
+ if (ary[ix] != expected)
+ {
+ exit = 1;
+ printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected);
+ }
+ }
+
+ return exit;
+}
