similar to: [LLVMdev] PROPOSAL: Rename Target R600 -> AMDGPU

Hi, I'm finally going to do the R600->AMDGPU rename this Friday. This is something I originally proposed last July [1], but had to put off in order to avoid creating really bad merge headaches for some users. The only change from my original proposal is that I'll just keep the existing r600 and amdgcn triples rather than adding a new one for amdgpu. If anyone has any strong

Moving this thread to llvmdev. On Tue, Oct 30, 2012 at 11:09:34PM -0700, Chris Lattner wrote: > On Oct 30, 2012, at 11:35 AM, Tom Stellard <tom at stellard.net> wrote: > >> Hi Tom, > >> > >> Time is running short, but this would be great. The best place to start is to begin decomposing the mega-patch into individual pieces that makes sense. Do you have

On 01.11.2012, at 14:44, Tom Stellard <tom at stellard.net> wrote: > Moving this thread to llvmdev. > > On Tue, Oct 30, 2012 at 11:09:34PM -0700, Chris Lattner wrote: >> On Oct 30, 2012, at 11:35 AM, Tom Stellard <tom at stellard.net> wrote: >>>> Hi Tom, >>>> >>>> Time is running short, but this would be great. The best place to

On Sat, Nov 17, 2012 at 10:56:26PM +0100, Benjamin Kramer wrote: > > On 01.11.2012, at 14:44, Tom Stellard <tom at stellard.net> wrote: > > > Moving this thread to llvmdev. > > > > On Tue, Oct 30, 2012 at 11:09:34PM -0700, Chris Lattner wrote: > >> On Oct 30, 2012, at 11:35 AM, Tom Stellard <tom at stellard.net> wrote: > >>>> Hi

Hi Mr. Liu, Thanks for your quick reply. I compiled the code with the libclc_trunk and linked the bitcode file under $LIBCLC_DIR/built_libs/tahiti-amdgcn--.bc. After looking into the libclc, it is currently using the new workitem intrinsics (commit ba9858caa1e927a6fcc601e3466faa693835db5e). In the linked bitcode ($LIBCLC_DIR/built_libs/tahiti-amdgcn--.bc), it has the following code segment,

Is there a version of the AMDIL back-end that is compatible with LLVM 3.0/3.1? On Tue, May 29, 2012 at 8:33 AM, Villmow, Micah <Micah.Villmow at amd.com>wrote: > > > > -----Original Message----- > > From: Stellard, Thomas > > Sent: Monday, May 28, 2012 9:07 AM > > To: Justin Holewinski > > Cc: Villmow, Micah; Tom Stellard; llvmdev at cs.uiuc.edu >

> -----Original Message----- > From: Stellard, Thomas > Sent: Monday, May 28, 2012 9:07 AM > To: Justin Holewinski > Cc: Villmow, Micah; Tom Stellard; llvmdev at cs.uiuc.edu > Subject: Re: [LLVMdev] RFC: R600, a new backend for AMD GPUs > > On Mon, May 28, 2012 at 08:54:41AM -0700, Justin Holewinski wrote: > > On May 28, 2012 6:44 AM, "Tom Stellard"

Dear Developers, I compiled a OpenCL kernel before (on Nov. last year) like __kernel void g(__global float* array) { array[get_global_id(0)] = 1; } with libclc, which would originally use the instrinsics like llvm.r600.read.local.size.x(). I executed the generated object file with one version of the hsa-runtime [1] provided by Mr. Stellard, when there was more than one workgroup, the output

On Mon, May 28, 2012 at 08:54:41AM -0700, Justin Holewinski wrote: > On May 28, 2012 6:44 AM, "Tom Stellard" <thomas.stellard at amd.com> wrote: > > > > On Fri, May 25, 2012 at 02:37:26PM -0700, Justin Holewinski wrote: > > > Hi Tom, > > > > > > I have a higher-level question regarding this back-end. If I have an > LLVM > > >

On Mon, Apr 09, 2012 at 09:37:37AM -0400, Tom Stellard wrote: > On Mon, Mar 26, 2012 at 12:50:07PM -0400, Tom Stellard wrote: > > Hi, > > > > We've been working on an LLVM backend for the previous generation of AMD > > GPUs (HD 2XXX - HD 6XXX) and we would like submit it for inclusion in the > > main LLVM tree. The latest code can be found in this git

On May 28, 2012 6:44 AM, "Tom Stellard" <thomas.stellard at amd.com> wrote: > > On Fri, May 25, 2012 at 02:37:26PM -0700, Justin Holewinski wrote: > > Hi Tom, > > > > I have a higher-level question regarding this back-end. If I have an LLVM > > IR module and run it through this back-end, it seems like the only output > > option is a binary

On Mon, Mar 26, 2012 at 12:50:07PM -0400, Tom Stellard wrote: > Hi, > > We've been working on an LLVM backend for the previous generation of AMD > GPUs (HD 2XXX - HD 6XXX) and we would like submit it for inclusion in the > main LLVM tree. The latest code can be found in this git repository: > http://cgit.freedesktop.org/~tstellar/llvm/ in the r600-initial-review > branch

Tom, Two things. One is missing tests. I have some I could send you, but they are mainly OpenCL based for the AMDIL backend, not for the R600. That brings me to the second thing. Are the AMDIL backend and the R600 backend the same, or not? At this point, they really do feel like they are separate back ends, with one dependent on the other. As there is no other backend that is dependent on

On 26.11.2012, at 23:37, Tom Stellard <tom at stellard.net> wrote: > On Sat, Nov 17, 2012 at 10:56:26PM +0100, Benjamin Kramer wrote: >> >> On 01.11.2012, at 14:44, Tom Stellard <tom at stellard.net> wrote: >> >>> Moving this thread to llvmdev. >>> >>> On Tue, Oct 30, 2012 at 11:09:34PM -0700, Chris Lattner wrote: >>>> On

On Fri, May 25, 2012 at 02:37:26PM -0700, Justin Holewinski wrote: > Hi Tom, > > I have a higher-level question regarding this back-end. If I have an LLVM > IR module and run it through this back-end, it seems like the only output > option is a binary format. Is this a device binary, or another > intermediate format? > > If the input LLVM IR module was a compute kernel,

On Wed, Oct 09, 2013 at 08:06:42PM -0500, Jay Cornwall wrote: > Hi, > > This patch adds R600/SI disassembly text to compiled object files, when > a code dump is requested, to assist debugging in Mesa clients. > > Here's an example of the output in a Mesa client with a corresponding > patch and RADEON_DUMP_SHADERS set: > > Shader Disassembly: > >

Hi Tom, I have a higher-level question regarding this back-end. If I have an LLVM IR module and run it through this back-end, it seems like the only output option is a binary format. Is this a device binary, or another intermediate format? If the input LLVM IR module was a compute kernel, how would I go about executing it on an AMD GPU? Can I use the APP SDK to load the binary, perhaps

----- Original Message ----- > From: "Tom Stellard" <tom at stellard.net> > To: "Hal Finkel" <hfinkel at anl.gov> > Cc: "LLVM Developers Mailing List" <llvmdev at cs.uiuc.edu> > Sent: Thursday, March 13, 2014 9:46:22 AM > Subject: Re: [LLVMdev] Be Careful with Positionally-Encoded Operands (AArch64, Mips, AMDGPU, etc.) > > On

On Mon, Mar 26, 2012 at 12:50:07PM -0400, Tom Stellard wrote: > Hi, > > We've been working on an LLVM backend for the previous generation of AMD > GPUs (HD 2XXX - HD 6XXX) and we would like submit it for inclusion in the > main LLVM tree. The latest code can be found in this git repository: > http://cgit.freedesktop.org/~tstellar/llvm/ in the r600-initial-review > branch

Hi, This patch adds R600/SI disassembly text to compiled object files, when a code dump is requested, to assist debugging in Mesa clients. Here's an example of the output in a Mesa client with a corresponding patch and RADEON_DUMP_SHADERS set: Shader Disassembly: S_WQM_B64 EXEC, EXEC ; BEFE0A7E S_MOV_B32 M0, SGPR6 ; BEFC0306