Почему расположение указателя разделяемой памяти в двух блоках одинаково?

Question

Я думал, что общая память устройства cuda является частной для блока. Тем не менее, мне кажется, что указатель общей памяти между двумя блоками идентичен:

#include <stdio.h>

__global__ void foo() {
    __shared__ int ar[8];
    printf("shared memory pointer %p at blockidx %i\n", ar, blockIdx.x);
}

int main() {
    dim3 blockDim(1);
    dim3 gridDim(2);
    foo<<<gridDim, blockDim>>>();
    cudaDeviceSynchronize();
}

Выполнение кода выше дает:

shared memory pointer 0x7f88f5000000 at blockidx 0
shared memory pointer 0x7f88f5000000 at blockidx 1

С этой программой я ожидал создать два разных блока, инициализируйте разделяемую память в каждом блоке и получите два разных места для памяти. Я что-то неправильно понимаю? Или эти указатели действительно имеют другое физическое местоположение, но внутри блока адреса кажутся одинаковыми?

Answer 1

Общая память является приватной, то есть потоки из одного блока не могут получить доступ к общей памяти другого блока.

... именно по этой причине на самом деле следует ожидать, что адрес диапазон для разделяемой памяти будет одинаковым для всех блоков - но в каждом блоке загрузка или сохранение по этим адресам влияет на локальную разделяемую память блока.

Для интуиции: это несколько подобно тому, как в ЦП код в двух процессах может использовать идентичные адреса указателей, но они фактически будут обращаться к разным физическим местам в памяти (обычно).