Как можно использовать shared memory в device функциях

Как я понимаю указатели можно использовать только на область глобальной памяти. Как тогда изменять значение shared memory в device функции?

Forums: 

в своем кернеле (который

в своем кернеле (который __global__ void) можно создать переменную (например матрицу) с __shared__ перед ней. затем при вызове __device__ функции передавать созданную переменную как параметр, так в этой функции можно "изменять значение shared memory в device функции".

__device__ void initMatrix(double matrix[16]){
// здесь что нибудь делать с этой матрицей
}

__global__ void blah(){
__shared__ double matrix[16];
initMatrix(matrix);
// дальнейшие действия с матрицей
}

Только не следует забывать что эта переменная(матрица, все что угодно) объявленная с помощью __shared__ будет одна на все потоки блока (один диапазон памяти), и потому нужно разобраться с индексацией и размером этой матрицы (чтобы каждый поток читал/писал только свою часть матрицы). А так же не забыть про т.н. bank conflicts (подробнее см в Whitepaper прилагаемый к New Matrix Transpose в SDK вроде с 2.3)

п.с. мы ж про куду говорим? :)))))

То есть если я сделаю как в

То есть если я сделаю как в ниже приведеном примере - то в функцию d_foo фактически произойдет передача указателья на область shared памяти а не передача по значению ..... или я чего то не так понял

  1. __global__  g_foo(...){
  2. extern  __shared__  float sdata[];
  3. .....
  4. d_foo(sdata[threadIdx.x])
  5. .....
  6. }
  7.  
  8. __device__ int d_foo(float shared){
  9. ....
  10. shared = ......;
  11. .....
  12. }

> мы ж про куду говорим
)))))) ага про нее

я начинаю подозревать что это

я начинаю подозревать что это я чего то не так понял :)
Насколько мне известно, shared не может быть extern - потому то он и shared, то есть один на мультипроцессор, ниоткуда извне он не возьмется. но для пущей уверенности проверил: не компилируется, говорит "__local__ and __shared__ variables cannot have external linkage". Что вполне логично.
а вопрос о передаче вообще не стоит - device-функция "gets inlined", то есть на каком то этапе компиляции она просто вставляется в твой global кернел. типа как #define работает.

Так вроде с gets inlined

Так вроде с gets inlined разобрался

Странно, но в $SDK_CUDA/template/template_kernel.cu:58: extern __shared__ float sdata[]; Да и во многих других проэктах СДК extern __shared__ используется.

Таааак где в доке написано про extern :-[.

не разбираясь, вспоминаю про

не разбираясь, вспоминаю про динамически и статически выделяемую общую (shared) память. одна из них - которая декларируется внутри кернела, вторая - количество которой передается как параметр внутри синтаксиса <<< >>>. нужно искать в доках и вправду.

В разделе B.2.3 Programming

В разделе B.2.3 Programming Guide (номер раздела - для версии 2.3, в других может быть другим) про это все написано

Т.е. вы можете завести extern __shared__ array[] (один!) размер которого задается динамически, а потом его куски использовать для чего хотите.

Но он тоже будет не инициализированным и его надо самому инициализировать в блоке.

В этом примере должна

В этом приведенном выше примере должна произойти передача по значению. То есть я предполагаю что значения sdata[threadIdx.x] должно скопироваться в регистровую память и там произвести изменения. Таким образом в d_foo() значение sdata не поменяется. По логике Си мне бы надо передать указатель на эту область памяти &sdata[threadIdx.x] и в device уже писать в эту область памяти, те я бы сделал так

  1. __global__  g_foo(...){
  2. extern  __shared__  float sdata[];
  3. .....
  4. d_foo(&sdata[threadIdx.x])
  5. .....
  6. }
  7.  
  8. __device__ int d_foo(float* shared){
  9. ....
  10. *shared = ......;
  11. .....
  12. }

но указатели можно использовать только на область глобальной памяти.

Это же C: - если вы

Это же C:
- если вы передаете float array[] - то передача идет по ссылке
- а если float array[i] - то по значению

Так и изменять - записывать

Так и изменять - записывать туда что-то.

Но эти изменения будут видны только внутри блока, а треды из другого блока будут читать/писать в свою shared mem.