From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 6671 invoked by alias); 27 Jan 2014 13:13:09 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 6661 invoked by uid 89); 27 Jan 2014 13:13:08 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=AWL,BAYES_00 autolearn=ham version=3.3.2 X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Mon, 27 Jan 2014 13:13:06 +0000 Received: from svr-orw-fem-01.mgc.mentorg.com ([147.34.98.93]) by relay1.mentorg.com with esmtp id 1W7lzq-0001lB-74 from Bernd_Schmidt@mentor.com ; Mon, 27 Jan 2014 05:13:02 -0800 Received: from SVR-IES-FEM-02.mgc.mentorg.com ([137.202.0.106]) by svr-orw-fem-01.mgc.mentorg.com over TLS secured channel with Microsoft SMTPSVC(6.0.3790.4675); Mon, 27 Jan 2014 05:13:01 -0800 Received: from [127.0.0.1] (137.202.0.76) by SVR-IES-FEM-02.mgc.mentorg.com (137.202.0.106) with Microsoft SMTP Server id 14.2.247.3; Mon, 27 Jan 2014 13:12:59 +0000 Message-ID: <52E65B56.7010107@codesourcery.com> Date: Mon, 27 Jan 2014 13:13:00 -0000 From: Bernd Schmidt User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:24.0) Gecko/20100101 Thunderbird/24.2.0 MIME-Version: 1.0 To: Andrey Turetskiy CC: "Michael V. Zolotukhin" , Jakub Jelinek , Thomas Schwinge , Richard Biener , Kirill Yukhin , Ilya Verbin , Ilya Tocar , gcc Subject: Re: [RFC][gomp4] Offloading patches (1/3): Add '-fopenmp_target' option References: <20131217113522.GA23078@msticlxl57.ims.intel.com> <52DE86DC.10103@codesourcery.com> In-Reply-To: Content-Type: text/plain; charset="windows-1252"; format=flowed Content-Transfer-Encoding: 8bit X-SW-Source: 2014-01/txt/msg01687.txt.bz2 On 01/22/2014 11:53 AM, Andrey Turetskiy wrote: > We have some testcases, but they require XeonPhi hardware and a > working libgomp plugin. Our current version of the plugin depends on > some libraries, that are not open-sourced yet, so currently we can’t > share it. > > However, you could examine what these patches do, making the following steps: > 1) Build GCC with patches: > http://gcc.gnu.org/ml/gcc-patches/2013-12/msg01484.html > http://gcc.gnu.org/ml/gcc-patches/2013-12/msg01485.html > http://gcc.gnu.org/ml/gcc-patches/2013-12/msg01486.html > http://gcc.gnu.org/ml/gcc-patches/2013-12/msg01896.html > 2) Set environment variables (e.g. for two ‘targets’): > export OFFLOAD_TARGET_NAMES=mic:hsail (for now > names don’t really matter) > export OFFLOAD_TARGET_COMPILERS=./gcc:./gcc (use GCC with > patches above as target compiler, because it must support the > -fopenmp_target option) > 3) Build any example with #pragma omp target (e.g. see attachment): > ./gcc -flto -fopenmp test.c -o test.exe > Options -flto and -fopenmp are necessary for using. > > Now you have a binary with target images embedded and tables properly > filled. You can’t run it due to reasons mentioned above, though you > could examine it with objdump/nm/readelf to see new sections and their > content: there will be .offload_image_section with ‘target’ code and > .offload_func_table_section with ‘target’ function table. I played around with this for a while last week. To have a slightly more realistic scenario where the offload compiler is for a different target, I built an aarch64-linux compiler and used that in OFFLOAD_TARGET_COMPILERS. This exposed some problems. > + /* Run gcc for target. */ > + obstack_init (&argv_obstack); > + obstack_ptr_grow (&argv_obstack, compiler); > + obstack_ptr_grow (&argv_obstack, "-shared"); > + obstack_ptr_grow (&argv_obstack, "-fPIC"); > + obstack_ptr_grow (&argv_obstack, "-xlto"); > + obstack_ptr_grow (&argv_obstack, "-fopenmp_target"); > + obstack_ptr_grow (&argv_obstack, "-o"); > + obstack_ptr_grow (&argv_obstack, target_image_file_name); Since environment variables such as GCC_EXEC_PREFIX and COMPILER_PATH are set at this point, the compiler we're running here won't find the correct lto1 - best case it doesn't find anything, worst case it finds the lto1 for the host compiler and produces an image for the host, not the target (this fails with an arm compiler since the host assembler doesn't understand -meabi=5, but it could silently do the wrong thing with other offload toolchains). Once I worked around this by unsetting the environment variables around this compiler invocation here, the next problem is exposed - the code tries to link together files compiled for the target (created by the code quoted above) and the host (the _omp_descr file, I believe). Linker errors ensue. As mentioned before, I think all this target-specific code has no place in lto-wrapper to begin with. For ptx, we're going to require some quite different mechanisms, so I think it might be best to invoke a new tool, maybe called $target-gen-offload, which knows how to produce an image that can be linked into the host executable. Different offload targets can then use different strategies to produce such an image. Probably each such image should contain its own code to register itself with libgomp, so that we don't have to construct a table. Some other observations: * is OFFLOAD_TARGET_NAMES actually useful, or would any string generated at link time suffice? * Is the user expected to set OFFLOAD_TARGET_COMPILERS, or should this be done by the gcc driver, possibly based on command line options (I'd much prefer that)? * Do we actually need an -fopenmp-target option? The way I imagine it (and which was somewhat present in the Makefile patches I posted last year) is that an offload compiler is specially configured to know that that's how it will be used, and to know what the host architecture is. A $target-gen-offload could then be built with knowledge of the host architecture and installed in the host compiler's libexec install directory. I think I'll need to implement my own set of mechanisms for ptx, since this code doesn't seem suitable for inclusion in its current state. I'll try to take on board some of the ideas I've found here in the hope that we'll converge on something that works for everybody. Bernd