장치 기능 포인터

Question

난 다음 호스트 코드 디바이스 버전이 필요) __global__ 커널.

은 내가 NVIDIA 지포스 GTS 450 (연산 능력 2.1) 이

미셸

사전 에 감사 =======================이 =================================

작동 용액

#define REAL double 

typedef REAL (*func)(REAL x); 

__host__ __device__ REAL func1(REAL x) 
{ 
    return x+1.0f; 
} 

__host__ __device__ REAL func2(REAL x) 
{ 
    return x+2.0f; 
} 

__host__ __device__ REAL func3(REAL x) 
{ 
    return x+3.0f; 
} 

__device__ func func_list_d[3]; 
func func_list_h[3]; 

__global__ void assign_kernel(void) 
{ 
    func_list_d[0]=func1; 
    func_list_d[1]=func2; 
    func_list_d[2]=func3; 
} 

void assign(void) 
{ 
    func_list_h[0]=func1; 
    func_list_h[1]=func2; 
    func_list_h[2]=func3; 
} 


__global__ void test_kernel(void) 
{ 
    REAL x; 
    for(int i=0;i<3;++i){ 
     x=func_list_d[i](2.0); 
     printf("%g\n",x); 
    } 
} 

void test(void) 
{ 
    REAL x; 
    printf("=============\n"); 
    for(int i=0;i<3;++i){ 
     x=func_list_h[i](2.0); 
     printf("%g\n",x); 
    } 
} 

int main(void) 
{ 
    assign_kernel<<<1,1>>>(); 
    test_kernel<<<1,1>>>(); 
    cudaThreadSynchronize(); 

    assign(); 
    test(); 

    return 0; 
} 

Answer 1

함수 포인터는 페르미에 허용. 다음과 같이 수행 할 수 있습니다.

typedef double (*func)(double x); 

__device__ double func1(double x) 
{ 
return x+1.0f; 
} 

__device__ double func2(double x) 
{ 
return x+2.0f; 
} 

__device__ double func3(double x) 
{ 
return x+3.0f; 
} 

__device__ func pfunc1 = func1; 
__device__ func pfunc2 = func2; 
__device__ func pfunc3 = func3; 

__global__ void test_kernel(func* f, int n) 
{ 
    double x; 

    for(int i=0;i<n;++i){ 
    x=f[i](2.0); 
    printf("%g\n",x); 
    } 
} 

int main(void) 
{ 
    int N = 5; 
    func* h_f; 
    func* d_f; 
    h_f = (func*)malloc(N*sizeof(func)); 
    cudaMalloc((void**)&d_f,N*sizeof(func)); 

    cudaMemcpyFromSymbol(&h_f[0], pfunc1, sizeof(func)); 
    cudaMemcpyFromSymbol(&h_f[1], pfunc1, sizeof(func)); 
    cudaMemcpyFromSymbol(&h_f[2], pfunc2, sizeof(func)); 
    cudaMemcpyFromSymbol(&h_f[3], pfunc3, sizeof(func)); 
    cudaMemcpyFromSymbol(&h_f[4], pfunc3, sizeof(func)); 

    cudaMemcpy(d_f,h_f,N*sizeof(func),cudaMemcpyHostToDevice); 

    test_kernel<<<1,1>>>(d_f,N); 

    cudaFree(d_f); 
    free(h_f); 

    return 0; 
}