Comment utiliser la mémoire constante CUDA dans un programmeur de manière agréable?

Question

Je travaille sur une application crunching en utilisant le framework CUDA. J'ai des données statiques qui doivent être accessibles à tous les threads, donc je l'ai mis en mémoire constante comme ceci:

__device__ __constant__ CaseParams deviceCaseParams; 

J'utilise le cudaMemcpyToSymbol d'appel pour transférer ces params de l'hôte vers l'appareil:

void copyMetaData(CaseParams* caseParams) 
{ 
    cudaMemcpyToSymbol("deviceCaseParams", caseParams, sizeof(CaseParams)); 
} 

qui fonctionne.

De toute façon, il semble (par essais et erreurs, et aussi de la lecture des messages sur le net) que pour une raison quelconque, la déclaration de deviceCaseParams et l'opération de copie (l'appel à cudaMemcpyToSymbol) doit être dans le même fichier. Pour le moment j'ai ces deux dans un fichier .cu, mais je veux vraiment avoir le paramètre struct dans un fichier .cuh afin que toute implémentation puisse le voir si elle le veut. Cela signifie que je dois aussi avoir la fonction copyMetaData dans un fichier d'en-tête, mais cela perturbe le lien (symbole déjà défini) puisque les deux fichiers .cpp et .cu incluent cet entête (et donc le compilateur MS C++ et nvcc le compile).

Quelqu'un a-t-il des conseils sur la conception ici?

Mise à jour: Voir les commentaires

Answer 1

Avec une CUDA mise à jour (par exemple 3.2), vous devriez être en mesure de faire la memcpy à partir d'une unité de traduction différente si vous cherchez le symbole à runtime (c'est-à-dire en passant une chaîne comme premier argument à cudaMemcpyToSymbol comme vous l'êtes dans votre exemple). En outre, avec les périphériques de classe Fermi, vous pouvez simplement déplacer la mémoire (cudaMalloc), la copier dans la mémoire de l'unité, puis passer l'argument sous la forme d'un pointeur const. Le compilateur reconnaîtra si vous accédez aux données uniformément à travers les déformations et si oui utilisera le cache constant. Voir le Guide de programmation CUDA pour plus d'informations. Note: vous devrez compiler avec -arch=sm_20.

Answer 2

Si vous utilisez CUDA pré-Fermi, vous avez déjà découvert que ce problème ne s'applique pas seulement à la mémoire constante, il s'applique à tout ce que vous voulez du côté CUDA. Les deux seules manières que j'ai trouvé autour de ce sont soit:

1. Ecrire tout CUDA dans un seul fichier (.cu), ou
2. Si vous devez sortir le code dans des fichiers séparés, vous limiter aux en-têtes que votre seul fichier .cu comprend alors.

Si vous avez besoin de partager du code entre CUDA et C/C++, ou si vous partagez du code commun entre projets, l'option 2 est le seul choix. Il semble très anormal pour commencer, mais il résout le problème. Vous continuez à structurer votre code, mais pas de la même manière que C. La surcharge principale est que chaque fois que vous faites une compilation, vous compilez tout. Le côté positif de ceci (ce qui est peut-être la raison pour laquelle cela fonctionne de cette façon) est que le compilateur CUDA a accès à tout le code source en une seule fois ce qui est bon pour l'optimisation.