0

I'm trying to do something very very simple with shared memory in jcuda. My kernel:

__shared__ int testMe;

extern "C"
__global__ void test() {
  testMe = 5;
}

Making shared memory global allows me to use it in device functions, unfortunately, I can't even declare shared memory global. Cuda just crashes with kernel cannot launch issue. I've also tried doing:

extern __shared__ int test

However nvcc compiler gives, cannot externally link shared memory error. I really want to get this to work to make writing kernels easier, but I don't see a way.

Dr.Knowitall
  • 10,080
  • 23
  • 82
  • 133
  • That makes no sense. Shared memory can only be declared *inside* a kernel function, .e.g. [read here](http://devblogs.nvidia.com/parallelforall/using-shared-memory-cuda-cc/). – Kerrek SB Jun 26 '15 at 08:54
  • @KerrekSB: That is not true. It is perfectly legal to declare shared memory at translation unit scope, when the memory will be dynamically allocated at runtime – talonmies Jun 26 '15 at 12:00
  • This question doesn't make much sense. What is stopping you from passing the shared variable by reference to a device function? The compiler is perfectly capable of handling that usage case. – talonmies Jun 26 '15 at 13:10

1 Answers1

1

You can't define static shared memory at translation unit scope, you have to declare it at kernel scope. If you do so, your shared memory will be the same memory for every thread in the same block. If you want to use memory that is same for every thread in all blocks, you have to use global memory (call cudaMalloc() in host-code and pass a pointer of the memory to your kernel).

talonmies
  • 70,661
  • 34
  • 192
  • 269
Klaus
  • 538
  • 8
  • 26
  • 1
    Actually, you can declare global shared memory, but only if it is dynamically allocated pointer or array. – talonmies Jun 26 '15 at 13:12
  • @talonmies ok but then it's another thing and in no way related to the `__shared__` keyword – Klaus Jun 26 '15 at 14:41
  • 1
    No. It is related to the `__shared__` keyword. There are two ways to declare shared memory, static or dynamic, and for the dynamic case it is perfectly acceptable to declare outside of a kernel (at translation unit scope) – talonmies Jun 26 '15 at 14:43