[LLVMdev] UNREACHABLE executed! error while trying to generate PTX
justin.holewinski at gmail.com
Wed Mar 20 10:00:20 PDT 2013
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 definition is here
> #define __constant__ \
> 1 warning generated.
> Another question is
> What about extern __shared__ ?
> I can see that the error goes away if I replace "extern __shared__ float4
> sharedPos" with "__shared__ float4* sharedPos;". Do I have to dynamically
> allocate the shared memory by specifying size in kernel Launch? If so, why
> doesn't the second use of the same statement in another function cause the
> error ?
> I am using 3.2.
I would just do away with the toolkit headers. I may try to put together
some minimalistic headers for clang w/ nvptx at some point. Your best bet
is to just define what you need yourself for now.
__shared__ would be address space 3, so:
#define __shared__ __attribute__((address_space(3)))
Either using  or * should work. Just be aware that you will need to
specify a shared size when you launch the kernel. You can get the address
space mapping from lib/Target/NVPTX/MCTargetDesc/NVPTXBaseInfo.h.
> View this message in context:
> Sent from the LLVM - Dev mailing list archive at Nabble.com.
> LLVM Developers mailing list
> LLVMdev at cs.uiuc.edu http://llvm.cs.uiuc.edu
-------------- next part --------------
An HTML attachment was scrubbed...
More information about the llvm-dev