3

I need to know something about CUDA shared memory. Let's say I assign 50 blocks with 10 threads per block in a G80 card. Each SM processor of a G80 can handle 8 blocks simultaneously. Assume that, after doing some calculations, the shared memory is fully occupied.

What will be the values in shared memory when the next 8 new blocks arrive? Will the previous values reside there? Or will the previous values be copied to global memory and the shared memory refreshed for next 8 blocks?

Improve this question
3
  • 1
    can you explain a little better? i am not sure i'm following you Commented Feb 17, 2011 at 18:00
  • 1
    My interpretation/simplification: Each block requires all of the shared memory of one processor, ie each processor can host only one block at a time. Enough blocks are launched that the processors will host two blocks in order to complete the kernel execution. Now, looking at one processor... after it executes the first block, the shared memory has been used and may have meaningful values in it. Will these values still be there for the second block, and will they be at the same address as they were the for the previous block? Commented Feb 18, 2011 at 0:18
  • yes that is the answer i want Commented Feb 18, 2011 at 12:19

2 Answers 2

Reset to default
6

It states about the type qualifiers:

  1. Variables in registers for a thread, only stays in kernel
  2. Variables in global memory for a thread, only stays in kernel
  3. __device__ __shared__ type variable in shared memory for a block, only stays in kernel
  4. __device__ type variable in global memory for a grid, stays until the application exits
  5. __device__ __constant__ type variable for a grid, stays until the application exits

thus from this reference, the answer to your question is the memory should be refreshed for the next 8 blocks if they reside in shared memory of your device.

Improve this answer
Sign up to request clarification or add additional context in comments.

Comments

2

For kernel blocks, the execution order and SMs are randomly assigned. In that sense, even if the old value or address preserves, it is hard to keep things in track. I doubt there is even a way to do that. Communication between blocks are done via off chip memory. The latency associated with off chip memory is the performance killer, which makes gpu programming tricky. In Fermi cards, blocks share some L2 cache, but one can't alter the behavior of these caches.

Improve this answer

Comments

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service and acknowledge you have read our privacy policy.

Start asking to get answers

Find the answer to your question by asking.

Ask question

Explore related questions

See similar questions with these tags.