similar to: Debugging INVALID_OPCODE / MULTIPLE_WARP_ERRORS ?

Also, where's the exit op? Perhaps what's happening is that you don't have an exit and it just goes off executing into the ether? On Tue, Dec 15, 2015 at 12:00 PM, Ilia Mirkin <imirkin at alum.mit.edu> wrote: > A few things that stand out: > > 0: ld u32 %r219 c0[0x0000000000000000+0x0] (0) > > wtf is that 0x0000000000000 thing doing there? Was it a %rX which got

I believe that your problem is this: /*01a0*/ LD R8, [R8]; /* 0x8000000000821c85 */ That needs to be LD.E (and your ST's need to be ST.E). You're using a 32-bit gmem address, but you need to be using a 64-bit one. I believe the 32-bit ones work on fermi, but afaik not on Kepler. Cheers, -ilia On Wed, Dec 16, 2015 at 12:06 PM, Hans de Goede

A few things that stand out: 0: ld u32 %r219 c0[0x0000000000000000+0x0] (0) wtf is that 0x0000000000000 thing doing there? Was it a %rX which got constant-folded into 0? That indirectness should have then been removed... that said, the final encoding looks fine. I believe that kepler has this launch descriptor thing too... is that being set correctly? Please generate a mmt trace, and we can

Hi, On 15-12-15 20:04, Ilia Mirkin wrote: > Also, where's the exit op? Perhaps what's happening is that you don't > have an exit and it just goes off executing into the ether? Sorry I only included a small bit of the program in my original mail because I found the use of "MOV" instructions to load constants suspicious, is that normal ? I've put a log with

BTW, you may be interested in https://github.com/imirkin/mesa/commits/atomic3 which has working ARB_shader_atomic_counters and ARB_shader_storage_buffer_object support (while ripping out things like TGSI_FILE_RESOURCE). Still working on proper memory qualifier support, and obviously need to do some cleanup before upstreaming. Should be getting into a pushable state probably early January. Cheers,

Hi, On 16-12-15 18:24, Ilia Mirkin wrote: > I believe that your problem is this: > > /*01a0*/ LD R8, [R8]; > /* 0x8000000000821c85 */ > > That needs to be LD.E (and your ST's need to be ST.E). You're using a > 32-bit gmem address, but you need to be using a 64-bit one. I believe > the 32-bit ones work on fermi, but afaik not

Hi Micah, > In X86InstrInfo.td for Call Instructions, it mentions that Uses for > argument registers are added manually. Can someone point me to the > location where they are added as the comment doesn't reference a > where or how? the register uses are added by the function X86TargetLowering::LowerCall() during the DAG Lowering phase. This is the relevant code segment: // Add

Thanks Cristoph, I have that code in my backend, but unless I do the following, the registers are never considered 'live' into the call. / Handle a function call let isCall = 1, Defs = [ R0, R1, R2, R3, R4, R5, R6, R7, R8, R9, R10, R11, R12, R13, R14, R15, R16, R17, R18, R19, R20, R21, R22, R23, R24, R25, R26, R27, R28, R29, R30, R31, R32, R33, R34, R35, R36, R37, R38, R39, R40, R41,

Hi Robert, Thanks for the quick response! That goes in line with my observations which is that these things happen when using an ATOM/RED instruction. I've checked and rechecked that I'm generating ops with identical bits as what the proprietary driver does, however (and nvdisasm prints identical output). Could you advise what the proper way of indicating that the memory is

Hello, I've recently come across an error reported by the GPU and would like to know what it means and especially what causes it to be triggered. Any information would be very useful: I'm seeing MP warp error 0x10 (appears in MP register 0x48). This is what we currently have in nouveau: <reg32 offset="0x048" name="TRAP_WARP_ERROR"> <!-- ctx-switched -->

I am trying to generate PTX code for 'nbody' sample program's kernel (nbody_kernel.cu) using clang/LLVM version 3.2. The nbody CUDA program is available in Nvidia's SDK. I am referring to https://github.com/jholewinski/llvm-ptx-samples project. Following are my commands, clang++ -O4 -S -I/usr/local/cuda/include -emit-llvm -target nvptx64 nbody_kernel.cu -o nbody_kernel.ll

Well, I tried the command line given by you and I get the following error clang++ nbody.kernel.cu -Xclang -fcuda-is-device -I/home/upitamba/llvm-3.2.src/tools/clang/test/SemaCUDA/ -Xclang -triple -Xclang nvptx64 -Xclang -target-cpu -Xclang sm_10 -S fatal error: error in backend: Cannot select: 0x334a870: v4f32 = NVPTXISD::MoveParam 0x334a770 [ORD=1] [ID=22] 0x334a770: v4f32 =

Please find the .ll attached below . Yes, I am using the cuda_runtime.h from the toolkit. nbody.kernel.ll <http://llvm.1065342.n5.nabble.com/file/n56048/nbody.kernel.ll> - Uday -- View this message in context: http://llvm.1065342.n5.nabble.com/UNREACHABLE-executed-error-while-trying-to-generate-PTX-tp56026p56048.html Sent from the LLVM - Dev mailing list archive at Nabble.com.

That particular error is fixed in trunk, but with your code I'm now hitting a new issue. I'll get the fix in soon. On Fri, Mar 22, 2013 at 1:08 AM, upit <uday_pitambare at yahoo.com> wrote: > Well, I tried the command line given by you and I get the following error > > clang++ nbody.kernel.cu -Xclang -fcuda-is-device >

I noticed you're using cuda_runtime.h in the source file. Where are you getting this file? From the CUDA toolkit? Since the error is in the back-end, can you just post the .ll or .bc file you are trying to compile? On Mon, Mar 18, 2013 at 12:42 AM, upit <uday_pitambare at yahoo.com> wrote: > I am trying to generate PTX code for 'nbody' sample program's kernel >

The problem you are seeing is because clang is putting the global variables in address space 0, which in NVPTX means the generic address space. PTX does not allow this, so the back-end *should* be printing an error for you. Are you using trunk or 3.2? Generally, clang won't be compatible with the CUDA Toolkit headers. If you want to use the __constant__ modifier from CUDA in Clang, define

OK. That helps. It does flash a warning though [DEVICE-C++] nbody.kernel.cpp nbody.kernel.cpp:29:9: warning: '__constant__' macro redefined #define __constant__ __attribute__((address_space(2))) ^ /opt/cuda/include/host_defines.h:183:9: note: previous definition is here #define __constant__ \ ^ 1 warning generated. Another question is What about extern __shared__ ? I

https://bugs.freedesktop.org/show_bug.cgi?id=80865 Ilia Mirkin <imirkin at alum.mit.edu> changed: What |Removed |Added ---------------------------------------------------------------------------- Summary|Hard hang (GPC0/TPC0/MP |[NVE7] Hard hang |trap: MULTIPLE_WARP_ERRORS |(GPC0/TPC0/MP trap:

https://bugs.freedesktop.org/show_bug.cgi?id=80865 Aaron Plattner <aplattner at nvidia.com> changed: What |Removed |Added ---------------------------------------------------------------------------- Assignee|aplattner at nvidia.com |nouveau at lists.freedesktop.o | |rg

https://bugs.freedesktop.org/show_bug.cgi?id=80865 --- Comment #6 from tethys at gmail.com --- localhost:~# rpm -qa | fgrep mesa mesa-libGL-10.1.5-1.20140607.fc20.i686 mesa-libglapi-10.1.5-1.20140607.fc20.x86_64 mesa-filesystem-10.1.5-1.20140607.fc20.x86_64 mesa-libglapi-10.1.5-1.20140607.fc20.i686 mesa-libGLU-9.0.0-5.fc20.x86_64 mesa-libgbm-10.1.5-1.20140607.fc20.x86_64