search for: sharedpos

...ribute__((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:...

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

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

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.