Я имею в виду, разве функция не является символом?
Нет, это не так.
У меня нет особого понимания, но нет сомнений в том, что отчасти это объясняется историей: когда были изобретены API CUDA, функции __device__ были всего лишь средством программирования. Не было ни ABI, ни поддержки указателей на функции, и все функции устройства были встроены в компилятор. Единственными статическими символами устройства c, которые были отправлены, были __global__ функции, ссылки на текстуры и __device__ переменные. Таким образом, не было абсолютно никакого способа, чтобы такое использование было предусмотрено или возможно, когда язык и API были объединены в течение 15 лет go.
Даже с сегодняшней версией ABI и инструментарием устройства формата ELF (первоначально все было простым текстом со встроенными строками), вы не найдете функций __device__, предоставляемых интерфейсом ELF объектного файла устройства. Невозможно извлечь произвольную функцию __device__ через любой из API хоста, в отличие от функций __global__ и других символов устройства.
И если адрес на стороне устройства глобального указателя может быть " Известно "моему коду хоста, почему не может сама функция?
См. Выше. API никогда не раскрывали это.
И если это не работает - почему он компилируется, а не жалуется?
Из-за траектории компиляции. Внешний интерфейс CUDA делает это с вашей __device__ функцией в хост-коде (и здесь нет различия, он делает это с каждой __device__ функцией, включая внутренние функции набора инструментов и библиотеки устройств):
# 3 "unobtainium.cu"
__attribute__((unused)) int f1() {int volatile ___ = 1;::exit(___);}
#if 0
# 3
{ printf("dev f1\n"); return 0; }
#endif
т.е. он создает фиктивную заглушку хоста, чтобы все компилировалось. Ядра и символы устройств также получают заглушки, но с другим шаблоном. Эти стандартные заглушки совпадают с тегами, которые используются внутренними функциями времени выполнения для обеспечения работы API времени выполнения на стороне хоста. Но функции устройства этого не делают, потому что они не предоставляются API-интерфейсами кода устройства CUDA.
И, наконец, ваш оригинальный вопрос:
Почему это работает:
#include <stdio.h>
__device__ int f1(){ printf("dev f1\n"); return 0;}
__device__ void *fptrf1 = (void*) f1;
__global__ void mykernel(int (*fptr)()) {
fptr();
printf("executed\n");
}
int main() {
void *hf1;
cudaMemcpyFromSymbol(&hf1, fptrf1, sizeof(int *));
mykernel<<<1,1>>>((int (*)())hf1);
cudaDeviceSynchronize();
}
Интересно то, что это не всегда работало. Когда-то вам приходилось запускать установочное ядро для инициализации указателя функции на стороне устройства. Где-то около CUDA 5 он начал работать таким образом. Почему это относительно просто - переменная области действия __device__ модуля компиляции является допустимым символом устройства, поэтому предоставляется интерфейсом API хоста, и компоновщик на стороне устройства может (сейчас) статически назначить правильное значение во время связывания, так что при инициализации среды выполнения значение является правильным. Но обратите внимание, что это статическое присвоение c, ничего, что происходит во время выполнения.