From patchwork Mon May 16 17:45:19 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Martin Jambor X-Patchwork-Id: 622690 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 3r7ntd51mYz9t6x for ; Tue, 17 May 2016 03:45:43 +1000 (AEST) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=FzGSXFuo; dkim-atps=neutral DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:date :from:to:cc:subject:message-id:mime-version:content-type; q=dns; s=default; b=m7WEmmuNGoFvRmni5+8YpKlPUmDt48bUDfkYT42jHvZXwa67Lj OnUcchp7EPCVa7Zr3pvKVhJDJckQctd1OUW6IeRpMgZduPneawhOJRYMN69n+Xy5 OiaQnw4y+gSIoZSOpxvk/lkREZ4xH9VVi5v1+xRDiO0R+EQup0Lwu3rJU= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:date :from:to:cc:subject:message-id:mime-version:content-type; s= default; bh=edQN6NHrJwaeeSlZQNI3mlpzcxc=; b=FzGSXFuobYOcey1CsudU TXyssmBAZ/djbXkS+DP7RQFPp++gpaRGp9/XtLzlDawMShBHvhaBnl5+s6os3Idp EPs94HmE2hp5bvpr1b2gM9x3ISsaCfLXTJdlGYQyDx8ptQIBrGv4Fk+zGn/e7skE Hv5LaMNEc7DH6qa5+fP7lcw= Received: (qmail 21358 invoked by alias); 16 May 2016 17:45:33 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 21326 invoked by uid 89); 16 May 2016 17:45:32 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=BAYES_00, SPF_PASS autolearn=ham version=3.3.2 spammy=finalize X-HELO: mx2.suse.de Received: from mx2.suse.de (HELO mx2.suse.de) (195.135.220.15) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (CAMELLIA256-SHA encrypted) ESMTPS; Mon, 16 May 2016 17:45:22 +0000 Received: from relay1.suse.de (charybdis-ext.suse.de [195.135.220.254]) by mx2.suse.de (Postfix) with ESMTP id 44555AB1A for ; Mon, 16 May 2016 17:45:19 +0000 (UTC) Date: Mon, 16 May 2016 19:45:19 +0200 From: Martin Jambor To: GCC Patches Cc: Martin Liska Subject: [hsa] Increase hsa symbol alignment to a natural one Message-ID: <20160516174518.GN5580@virgil.suse.cz> Mail-Followup-To: GCC Patches , Martin Liska MIME-Version: 1.0 Content-Disposition: inline User-Agent: Mutt/1.6.0 (2016-04-01) X-IsSubscribed: yes Hi, in the last round fo alignment fixes, we have forgot to make sure that all symbols are at least naturally aligned, which is a hard HSAIL requirement. This caused problems when emitting a symbol for a private complex number, as the natural alignment as defined by HSAIL is twice the one of the component, which was selected by gcc. The following patch addresses this in two ways. First, it simply increases the alignment of symbols that are only accessible from within HSAIL. If however a symbol that is shared in between host and an HSA accelerator is under-aligned (in my experience it only happens if the user uses the aligned attribute), we have no option but to abort HSAIL generation because even if we did generate it, it would not finalize. Bootstrapped and tested on x86_64-linux with hsa enabled. I will commit it to trunk and the gcc-6 branch shortly. Martin 2016-05-16 Martin Jambor * hsa-gen.c (fillup_for_decl): Increase alignment to natural one. (get_symbol_for_decl): Sorry if a global symbol in under-aligned. libgomp/ * testsuite/libgomp.hsa.c/complex-align-2.c: New test. --- gcc/hsa-gen.c | 19 ++++++++++++---- libgomp/testsuite/libgomp.hsa.c/complex-align-2.c | 27 +++++++++++++++++++++++ 2 files changed, 42 insertions(+), 4 deletions(-) create mode 100644 libgomp/testsuite/libgomp.hsa.c/complex-align-2.c diff --git a/gcc/hsa-gen.c b/gcc/hsa-gen.c index 5baf607..697d599 100644 --- a/gcc/hsa-gen.c +++ b/gcc/hsa-gen.c @@ -203,9 +203,13 @@ hsa_symbol::fillup_for_decl (tree decl) { m_decl = decl; m_type = hsa_type_for_tree_type (TREE_TYPE (decl), &m_dim, false); - if (hsa_seen_error ()) - m_seen_error = true; + { + m_seen_error = true; + return; + } + + m_align = MAX (m_align, hsa_natural_alignment (m_type)); } /* Constructor of class representing global HSA function/kernel information and @@ -929,6 +933,14 @@ get_symbol_for_decl (tree decl) BRIG_LINKAGE_PROGRAM, true, BRIG_ALLOCATION_PROGRAM, align); hsa_cfun->m_global_symbols.safe_push (sym); + sym->fillup_for_decl (decl); + if (sym->m_align > align) + { + sym->m_seen_error = true; + HSA_SORRY_ATV (EXPR_LOCATION (decl), + "HSA specification requires that %E is at least " + "naturally aligned", decl); + } } else { @@ -944,12 +956,11 @@ get_symbol_for_decl (tree decl) sym = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_PRIVATE, BRIG_LINKAGE_FUNCTION); sym->m_align = align; + sym->fillup_for_decl (decl); hsa_cfun->m_private_variables.safe_push (sym); } - sym->fillup_for_decl (decl); sym->m_name = hsa_get_declaration_name (decl); - *slot = sym; return sym; } diff --git a/libgomp/testsuite/libgomp.hsa.c/complex-align-2.c b/libgomp/testsuite/libgomp.hsa.c/complex-align-2.c new file mode 100644 index 0000000..b2d7acf --- /dev/null +++ b/libgomp/testsuite/libgomp.hsa.c/complex-align-2.c @@ -0,0 +1,27 @@ +#pragma omp declare target + _Complex int *g; +#pragma omp end declare target + + + +_Complex float f(void); + +int +main () +{ + _Complex int y; +#pragma omp target map(from:y) + { + _Complex int x; + g = &x; + __imag__ x = 1; + __real__ x = 2; + y = x; + } + + if ((__imag__ y != 1) + || (__real__ y != 2)) + __builtin_abort (); + return 0; +} +