similar to: [LLVMdev] R600 -> AMDGPU rename coming on Friday

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,

Hi, Now that the 3.5 branch has been made, I would like to propose renaming the R600 target to AMDGPU. R600 is the name of an AMDGPU GPU that was released about 8 years ago. The R600 backend supports this GPU and also all other GPUs which have been made since then. When people see the name R600 they often assume that only the older GPU family is supported which is not true. The reason that the

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

Hello everyone I was wondering, what the correct way of using an AMDGPU generated kernel code for OpenCL was. I am trying to provide Polly's GPGPU Code generation with the ability to run on different GPU devices, such as AMD GPUs. For NVIDIA, I simply retrieve a pre-compiled PTX string from the NVPTX backend and pass that to OpenCL's 'clCreateProgramWithBinary' function. However,

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

Hi! I'm trying to figure out how to access the workgroup id from within the LLVM IR language when lowering with the AMDGPU backend. Looking at the 'llvm/include/llvm/IR/IntrinsicsAMDGPU.td' file there are intrinsics defined to access the workitem index (thread index), but this file lives in 'llvm/include':

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, 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

> On Dec 5, 2017, at 13:53, Matt Arsenault <arsenm2 at gmail.com> wrote: > > > >> On Dec 5, 2017, at 02:51, Haidl, Michael via llvm-dev <llvm-dev at lists.llvm.org <mailto:llvm-dev at lists.llvm.org>> wrote: >> >> Hi dev list, >> >> I am currently exploring the integration of AMDGPU/ROCm into the PACXX project and observing some

I'm wondering about the focus on bound_cntl. Any cleared bit in the row_mask or bank_mask will also disable updating the result. Brian -----Original Message----- From: Connor Abbott [mailto:cwabbott0 at gmail.com] Sent: Wednesday, June 14, 2017 6:13 PM To: tstellar at redhat.com Cc: Matt Arsenault; llvm-dev at lists.llvm.org; Kolton, Sam; Sumner, Brian; Pykhtin, Valery Subject: Re:

Hi, I am wondering why we dont support PPC32 LE? Here is the output of llvm-mc --version, in which only PPC32, PPC64 & PPC64LE are supported. $ llvm-mc --version LLVM (http://llvm.org/): LLVM version 3.6.2 Optimized build with assertions. Built Aug 2 2015 (11:39:46). Default target: x86_64-apple-darwin15.4.0 Host CPU: core-avx2 Registered Targets: aarch64 - AArch64

Sorry about the formatting... Anyway, I think there may be a misinterpretation of bound_cntl. My understanding is that: 0 => if the source is invalid or disabled, do not write a result 1 => if the source is invalid or disabled, use a 0 instead So the problematic case is where bound_cntl is 0, not when it is 1. -----Original Message----- From: Tom Stellard [mailto:tstellar at redhat.com]

Dear Developers, I compile a OpenCL kernel, FFT, in AMDAPP SDK v2.5 using clang 3.8 + libclc and assembling the code with lld (The LLVM linker). The assembly code contains the following assembly codes (and lots of other similar format assembly) that fails to be parsed by AMDGPUAsmParser. It seems to me that both are valid instructions after looking at the SI instruction spec. s_mov_b32 s0,

Hi, I've prepared patches for both LLVM and Clang to change the datalayout for R600. This may seem like a bold move, but I think it is warranted. R600/SI is a strange architecture in that it uses 64bit pointers but does not support 64 bit arithmetic except for load/store operations that roughly map onto getelementptr. The current datalayout for r600 includes n32:64, which is odd

On 06/14/2017 05:05 PM, Connor Abbott wrote: > On Tue, Jun 13, 2017 at 6:13 PM, Tom Stellard <tstellar at redhat.com> wrote: >> On 06/13/2017 07:33 PM, Matt Arsenault wrote: >>> >>>> On Jun 12, 2017, at 17:23, Tom Stellard <tstellar at redhat.com <mailto:tstellar at redhat.com>> wrote: >>>> >>>> On 06/12/2017 08:03 PM, Connor

> On Dec 6, 2017, at 02:28, Haidl, Michael <michael.haidl at uni-muenster.de> wrote: > > The IR goes through a backend agnostic preparation phase that brings it into SSA from and changes the AS from 0 to 1. This sounds possibly problematic to me. The IR should be created with the correct address space to begin with. Changing this in the middle sounds suspect. > After this

Hi dev list, I am currently exploring the integration of AMDGPU/ROCm into the PACXX project and observing some strange behavior of the AMDGPU backend. The following IR is generated for a simple address space test that copies from global to shared memory and back to global after a barrier synchronization. Here is the IR is attached as as1.ll The output is as follows: 0 0 0 0 0 0 0 0 0 0 0 0 0

On 06/13/2017 07:33 PM, Matt Arsenault wrote: > >> On Jun 12, 2017, at 17:23, Tom Stellard <tstellar at redhat.com <mailto:tstellar at redhat.com>> wrote: >> >> On 06/12/2017 08:03 PM, Connor Abbott wrote: >>> On Mon, Jun 12, 2017 at 4:56 PM, Tom Stellard <tstellar at redhat.com <mailto:tstellar at redhat.com>> wrote: >>>> On

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: > >