Как я понимаю указатели можно использовать только на область глобальной памяти. Как тогда изменять значение shared memory в device функции?
Как можно использовать shared memory в device функциях
[x]
Вход
Amazon
AMD
ATI
brute force
bruteforce
cloud
CUDA
GPGPU
gpgpu.ru
GPU Gems
Intel
Larrabee
Linpack
MapReduce
MD5 crack
Nexus
NVidia
NVidia 8800
NVidia CUDA
NVidia G200
NVidia GTX280
NVidia Nexus
OpenCL
Parallel Nsight
signal processing
sparse matrices
Stream SDK
VISPL
VMWare
web
ВМиК МГУ
МГУ
Москва
Т-Платформы
Физфак МГУ
бенчмарки
блогосфера
вычисления
конкурсы
курсы
новости сайта
обработка изображений
подбор паролей
поиск
программирование GPU
работа
разное
сортировка
фильтрация трафика
численные методы
Navigation
Cвежие комментарии
-
3 days 8 hours ago
-
4 days 6 hours ago
-
4 days 6 hours ago
-
4 days 7 hours ago
-
4 days 7 hours ago
-
4 days 8 hours ago
-
4 days 8 hours ago
-
1 week 4 hours ago
-
2 weeks 6 days ago
-
3 weeks 1 day ago
Новое на форуме
Популярно
- Как начать с самого начала работу с CUDA (33,691)
- Форумы NVidia CUDA: обзор за май (31,715)
- GPGPU и видеокарты AMD (18,141)
- NVidia GTX 280, Tesla T10P (15,735)
- SGEMM на видеокарте и CPU, серия 6 (14,878)
Comments
Так и изменять - записывать туда что-то.
Но эти изменения будут видны только внутри блока, а треды из другого блока будут читать/писать в свою shared mem.
в своем кернеле (который __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 памяти а не передача по значению ..... или я чего то не так понял
> мы ж про куду говорим
)))))) ага про нее
В этом приведенном выше примере должна произойти передача по значению. То есть я предполагаю что значения sdata[threadIdx.x] должно скопироваться в регистровую память и там произвести изменения. Таким образом в d_foo() значение sdata не поменяется. По логике Си мне бы надо передать указатель на эту область памяти &sdata[threadIdx.x] и в device уже писать в эту область памяти, те я бы сделал так
но указатели можно использовать только на область глобальной памяти.
Это же C:
- если вы передаете float array[] - то передача идет по ссылке
- а если float array[i] - то по значению
я начинаю подозревать что это я чего то не так понял :)
Насколько мне известно, shared не может быть extern - потому то он и shared, то есть один на мультипроцессор, ниоткуда извне он не возьмется. но для пущей уверенности проверил: не компилируется, говорит "__local__ and __shared__ variables cannot have external linkage". Что вполне логично.
а вопрос о передаче вообще не стоит - device-функция "gets inlined", то есть на каком то этапе компиляции она просто вставляется в твой global кернел. типа как #define работает.
Так вроде с gets inlined разобрался
Странно, но в $SDK_CUDA/template/template_kernel.cu:58: extern __shared__ float sdata[]; Да и во многих других проэктах СДК extern __shared__ используется.
Таааак где в доке написано про extern :-[.
не разбираясь, вспоминаю про динамически и статически выделяемую общую (shared) память. одна из них - которая декларируется внутри кернела, вторая - количество которой передается как параметр внутри синтаксиса <<< >>>. нужно искать в доках и вправду.
В разделе B.2.3 Programming Guide (номер раздела - для версии 2.3, в других может быть другим) про это все написано
Т.е. вы можете завести extern __shared__ array[] (один!) размер которого задается динамически, а потом его куски использовать для чего хотите.
Но он тоже будет не инициализированным и его надо самому инициализировать в блоке.