similar to: [LLVMdev] UNREACHABLE executed! error while trying to generate PTX

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 >

Hi all, As part of my compute work I'm trying to get some TGSI compute code to work. The code from mesa/src/gallium/tests/trivial.c works. So now I'm trying to get a "native" tgsi kernel to run via clover, I'm using Francisco's nbody.c example for this: https://fedorapeople.org/~jwrdegoede/nbody.c Which does not work, at first I thought there was an issue with the

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 =

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

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 >

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.

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

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

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

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,

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 Steven- Nice, but will this not break Windows? From an initial glance over your patch, it seems to assume the existence of mmap() in some form or other. Alistair On 8 Aug 2010, at 03:05, Steven Noonan wrote: > Hi folks, > > I've been doing work on memory reduction in Unladen Swallow, and > during testing, LiveRanges seemed to be consuming one of the largest > chunks of

Hi folks, I've been doing work on memory reduction in Unladen Swallow, and during testing, LiveRanges seemed to be consuming one of the largest chunks of memory. I wrote a replacement allocator for use by BumpPtrAllocator which uses mmap()/munmap() in place of malloc()/free(). It has worked flawlessly in testing, and reduces memory usage quite nicely in Unladen Swallow. The code is available

Not really. Clang does not have a way to annotate device vs. kernel functions in C/C++ mode. You're probably better off trying to use OpenCL or CUDA mode in clang. In the clang unit tests, there is a cuda.h header that provides very basic support for these keywords: tests/SemaCUDA/cuda.h If you compile as CUDA (use .cu extension, or "-x cuda") and use this header, you will have

Thanks a lot Justin, I will remove the toolkit header. Just one last question..(maybe ;) ) If I do away with toolkit headers it says unknown type name '__device__'. Does this function qualifier have an alternative ? or I can just do away with ? -- View this message in context: http://llvm.1065342.n5.nabble.com/UNREACHABLE-executed-error-while-trying-to-generate-PTX-tp56026p56093.html

On Wed, Mar 20, 2013 at 11:29 AM, upit <uday_pitambare at yahoo.com> wrote: > 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

Ralf, What version of the SDK were you using for your analysis? I don't see that in the slides/pdf. Thanks, Micah > -----Original Message----- > From: llvmdev-bounces at cs.uiuc.edu [mailto:llvmdev-bounces at cs.uiuc.edu] > On Behalf Of Ralf Karrenberg > Sent: Wednesday, October 19, 2011 2:13 PM > To: llvmdev at cs.uiuc.edu > Subject: Re: [LLVMdev] ANN: libclc (OpenCL C

Hi Micah, The numbers from the paper were measured with the ATI Stream SDK v2.1 (it's only mentioned in the references I think). The most recent measurements I have were done with the current v2.5. Best, Ralf Am 19.10.2011 23:43, schrieb Villmow, Micah: > Ralf, > What version of the SDK were you using for your analysis? I don't see that in the slides/pdf. > > Thanks, >