search for: r600

...n, so I might be able to triage at least the code that hangs with this patch On Wed, 2014-09-24 at 20:27 -0400, Jan Vesely wrote: > v2: tighten the sub64 tests > v3: rename to CARRY/BORROW > > Signed-off-by: Jan Vesely <jan.vesely at rutgers.edu> > > --- > lib/Target/R600/AMDGPUISelLowering.h | 2 + > lib/Target/R600/AMDGPUInstrInfo.td | 6 ++ > lib/Target/R600/AMDGPUSubtarget.h | 8 ++ > lib/Target/R600/EvergreenInstructions.td | 3 + > lib/Target/R600/R600ISelLowering.cpp | 39 +++++++- > test/CodeGen/R600/add.ll...

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

..."yeah, the structure looks fine and there aren't any obviously gratuitous code style violations" is the sort of review I'm looking for. I'm not sure a great way to get this, but asking on llvmdev is probably a good start. > Hi LLVM Developers, I would like to merge the R600 backend from branches/R600 into TOT prior to the LLVM 3.2 release. As you can see in the above paragraph, Chris has asked me to find someone to review the code. If there is anyone that would be willing to take the time and look at the code, I would really appreciate it. For your convenience, I h...

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

Hi, When trying to compile a trivial opencl kernel such as: __kernel void if_eq(__global int * out, int arg0, int arg1){ out[0] = arg0==arg1?0:1; } Clang generates IR like: %1 = icmp eq i32 %arg0, %arg1 %. = zext i1 %1 to i32 This eventually crashes ISel on R600. Attached patch adds a selector so it will compile. Regards, Jon Pry jonpry at gmail.com -------------- next part -------------- An HTML attachment was scrubbed... URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20131230/057c13c5/attachment.html> -------------- next part --------...

Hi Tom, Thanks a lot for your explanations, now it makes a lot more sense ;) I had a slightly closer look at the R600 packetizer, and the issue is that the third LSHL instruction has both an implicit use and *afterwards* an implicit def of T1_XYZW. The latter def causes the current ScheduleDAGInstrs implementation to ignore the implicit use, thus the ScheduleDAG only contains an anti-dependency from the second...

...re looks fine and there aren't any obviously gratuitous code style violations" is the sort of review I'm looking for. I'm not sure a great way to get this, but asking on llvmdev is probably a good start. >> > > Hi LLVM Developers, > > I would like to merge the R600 backend from branches/R600 into TOT > prior to the LLVM 3.2 release. As you can see in the above paragraph, > Chris has asked me to find someone to review the code. If there is > anyone that would be willing to take the time and look at the code, > I would really appreciate it. For y...

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 another backend in the tree, how would that wo...

...e I patched the buildSchedGraph() function to iterate over all defs first, then over all uses in a second iteration (maybe not the most efficient solution, but it solves the issue). I am attaching the corresponding patch. However, while upgrading our code to LLVM 3.4, I found that the test/CodeGen/R600/store.ll test fails for -mcpu=redwood due to this change. The difference is in the following line in @store_i8: With my patch applied I get: LSHL * T0.W, PV.W, literal.x, 3(4.203895e-45), 0(0.000000e+00) LSHL * T1.X, T1.W, PV.W, LSHL...

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 or if you prefer you can download the entire tree with this link: http://cgit.freedesktop.org/~tstellar/llvm/snapshot/llvm-r600-initial-review.tar.gz The R600 backend is located in lib/Target/AMDIL First, a brief description of the backend: The r600 backend is being develope...

...ren't any obviously gratuitous code style violations" is the sort of review I'm looking for. I'm not sure a great way to get this, but asking on llvmdev is probably a good start. > >> > > > > Hi LLVM Developers, > > > > I would like to merge the R600 backend from branches/R600 into TOT > > prior to the LLVM 3.2 release. As you can see in the above paragraph, > > Chris has asked me to find someone to review the code. If there is > > anyone that would be willing to take the time and look at the code, > > I would really a...

> -----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" <thomas.stellard at amd.com> > wrote: > > > > > > On Fri, May 25, 2012 at 02:37:26PM -0700, Justin Holewin...

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 ;...

...bviously gratuitous code style violations" is the sort of review I'm looking for. I'm not sure a great way to get this, but asking on llvmdev is probably a good start. >>>> >>> >>> Hi LLVM Developers, >>> >>> I would like to merge the R600 backend from branches/R600 into TOT >>> prior to the LLVM 3.2 release. As you can see in the above paragraph, >>> Chris has asked me to find someone to review the code. If there is >>> anyone that would be willing to take the time and look at the code, >>> I wo...

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

...t;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 > > 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" <thomas.stellard at amd.com> > > wrote: > > > > > > > > On Fri, May 25, 2012 at 02:...

...compiled llvm-3.3 from sources. My machine has these: --cpu -amd64 --os- clbfs 64bit -linux3 series kernel --C/C++ compiler -modern versions gcc My preferred build tool is cmake. llvm is needed for MesaLib and the autotools build of llvm allows this option:- --enable-experimental-targets=R600 ( for example described here: http://www.linuxfromscratch.org/blfs/view/svn/general/llvm.html ) I have been unable to locate a equivalent option in the CMakeLists.txt in the base llvm-3.3 directory I have also searched the llvm cmake documentation here:- http://llvm.org/releases/3.3/docs/CMake...

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 addi...

On Mon, Mar 26, 2012 at 01:22:01PM -0400, Villmow, Micah wrote: > 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. > I've started working on the tests, and I've pushed a few up to my llvm repo. I can add the AMDIL tests too if you want to send them. I think they will be helpful. > That brings me to the second thing. Are the AMDIL backend and the R600 backend the same, or not? At this point,...

..., > > 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 or if you prefer you can download the entire tree with this link: > http://cgit.freedesktop.org/~tstellar/llvm/snapshot/llvm-r600-initial-review.tar.gz > The R600 backend is located in lib/Target/AMDIL > > First, a brief description of the backend: > >...