public inbox for gcc-bugs@sourceware.org
help / color / mirror / Atom feed
From: "burnus at gcc dot gnu.org" <gcc-bugzilla@gcc.gnu.org>
To: gcc-bugs@gcc.gnu.org
Subject: [Bug c/64748] New: OpenACC: "is not a variable" error with deviceptr()
Date: Fri, 23 Jan 2015 13:36:00 -0000 [thread overview]
Message-ID: <bug-64748-4@http.gcc.gnu.org/bugzilla/> (raw)
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #1: Type: text/plain; charset="UTF-8", Size: 5157 bytes --]
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=64748
Bug ID: 64748
Summary: OpenACC: "is not a variable" error with deviceptr()
Product: gcc
Version: 5.0
Status: UNCONFIRMED
Keywords: openacc, rejects-valid
Severity: normal
Priority: P3
Component: c
Assignee: unassigned at gcc dot gnu.org
Reporter: burnus at gcc dot gnu.org
From https://github.com/jefflarkin/openacc-interoperability/
The following program seems to compile with PGI's and Cray's compilers. With
GCC, it shows the odd:
foo.c:3:30: error: âarrâ is not a variable
#pragma acc kernels deviceptr(arr)
^
void set(int n, float val, float * restrict arr)
{
#pragma acc kernels deviceptr(arr)
{
for(int i=0; i<n; i++)
{
arr[i] = val;
}
}
}
>From gcc-bugs-return-474563-listarch-gcc-bugs=gcc.gnu.org@gcc.gnu.org Fri Jan 23 13:40:13 2015
Return-Path: <gcc-bugs-return-474563-listarch-gcc-bugs=gcc.gnu.org@gcc.gnu.org>
Delivered-To: listarch-gcc-bugs@gcc.gnu.org
Received: (qmail 17562 invoked by alias); 23 Jan 2015 13:40:08 -0000
Mailing-List: contact gcc-bugs-help@gcc.gnu.org; run by ezmlm
Precedence: bulk
List-Id: <gcc-bugs.gcc.gnu.org>
List-Archive: <http://gcc.gnu.org/ml/gcc-bugs/>
List-Post: <mailto:gcc-bugs@gcc.gnu.org>
List-Help: <mailto:gcc-bugs-help@gcc.gnu.org>
Sender: gcc-bugs-owner@gcc.gnu.org
Delivered-To: mailing list gcc-bugs@gcc.gnu.org
Received: (qmail 17411 invoked by uid 48); 23 Jan 2015 13:39:47 -0000
From: "adrien at guinet dot me" <gcc-bugzilla@gcc.gnu.org>
To: gcc-bugs@gcc.gnu.org
Subject: [Bug target/64749] New: "truncating" instructions generated instead of a load one using SSE & AVX2 intrinsics
Date: Fri, 23 Jan 2015 13:40:00 -0000
X-Bugzilla-Reason: CC
X-Bugzilla-Type: new
X-Bugzilla-Watch-Reason: None
X-Bugzilla-Product: gcc
X-Bugzilla-Component: target
X-Bugzilla-Version: 4.8.4
X-Bugzilla-Keywords:
X-Bugzilla-Severity: normal
X-Bugzilla-Who: adrien at guinet dot me
X-Bugzilla-Status: UNCONFIRMED
X-Bugzilla-Priority: P3
X-Bugzilla-Assigned-To: unassigned at gcc dot gnu.org
X-Bugzilla-Target-Milestone: ---
X-Bugzilla-Flags:
X-Bugzilla-Changed-Fields: bug_id short_desc product version bug_status bug_severity priority component assigned_to reporter attachments.created
Message-ID: <bug-64749-4@http.gcc.gnu.org/bugzilla/>
Content-Type: text/plain; charset="UTF-8"
Content-Transfer-Encoding: 7bit
X-Bugzilla-URL: http://gcc.gnu.org/bugzilla/
Auto-Submitted: auto-generated
MIME-Version: 1.0
X-SW-Source: 2015-01/txt/msg02557.txt.bz2
Content-length: 2410
https://gcc.gnu.org/bugzilla/show_bug.cgi?idd749
Bug ID: 64749
Summary: "truncating" instructions generated instead of a load
one using SSE & AVX2 intrinsics
Product: gcc
Version: 4.8.4
Status: UNCONFIRMED
Severity: normal
Priority: P3
Component: target
Assignee: unassigned at gcc dot gnu.org
Reporter: adrien at guinet dot me
Created attachment 34553
--> https://gcc.gnu.org/bugzilla/attachment.cgi?id4553&actioníit
test case
The code attached compiles and runs fine (that is the output of the program is
the good one) using GCC 4.9. When compiled with GCC 4.8, the output is
different and incorrect.
Indeed, when compiled with GCC 4.8, some kind of truncating is introduced at
the begginig of the loop (in f2). Here is the relevant assembly code (output of
GCC 4.8) :
xor eax, eax
mov rbp, rsp
and rsp, 0FFFFFFFFFFFFFFE0h
vbroadcastss ymm3, xmm6
add rsp, 10h
nop dword ptr [rax]
loc_400970:
vpmovzxwd ymm4, xmmword ptr [rdx+rax*4]
vpmovzxwd ymm2, xmmword ptr [rcx+rax*4]
vmovdqa [rsp-8+var_28], ymm4
; truncation here is done
vmovdqa xmm5, xmmword ptr [rsp-8+var_28]
vpmulld ymm0, ymm4, ymm2
; here it uses xmm5 which isn't thus the good value.
; xmm5 and ymm4 should be set like with something like this (like GCC 4.9
does):
; vmovqda xmm5, xmmword ptr [rdx+rax*4]
; vpmovzxwd ymm4, xmm5
vpmulhuw xmm1, xmm5, xmmword ptr [r8+rax*4]
vpmovzxwd ymm1, xmm1
vpmulld ymm1, ymm1, ymm3
vpsubd ymm0, ymm0, ymm1
vmovdqa xmmword ptr [rsi+rax*4], xmm0
add rax, 8
cmp rdi, rax
ja short loc_400970
GCC 4.9 indeed behaves correctly and generate this assembly code :
vbroadcastss ymm3, dword ptr [rbp-14h]
xor eax, eax
nop dword ptr [rax+00h]
loc_4009A8:
vmovdqa xmm0, xmmword ptr [rdx+rax*4] ; 128-bits load
vpmulhuw xmm2, xmm0, xmmword ptr [r8+rax*4] ; correctly uses xmm0
vpmovzxwd ymm2, xmm2 ; 16->32 bits conversion here
vpmulld ymm2, ymm2, ymm3
vpmovzxwd ymm1, xmm0
vpmovzxwd ymm0, xmmword ptr [rcx+rax*4]
vpmulld ymm0, ymm1, ymm0
vpsubd ymm0, ymm0, ymm2
vmovaps xmmword ptr [rsi+rax*4], xmm0
add rax, 8
cmp rdi, rax
ja short loc_4009A8
Thanks for any help about this!
P.S: sorry but I didn't manage to have a shorter test case :/
next reply other threads:[~2015-01-23 13:36 UTC|newest]
Thread overview: 5+ messages / expand[flat|nested] mbox.gz Atom feed top
2015-01-23 13:36 burnus at gcc dot gnu.org [this message]
2015-02-10 14:39 ` [Bug c/64748] " jakub at gcc dot gnu.org
2015-02-10 15:17 ` tschwinge at gcc dot gnu.org
2015-02-10 15:22 ` jakub at gcc dot gnu.org
2015-02-19 15:19 ` tschwinge at gcc dot gnu.org
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=bug-64748-4@http.gcc.gnu.org/bugzilla/ \
--to=gcc-bugzilla@gcc.gnu.org \
--cc=gcc-bugs@gcc.gnu.org \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).