编辑：
在我最初写这个答案的时候，我相信这些陈述是正确的：不可能在主机代码中获取内核地址。但是我相信从那时起 CUDA 发生了一些变化，所以现在（在 CUDA 8 中，也许是之前）可以在主机代码中获取 kernel 地址（仍然无法获取地址但是，在主机代码中使用 __device__ 函数。）
原始答案：
这个问题似乎时不时出现，尽管previous examples I can think of 与调用__device__ 函数而不是__global__ 函数有关。
一般来说，在主机代码中获取设备实体（变量、函数）的地址是非法的。
解决此问题的一种可能方法（尽管我不清楚它的实用性；似乎会有更简单的分派机制）是在“设备代码”中提取所需的设备地址并将该值返回给主机，用于调度使用。在这种情况下，我正在创建一个简单的示例，它将所需的设备地址提取到__device__ 变量中，但您也可以编写一个内核来执行此设置（即“给我一个在 GPU 上正确的指针地址”在您的词）。
这是一个粗略的示例，基于您显示的代码：
$ cat t746.cu
#include <stdio.h>
__global__ void ckernel1(){
printf("hello1\n");
}
__global__ void ckernel2(){
printf("hello2\n");
}
__global__ void ckernel3(){
printf("hello3\n");
}
__device__ void (*pck1)() = ckernel1;
__device__ void (*pck2)() = ckernel2;
__device__ void (*pck3)() = ckernel3;
template <typename TFunc, typename... TArgs>
__global__ void Test(TFunc func, int count, TArgs... args)
{
#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 350)
(*func)<< <1, 1 >> >(args...);
#else
printf("What are you doing here!?\n");
#endif
}
template <typename... TArgs>
__host__ void Iterate(void(*kernel)(TArgs...), const int sysInfo, int count, TArgs... args)
{
if(sysInfo >= 350)
{
printf("Iterate on GPU\n");
Test << <1, 1 >> >(kernel, count, args...);
}
else
{
printf("Iterate on CPU\n");
Test << <1, 1 >> >(kernel, count, args...);
}
}
int main(){
void (*h_ckernel1)();
void (*h_ckernel2)();
void (*h_ckernel3)();
cudaMemcpyFromSymbol(&h_ckernel1, pck1, sizeof(void *));
cudaMemcpyFromSymbol(&h_ckernel2, pck2, sizeof(void *));
cudaMemcpyFromSymbol(&h_ckernel3, pck3, sizeof(void *));
Iterate(h_ckernel1, 350, 1);
Iterate(h_ckernel2, 350, 1);
Iterate(h_ckernel3, 350, 1);
cudaDeviceSynchronize();
return 0;
}
$ nvcc -std=c++11 -arch=sm_35 -o t746 t746.cu -rdc=true -lcudadevrt
$ cuda-memcheck ./t746
========= CUDA-MEMCHECK
Iterate on GPU
Iterate on GPU
Iterate on GPU
hello1
hello2
hello3
========= ERROR SUMMARY: 0 errors
$
上述（__device__ 变量）方法可能无法用于模板化子内核，但可以创建一个模板化“提取器”内核，该内核返回（实例化）模板化子内核的地址.我链接的上一个答案中给出了“提取器”setup_kernel 方法的粗略概念。这是模板化子内核/提取器内核方法的粗略示例：
$ cat t746.cu
#include <stdio.h>
template <typename T>
__global__ void ckernel1(T *data){
int my_val = (int)(*data+1);
printf("hello: %d \n", my_val);
}
template <typename TFunc, typename... TArgs>
__global__ void Test(TFunc func, int count, TArgs... args)
{
#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 350)
(*func)<< <1, 1 >> >(args...);
#else
printf("What are you doing here!?\n");
#endif
}
template <typename... TArgs>
__host__ void Iterate(void(*kernel)(TArgs...), const int sysInfo, int count, TArgs... args)
{
if(sysInfo >= 350)
{
printf("Iterate on GPU\n");
Test << <1, 1 >> >(kernel, count, args...);
}
else
{
printf("Iterate on CPU\n");
Test << <1, 1 >> >(kernel, count, args...);
}
}
template <typename T>
__global__ void extractor(void (**kernel)(T *)){
*kernel = ckernel1<T>;
}
template <typename T>
void run_test(T init){
void (*h_ckernel1)(T *);
void (**d_ckernel1)(T *);
T *d_data;
cudaMalloc(&d_ckernel1, sizeof(void *));
cudaMalloc(&d_data, sizeof(T));
cudaMemcpy(d_data, &init, sizeof(T), cudaMemcpyHostToDevice);
extractor<<<1,1>>>(d_ckernel1);
cudaMemcpy((void *)&h_ckernel1, (void *)d_ckernel1, sizeof(void *), cudaMemcpyDeviceToHost);
Iterate(h_ckernel1, 350, 1, d_data);
cudaDeviceSynchronize();
cudaFree(d_ckernel1);
cudaFree(d_data);
return;
}
int main(){
run_test(1);
run_test(2.0f);
return 0;
}
$ nvcc -std=c++11 -arch=sm_35 -o t746 t746.cu -rdc=true -lcudadevrt
$ cuda-memcheck ./t746
========= CUDA-MEMCHECK
Iterate on GPU
hello: 2
Iterate on GPU
hello: 3
========= ERROR SUMMARY: 0 errors
$