From mboxrd@z Thu Jan  1 00:00:00 1970
Return-Path: <gcc-patches-return-412870-listarch-gcc-patches=gcc.gnu.org@gcc.gnu.org>
Received: (qmail 62839 invoked by alias); 6 Nov 2015 10:54:37 -0000
Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm
Precedence: bulk
List-Id: <gcc-patches.gcc.gnu.org>
List-Archive: <http://gcc.gnu.org/ml/gcc-patches/>
List-Post: <mailto:gcc-patches@gcc.gnu.org>
List-Help: <mailto:gcc-patches-help@gcc.gnu.org>
Sender: gcc-patches-owner@gcc.gnu.org
Received: (qmail 62824 invoked by uid 89); 6 Nov 2015 10:54:36 -0000
Authentication-Results: sourceware.org; auth=none
X-Virus-Found: No
X-Spam-SWARE-Status: No, score=-1.0 required=5.0 tests=AWL,BAYES_00,SPF_PASS autolearn=ham version=3.3.2
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; Fri, 06 Nov 2015 10:54:35 +0000
Received: from relay1.suse.de (charybdis-ext.suse.de [195.135.220.254])	by mx2.suse.de (Postfix) with ESMTP id B4D23AAD0;	Fri,  6 Nov 2015 10:54:11 +0000 (UTC)
Subject: Re: Merge of HSA branch
To: gcc-patches@gcc.gnu.org
References: <20151105215108.GC9264@virgil.suse.cz> <563C7D2B.6070806@redhat.com>
From: =?UTF-8?Q?Martin_Li=c5=a1ka?= <mliska@suse.cz>
Cc: bschmidt@redhat.com
Message-ID: <563C86E7.9000100@suse.cz>
Date: Fri, 06 Nov 2015 10:54:00 -0000
User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.3.0
MIME-Version: 1.0
In-Reply-To: <563C7D2B.6070806@redhat.com>
Content-Type: text/plain; charset=utf-8
Content-Transfer-Encoding: 7bit
X-IsSubscribed: yes
X-SW-Source: 2015-11/txt/msg00586.txt.bz2

On 11/06/2015 11:12 AM, Bernd Schmidt wrote:
> On 11/05/2015 10:51 PM, Martin Jambor wrote:
>> Individual changes are described in slightly more detail in their
>> respective messages.  If you are interested in how the HSAIL
>> generation works in general, I encourage you to have a look at my
>> Cauldron slides or presentation, only very few things have changed as
>> far as the general principles are concerned.  Let me just quickly stress
>> here that we do acceleration within a single compiler, as opposed to
>> LTO-ways of all the other accelerator teams.
> 
> Realistically we're probably not going to reject this work, but I still want to ask whether the approach was acked by the community before you started. I'm really not exactly thrilled about having two different classes of backends in the compiler, and two different ways of handling offloading.
> 
>> I also acknowledge that we should add HSA-specific tests to the GCC
>> testsuite but we are only now looking at how to do that and will
>> welcome any guidance in this regard.
> 
> Yeah, I was looking for any kind of new test, because...
> 
>> the class of OpenMP loops we can handle well is small,
> 
> I'd appreciate more information on what this means. Any examples or performance numbers?

Hello.

As mentioned by Martin Jambor, it was explained during his speech at the Cauldron this year.
It can be easily explained on the following simple case:

#pragma omp target teams
#pragma omp distribute parallel for private(j)
   for (j=0; j<N; j++)
      c[j] = a[j];

Which is simple vector copy, that's going to be transformed to:

_4 = omp_data.i_1(D).D.5301 (iteration space)
_5 = __builtin_omp_get_num_threads ();
_6 = __builtin_omp_get_thread_num ();
_7 = calculate_chunk_start (_4, _5, _6); // pseudocode
_8 = calculate_chunk_end (_4, _5, _6); // pseudocode

for(i = _7; i < _8; i++)
  dest[i] = src[i];

and such kernel is dispatched with default grid size (in our case 64), so that every
work item handles chunk of size N/64.

On the other hand, gridification is doing to transform to:

_7 = __builtin_omp_get_thread_num ();
dest[_7] = src[_7];

and the kernels is offloaded like this:
HSA debug: GOMP_OFFLOAD_run called with grid size 10000000 and group size 0

Performance numbers are in order of magnitude and can be seen on slides 27-30 in [1]

Martin

[1] https://gcc.gnu.org/wiki/cauldron2015?action=AttachFile&do=get&target=mjambor-hsa-slides.pdf

> 
> 
> Bernd