function as templated parameter in cuda

Question

Я пытаюсь написать функцию сокращения в cuda (это упражнение, я знаю, что я делаю то, что было сделано другими людьми), которое принимает двоичный ассоциативный оператор и массив и уменьшает массив с помощью оператора.

У меня возникли трудности с передачей функции. Я написал 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 

Предполагая, что я не делаю что-то ужасно глупое, как я проходил мимо fminf в deviceOp так, что не является незаконным доступом к памяти?

Если я делаю что-то ужасно глупое, что лучше?

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 
$