은 쿠다에서 템플릿 매개 변수로 사용됩니다

Question

나는 이진 연관 연산자를 사용하는 쿠다 (이것은 운동입니다. 나는 다른 사람들이 더 잘 수행하고있는 것을 알고 있습니다)에 감소 함수를 작성하려고합니다. 배열을 만들고 연산자를 사용하여 배열을 줄입니다.

나는 어떻게 함수를 전달하는 데 어려움을 겪고있다. hostOp()를 잘 작성한 호스트 기반 예제로 작성했습니다.

deviceOp()는 fminf()를 명시 적으로 호출하여 첫 ​​번째 명령문에서 작동하지만 함수 매개 변수를 호출 할 때 잘못된 메모리 액세스 오류가 발생합니다.

#include <iostream> 
#include <cstdio> 
#include <cmath> 
using namespace std; //for brevity 

__device__ float g_d_a = 9, g_d_b = 5; 
float g_h_a = 9, g_h_b = 5; 

template<typename argT, typename funcT> 
__global__ 
void deviceOp(funcT op){  
    argT result = fminf(g_d_a, g_d_b);     //works fine 
    printf("static function result: %f\n", result); 
    result = op(g_d_a,g_d_b);       //illegal memory access 
    printf("template function result: %f\n", result); 
} 

template<typename argT, typename funcT>     
void hostOp(funcT op){ 
    argT result = op(g_h_a, g_h_b); 
    printf("template function result: %f\n", result); 
} 

int main(int argc, char* argv[]){ 
    hostOp<float>(min<float>);       //works fine 
    deviceOp<float><<<1,1>>>(fminf); 

    cudaDeviceSynchronize(); 
    cout<<cudaGetErrorString(cudaGetLastError())<<endl; 
} 

출력 :

host function result: 5.000000 
static function result: 5.000000 
an illegal memory access was encountered 

불법적 인 메모리 액세스가되지 않도록 어떻게 deviceOp에 fminf을 통과해야한다, 내가 끔찍하게 바보 같은 일을하고 있지 않다 가정?

내가 끔찍한 짓을하고 있다면 더 좋은 방법이 무엇일까요?

Answer 1

장치에서 호출 할 함수는 __device__ (또는 커널로 사용하려면 __global__)으로 장식해야합니다. nvcc 컴파일러 드라이버는 호스트와 장치 코드를 분리하고, 장치 코드 (즉, 컴파일 된 장치 코드)에서 호출 될 때 장치 컴파일 된 버전의 함수를 사용하고 그렇지 않으면 호스트 버전을 사용합니다.

이 구조는 문제가있다 : 그것은 명확하지 않을 수 있지만

deviceOp<float><<<1,1>>>(fminf); 

, 이것은 기본적으로 모든 호스트 코드입니다. 예, 커널을 시작합니다 (호스트 코드에서 라이브러리 호출의 기본 시퀀스를 통해). 그러나 기술적으로 호스트 코드입니다. 따라서 fminf 함수 주소는 여기에 fminf 버전의 호스트 버전입니다 (실제로 포함되지 않은 CUDA math.h을 통해 사용할 수 있음).

일반적으로이 문제를 해결하기위한 어색한 방법은 장치 주소 을 장치 코드에 "캡처"한 다음 매개 변수로 커널에 전달하는 것입니다.

컴파일 타임에 추론 할 수있는 함수 주소를 약간 다른 템플릿 기술을 사용하여 전달하는 경우이 프로세스를 (약간) 회로 단락시킬 수도 있습니다. 이러한 개념은 this answer에서 다룹니다.

$ cat t1176.cu 
#include <iostream> 
#include <cstdio> 
#include <cmath> 
using namespace std; //for brevity 

__device__ float g_d_a = 9, g_d_b = 5; 
float g_h_a = 9, g_h_b = 5; 

template<typename argT, typename funcT> 
__global__ 
void deviceOp(funcT op){ 
    argT result = fminf(g_d_a, g_d_b);     //works fine 
    printf("static function result: %f\n", result); 
    result = op(g_d_a,g_d_b);       //illegal memory access 
    printf("template function result: %f\n", result); 
} 

__device__ float (*my_fminf)(float, float) = fminf; // "capture" device function address 

template<typename argT, typename funcT> 
void hostOp(funcT op){ 
    argT result = op(g_h_a, g_h_b); 
    printf("template function result: %f\n", result); 
} 

int main(int argc, char* argv[]){ 
    hostOp<float>(min<float>);       //works fine 
    float (*h_fminf)(float, float); 
    cudaMemcpyFromSymbol(&h_fminf, my_fminf, sizeof(void *)); 
    deviceOp<float><<<1,1>>>(h_fminf); 

    cudaDeviceSynchronize(); 
    cout<<cudaGetErrorString(cudaGetLastError())<<endl; 
} 
$ nvcc -o t1176 t1176.cu 
$ cuda-memcheck ./t1176 
========= CUDA-MEMCHECK 
template function result: 5.000000 
static function result: 5.000000 
template function result: 5.000000 
no error 
========= ERROR SUMMARY: 0 errors 
$