глобальный указатель на константную память

к которому получаешь доступ из девайса.
а чё, нельзя? :)
ругацца.

Forums: 

читать.я наивно полагал что

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

__constant__ double* device_blah;

__global__ void kernel(){
double a = device_blah[0]; // тут ошибочка доступа
}

Ну в общем это факт, я просто __констатирую__ (сорри за тавтологию :) )

Не совсем в тему

Можно ли каким-нибудь образом скопировать данные из девайсной памяти в константную (Хостом).

ни разу не встречался с

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

вроде можно,

вроде можно, функция
cudaMemcpyToSymbol ( const char * symbol, const void * src,
size_t count, size_t offset,
enum cudaMemcpyKind kind )

"Copies count bytes from the memory area pointed to by src to the memory area pointed to by offset bytes from the start of symbol symbol. The memory areas may not overlap. symbol can either be a variable that resides in global or constant memory space, or it can be a character string, naming a variable that resides in global or constant memory space. kind can be either cudaMemcpyHostToDevice or cudaMemcpyDeviceToDevice."

Используем cudaMemcpyDeviceToDevice, src - твоя память девайса, symbol - переменная в константной памяти.

Вот что имел ввиду: в

Вот что имел ввиду: в ProgrammingGuide - cudaMemcpyToSymbol(constData, data, sizeof(data));
в ReferenceManual - cudaMemcpyToSymbol (const char *symbol, const void *src, size_t count, size_t
offset, enum cudaMemcpyKind kind);
И так как при компиляции 1-го варианта не пишет "Error: function ... does not take 3 parameters", следовательно она либа перегружаемая, либо в прототипе написано: cudaMemcpyToSymbol (const char *symbol, const void *src, size_t count, size_t offset=0, enum cudaMemcpyKind kind=cudaMemcpyHostToDevice)

Разобрался с этим.

Разобрался с этим. Оказывается (надо было читать доки перед использованием) нельзя использовать константную память так как глобальную. Её не выделишь, так как она выделяется автоматически во время старта программы и уничтожается во время завершения программы.

__constant__ float constData[256]; // говорит программе что нужно столько то байт выделить
float data[256]; // исходные данные
cudaMemcpyToSymbol(constData, data, sizeof(data)); // копируем исходные данные в конст память

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

п.с. немного нубский вопрос
если у меня в константной памяти лежит массив double constant_a[16][16] и мне надо передать этот указатель в какую-нибудь функцию, то я не могу сделать это как:

  1. __constant__ double constant_a[16][16];</p>
  2. <p>__device__ void someFunction(double constant_a[16][16]){<br />
  3.     // blah<br />
  4. }</p>
  5. <p>__global__ void kernel(){<br />
  6.     someFunction(constant_a);<br />
  7. }<br />

Так же не могу и так:

  1. __constant__ double constant_a[16][16];</p>
  2. <p>__device__ void someFunction(double* constant_a){<br />
  3.     // blah<br />
  4. }</p>
  5. <p>__global__ void kernel(){<br />
  6.     someFunction(constant_a);<br />
  7. }

Но могу так:

  1. <br />
  2. __constant__ double constant_a[16][16];</p>
  3. <p>__device__ void someFunction(double* constant_a){<br />
  4.     // blah<br />
  5. }</p>
  6. <p>__global__ void kernel(){<br />
  7.     someFunction(constant_a[0]);<br />
  8. }

По сути то я передаю указатель на первый элемент первой строки. Гарантировано ли что вторая строка будет лежать сразу после первой?

Гарантировано ли что вторая

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

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

память под конст массив

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

Ну про то и вопрос -

Ну про то и вопрос - объявлено 16x16, если оно всегда так, то почему не просто [256] ?

А, да, и в последнем случае

А, да, и в последнем случае вы передали значение, а не указатель. - я невнимательно посмотрел, пардон А если вам нужно передать указатель на кусочек константного массива.... я бы не мучался (и не думал) и передал бы просто индекс в нем, ведь сам массив - глобален и из всех __device__ функций виден.

Значение? Хм, спасибо :) Мне

Значение? Хм, спасибо :)

Мне нужно передавать указатель потому что у меня несколько структур в конст памяти, которые обрабатываются одной и той же функцией.
Значит надо someFunction(&constant_a[0])?

Нет, я ошибся от

Нет, я ошибся от невнимательности, не увидел второе измерение в массиве 16x16

Да, передается адрес строки, по идее все правильно. У вас размер статический, всегда 16x16 или в зависимости от вызова может меняться?

подебагил - правильно

подебагил - правильно передается.
статический. вроде все правильно :)
только все равно unspecified launch failure ))) не факт конечно что в этом месте.