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

[Justin Holewinski]

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:
> http://llvm.1065342.n5.nabble.com/UNREACHABLE-executed-error-while-trying-to-generate-PTX-tp56026p56080.html
> 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
> http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
>



-- 

Thanks,

Justin Holewinski
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20130320/8466395d/attachment.html>