2012-01-25 25 views
11

aşağıdaki konak kod cihaz sürümü gerekir __global__ kernel.cihaz işlev işaretçileri

Bir NVIDIA GeForce GTS 450 (hesaplama yeteneği 2,1)

Michele

peşin teşekkürler ======================= var =================================

bir çalışma çözeltisi

#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; 
} 
+0

Fonksiyonlar işaretçileri bir cihaz kodunda desteklenmez. – Yappie

+0

@Yappie: bu yanlış - Fermi – talonmies

+0

üzerinde işlev göstergeleri destekleniyor CUDA SDK'da gönderilecek bir işlev işaretçisi örneği var ve sorunuza çok benzer bir örnek görebilirsiniz [bu yazı CUDA geliştirici forumlarında ] (http://forums.nvidia.com/index.php?showtopic=156792&view=findpost&p=1201985). – talonmies

cevap

19

işlev işaretçileri Fermi'ye izin verilir. Bu bunu nasıl geçerli:

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; 
} 
+0

Çok teşekkür ederim! Cevabınız benim için çok faydalı oldu. Func_list dizisini dinamik olarak ayırmak mümkün mü? – micheletuttafesta

+0

Dinamik ayırmayı nasıl kullanabileceğinizi göstermek için kodu düzenledim. – brano

+0

brano Yardımlarınız için size sonsuz minnettarım! Ancak bu çalışma çözümünü buldum ... doğru mu? Bir çekirdek – micheletuttafesta

İlgili konular