CODEWITHSUNDEEP
cudanvidia
Damian Andrysiak
Damian Andrysiak

Reputation: 75

Why variables in thread's block have the same memory address? Cuda

I am wondering why they have the same memory address, when If I remember correctly, each thread has a own copy of created variable in this way:

__global__ void
Matrix_Multiplication_Shared(
   const int* const Matrix_A, 
   const int* const Matrix_B, 
         int* const Matrix_C)
{   
    const int sum_value = threadIdx.x;
    printf("%p \n", &sum_value);
}

Output:

enter image description here

I am considering the case of one thread's block, for example with 2 and more threads.

Upvotes: 1

Views: 762

Answers (1)

Greg Smith
Greg Smith

Reputation: 11529

NVIDIA GPUs have multiple address spaces.

The primary virtual address spaced used by pointers is called the generic address space. Inside the generic address space are windows for local memory and shared memory. The rest of the generic address space is the global address space. PTX and the GPU instruction set support additional instructions for 0 based access to the local and shared memory address space.

Some automatic variables and stack memory is in the local memory address space. The primary difference between global memory and the local memory is that local memory is organized such that consecutive 32-bit words are accessed by consecutive thread IDs. If each thread reads or writes from the same local memory offset then the memory access is fully coalesced.

In PTX local memory is accessed via ld.local and st.local.

In GPU SASS the instructions have two forms:

  1. LDL, STL are direct access to local memory given as 0-based offset
  2. LD, ST can be used for local memory access through the generic local memory window.

When you take the address of the variable the generic address space address is returned. Each thread is seeing the same offset from the generic local memory window base pointer. The load store unit will covert the 0-based offset into to a unique per thread global address.

For more information see:

  • CUDA Programming Guide section on Local Memory

  • PTX ISA section on Generic Addressing. Details on local memory are scattered throughout the manual.

Upvotes: 5

Related Questions