Использование __shared__ памяти

Подскажите, пожалуйста, не пойму, как это работает.

У меня есть такой кернел:

  1. __device__ void func()
  2. {
  3. extern __shared__ unsigned char my_array[];
  4. int local_array_size = 256;
  5. unsigned char * local_array = my_array + local_array_size*threadIdx.x;
  6. //здесь используем local_array - shared memory только для этого треда
  7. }

Вызываю кернел я следующим образом:

  1. ...
  2. func<<<Dg,Db,MAX_SHARED_MEMORY_PER_BLOCK>>>();
  3. ...

Все замечательно работает.

Далее я изменяю кернел:

  1. __device__ void func()
  2. {
  3. __shared__ unsigned char my_array[MAX_SHARED_MEMORY_PER_BLOCK*10000];
  4. int local_array_size = 256;
  5. unsigned char * local_array = my_array + local_array_size*threadIdx.x;
  6. //здесь используем local_array - shared memory только для этого треда
  7. }

и его вызов:

  1. ...
  2. func<<<Dg,Db,0>>>();
  3. ...

Однако все работает, и кернел запускается. Как же так? Ведь я прошу под shared память в 10000 раз больше, чем доступно для блока. Почему кернел запускается? И как будет производиться вычисление смещения для локального массива?

Forums: 

Каким образом происходит

Каким образом происходит вызов func<<>>() если func объявлена как
__device__ void func()
__device__ ф-ии не могут быть вызваны с host'а.

Не понял вопроса. Ну

Не понял вопроса. Ну считайте, что <<>> - это и есть оператор вызова с хоста.

Надо понимать, кстати, что вызов __device__ из __device__ - это на самом деле компилятор сделает инлайнинг.

__device__ из __device__

__device__ из __device__ вызывать можно, а вот __device__ из __host__ нельзя. Поэтому для вызова func<<<>>> с хоста func должна быть объявлена как __global__

Извините, ошибся)

Ну конечно же, __global__
Вопрос же состоит в другом:
Почему вызывается кернел, если я объявляю __static__ массив с размером бОльшим, чем то число, которое указано как максимальное количество статической памяти для блока?