llvm dev - Mar 2013 - [LLVMdev] UNREACHABLE executed! error while trying to generate PTX

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 it like so:

#define __constant__ __attribute__((address_space(2)))


On Mon, Mar 18, 2013 at 6:31 PM, upit <uday_pitambare at yahoo.com> wrote:
> 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.
> _______________________________________________
> 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/e8443b82/attachment.html>

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. 





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