upit
2013-Mar-20 15:29 UTC
[LLVMdev] UNREACHABLE executed! error while trying to generate PTX
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.
Justin Holewinski
2013-Mar-20 17:00 UTC
[LLVMdev] UNREACHABLE executed! error while trying to generate PTX
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>
upit
2013-Mar-20 19:29 UTC
[LLVMdev] UNREACHABLE executed! error while trying to generate PTX
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 Sent from the LLVM - Dev mailing list archive at Nabble.com.
Maybe Matching Threads
- [LLVMdev] UNREACHABLE executed! error while trying to generate PTX
- [LLVMdev] UNREACHABLE executed! error while trying to generate PTX
- [LLVMdev] UNREACHABLE executed! error while trying to generate PTX
- [LLVMdev] UNREACHABLE executed! error while trying to generate PTX
- instrumenting device code with gpucc