__device__ && __global__

Начал писать программу в которой присутствуют спецификаторы функций: __device__ и __global__ но возникла проблема когда я пытаюсь вызвать функцию со спецификатором __device__ из функций со спецификаторами __device__ либо __global__, текст ошибки следующий:
a host function call can not be configured
вызов функций пишу таким же образом как и вызов функции со спецификатором __global__ тоесть:

  1.         int N=subLength;
  2.         int block_size=256;
  3.         int n_blocks=N/block_size+(N%block_size==0?0:1);
  4.         CountLast<<<n_blocks,block_size>>>(10);

Параметры передаю допустимые для работы использую VS C++.
Подскажите как быть.

Forums: 

__device__ и __host__

__device__ и __host__ -функции не нуждаются в конфигурировании запуска.
их выполнение происходит последовательно и синхронно.
только __global__ нужно конфигурировать, ибо только их вызов распараллеливается и асинхронен.

как я понял функция с

как я понял функция с модификатором __device__ запускается и выполняется параллельно?
тоесть в ней нету потоков, блоков, или же она запускается без конфигурирования и сама параллелится в зависимости от параметров запуска ядра?

пример: __device__ void

пример:

  1. __device__ void test_device(){}
  2. __global__ void test_global(){
  3.   test_device(); // КАЖДЫЙ поток в каждом блоке вызывает ф-цию test_device 1 раз
  4. }
  5.  
  6. int main(...){
  7.   test_global<<<2,3>>>(); // запускаем параллельно кернел в конфигурации <<<2,3>>>
  8. }

немного не понял тоесть мне

немного не понял тоесть мне нужно функцию __device__ void test_device() в цикле 6 раз вызвать в цикле?
наподобие:

  1. __device__ void test_device(char c){}
  2. __global__ void test_global(){
  3. char c[2,3];
  4. for(i=0;i<2;i++)
  5. for(i1=0;i1<3;i1++)
  6. test_device(c[i,i1]);
  7. }

я правильно понял?

)), я просто с девайс не

)), я просто с девайс не работал, а примера толкового найти не могу ((
судя по реакции я сделал не правильно тоесть так надо:

  1. __device__ void test_device(char *c){}
  2. __global__ void test_global(){
  3. char c[2,3];
  4. test_device(*c);
  5. }

И он сам её распаралелит?

нет. то что написано внутри

нет. то что написано внутри __global__ функции уже распараллелено автоматически (параметрами конфигурации вызова <<< >>>).
вот еще пример простой:

  1. __global__ void test_global(int *out){
  2.   int i = threadIdx.x; // индекс потока в блоке
  3.   out[i] = i;
  4. }
  5. int main(...){
  6.   int *out;
  7.   cudaMalloc((void **)&out, 4*sizeof(int));
  8.   test_global<<<4>>>(out);
  9. // сейчас массив out = {0, 1, 2, 3}
  10.   return 0;
  11. }

Как я понял распаралеливать

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

Примерно

Примерно правильно.

__device__ надо рассматривать как inline (или макрос с проверкой параметров) - ее тело подставит в место ее вызова компилятором.

Из Appendix'a E1 CUDA

Из Appendix'a E1 CUDA Programming Guide 3.0

By default, a __device__ function is always inlined. The __noinline__ function qualifier however can be used as a hint for the compiler not to inline the function if possible. The function body must still be in the same file where it is called. For devices of compute capability 1.x, the compiler will not honor the __noinline__ qualifier for functions with pointer parameters and for functions with large parameter lists.

Вообще говоря, можно сделать

Вообще говоря, можно сделать так:
1) Пишем функции __device__ в которых выполняется параллельный код
2) Пишем функцию __global__ в которой вызываем все нужные девайс ф-ии, но ф-ии девайс буду работать в конфигурации которая была указанна при запуске глобала.

Те они нужны для того чтобы каждый раз не вызывать глобал с хоста.

Как это нет, если оно так

Как это нет, если оно так работает)))

Да и сами подумайте, при запуске задаем конфигурацию глобал ф-ии, ф-ии девайс можно вызвать только из глобала. Как оно может работать иначе.

Возможно в предыдущим посте было написанно немного не ясно и вы не совсем поняли что я имел ввиду.

Я понял(ещё в том посте) что

Я понял(ещё в том посте) что вы имели в виду.
А теперь рассмотрим такой случай:
Нужно сделать две последовательные операции, выход первой операции подаётся на вход второй.
И что? Никто не даёт гарантий, в том что при вашей методике, все потоки первой операции закончатся одновременно - поэтому соответствующая картинка

Если учесть что до CC 2.0, __device__ были обычными инлайнами, то их первоначальное призвание это уменьшить копипасту, более лучше структурировать код.
Никто вручную подставлять код не запрещал(если бы не было __device__), и вызов с хоста в ЭТИХ случаях не обязателен.
Там где можно обойтись без вызова с хоста, нужно обходится без этих вызовов - но это делается и без __device__ .

С появлением CC 2.0, __device__ теперь могут быть не инлайнами(что уменьшает размер результирующего кода, но мешает оптимизации). Основным же нововведением CC 2.0 относительно __device__ , я считаю вызов функции по указателю (ещё не пробовал).

Насчет СС2 , насколько помню

Насчет СС2 , насколько помню он поддерживается начиная с ферми.
Хотелось бы узнать насколько вы разобрались в достоинствах этих чипов и если не против задал бы несколько вопросов.

Да с ферми, это не означает

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

не против задал бы несколько вопросов.

Против.
почему, ответ в конце сообщения
http://www.gpgpu.ru/node/238#comment-1726

да, __device__ функции ни чем

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