CUDA C++ Template de Paramètre de Noyau
Je suis en train de templatize un CUDA noyau basé sur une variable booléenne (comme illustré ici: Dois-je unifier les deux noyaux avec un "si", déclaration, au risque de perte de performance?), mais je reçois une erreur de compilateur qui dit que ma fonction n'est pas un modèle. Je pense que je suis juste en manque de quelque chose d'évident donc c'est assez frustrant.
Le code suivant ne fonctionne PAS:
util.cuh
#include "kernels.cuh"
//Utility functions
noyaux.cuh
#ifndef KERNELS
#define KERNELS
template<bool approx>
__global__ void kernel(...params...);
#endif
noyaux.cu
template<bool approx>
__global__ void kernel(...params...)
{
if(approx)
{
//Approximate calculation
}
else
{
//Exact calculation
}
}
template __global__ void kernel<false>(...params...); //Error occurs here
principal.cu
#include "kernels.cuh"
kernel<false><<<dimGrid,dimBlock>>>(...params...);
Le code suivant fonctionne:
util.cuh
#include "kernels.cuh"
//Utility functions
noyaux.cuh
#ifndef KERNELS
#define KERNELS
template<bool approx>
__global__ void kernel(...params...);
template<bool approx>
__global__ void kernel(...params...)
{
if(approx)
{
//Approximate calculation
}
else
{
//Exact calculation
}
}
#endif
principal.cu
#include "kernels.cuh"
kernel<false><<<dimGrid,dimBlock>>>(...params...);
Si je jette dans la
template __global__ void kernel<false>(...params...);
ligne à la fin de grains.cuh il fonctionne également.
J'obtiens les erreurs suivantes (les deux se référant à la ligne marquée ci-dessus):
kernel is not a template
invalid explicit instantiation declaration
Si cela fait une différence, je compile tout mon .cu fichiers en une seule ligne, comme:
nvcc -O3 -arch=sm_21 -I. main.cu kernels.cu -o program
OriginalL'auteur Adam27X | 2013-11-08
Vous devez vous connecter pour publier un commentaire.
Tous explicite de la spécialisation des déclarations doit être visible au moment de l'instanciation du modèle. Votre spécialisation explicite déclaration n'est visible que dans les grains.cu de l'unité de traduction, mais pas dans la main.cu.
Le code suivant est en effet de fonctionner correctement (à l'exception de l'ajout d'un
__global__
qualifier lors de l'instanciation explicite de l'instruction).MODIFIER
En C++, basé sur un modèle de fonctions ne sont pas compilés jusqu'à ce qu'une instanciation explicite de la fonction est rencontré. De ce point de vue, CUDA, qui supporte entièrement les modèles, se comporte exactement de la même manière qu'en C++.
Pour prendre un exemple concret, lorsque le compilateur trouve quelque chose comme
il vérifie juste la syntaxe de la fonction, mais ne produit pas de code objet. Ainsi, si vous compilez un fichier avec un simple basé sur un modèle de fonction comme ci-dessus, vous aurez un "vide", fichier de l'objet. C'est raisonnable, étant donné que le compilateur ne sais pas quel type d'assigner à
a
.Le compilateur génère un code objet uniquement lorsqu'il rencontre une instanciation explicite de la fonction de modèle. C'est à ce moment-là, combien de compilation basé sur un modèle de fonctions de travail et ce comportement introduit une restriction de fichiers multiples projets: la mise en œuvre (définition) basé sur un modèle de fonction doit être dans le même fichier que sa déclaration. Donc, vous ne pouvez pas séparer l'interface contenues dans
kernels.cuh
dans un fichier d'en-tête séparé dekernels.cu
, qui est la principale raison pour laquelle la première version de votre code ne compile pas. En conséquence, vous devez inclure à la fois l'interface et la mise en œuvre dans un fichier qui utilise les modèles, à savoir, vous devez inclure dansmain.cu
à la fois,kernels.cuh
etkernels.cu
.Depuis aucun code n'est généré sans une instanciation explicite, les compilateurs tolérer l'inclusion plus d'une fois le même fichier de modèle avec les deux déclarations et les définitions dans un projet sans générer d'erreurs de liaison.
Il y a plusieurs tutoriels sur l'utilisation de templates en C++. Un Idiot Guide pour les Modèles C++ - Partie 1, en dehors de l'irritant titre, vous donneront des instructions étape par étape introduction sur le sujet.
__global__
n'a pas de résoudre mon problème, en faisant cela et en mettant ma intégralité de la déclaration de la fonction danskernels.cuh
fait. Je préfère le garder danskernels.cu
si - est-ce possible?J'ai essayé d'expliquer pourquoi votre première version du code ne fonctionne pas dans mon édité réponse.
Juste pour être complet, vous pouvez remplacer
getch();
pourcudaDeviceReset();
à la place. Dans ce cas, ajouter#include <cuda_runtime.h>
OriginalL'auteur JackOLantern