NaN aléatoire et résultats incorrects avec le noyau OpenCL

Question

J'essaie d'implémenter un noyau OpenCL général de multiplication de matrices, qui soit conforme aux normes suivantes C = α*A*B + β*C .

Le noyau

J'ai fait quelques recherches en ligne et j'ai décidé d'utiliser un noyau modifié de ce site web comme point de départ. La principale modification que j'ai apportée est que l'allocation de la mémoire locale comme espace de travail est désormais dynamique. Voici le noyau que j'ai écrit :

__kernel
void clkernel_gemm(const uint M, const uint N, const uint K, const float alpha,
                   __global const float* A, __global const float* B, const float beta, 
                   __global float* C, __local float* Asub, __local float* Bsub) {

  const uint row = get_local_id(0);
  const uint col = get_local_id(1);
  const uint TS = get_local_size(0); // Tile size
  const uint globalRow = TS * get_group_id(0) + row; // Row ID of C (0..M)
  const uint globalCol = TS * get_group_id(1) + col; // Row ID of C (0..N)

  // Initialise the accumulation register
  float acc = 0.0f;

  // Loop over all tiles
  const int numtiles = K / TS;
  for (int t = 0; t < numtiles; t++) {
    const int tiledRow = TS * t + row;
    const int tiledCol = TS * t + col;
    Asub[col * TS + row] = A[tiledCol * M + globalRow];
    Bsub[col * TS + row] = B[globalCol * K + tiledRow];

    barrier(CLK_LOCAL_MEM_FENCE);

    for(int k = 0; k < TS; k++) {
      acc += Asub[k * TS + row] * Bsub[col * TS + k] * alpha;
    }

    barrier(CLK_LOCAL_MEM_FENCE);
  }

  C[globalCol * M + globalRow] = fma(beta, C[globalCol * M + globalRow], acc);
}

La taille des carreaux (TS) est désormais une valeur définie dans le code d'appel, qui se présente comme suit :

  // A, B and C are 2D matrices, their cl::Buffers have already been set up
  // and values appropriately set.

  kernel.setArg(0, (cl_int)nrowA);
  kernel.setArg(1, (cl_int)ncolB);
  kernel.setArg(2, (cl_int)ncolA);
  kernel.setArg(3, alpha);
  kernel.setArg(4, A_buffer);
  kernel.setArg(5, B_buffer);
  kernel.setArg(6, beta);
  kernel.setArg(7, C_buffer);
  kernel.setArg(8, cl::Local(sizeof(float) * nrowA * ncolB));
  kernel.setArg(9, cl::Local(sizeof(float) * nrowA * ncolB));

  cl::NDRange global(nrowA, ncolB);
  cl::NDRange local(nrowA, ncolB);

  status = cmdq.enqueueNDRangeKernel(kernel, cl::NDRange(0), global, local);

Le problème

Le problème que je rencontre est que les tests unitaires (écrits avec gtest de Google) que j'ai écrits échouent de manière aléatoire, mais seulement pour ce noyau particulier. (J'ai 20 autres noyaux dans les mêmes .cl qui réussissent les tests 100% du temps)

J'ai un test qui multiplie une matrice flottante 1x4 {0.0, 1.0, 2.0, 3.0} avec une version transposée de lui-même {{0.0}, {1.0}, {2.0}, {3.0}} . Le résultat attendu est {14.0} .

Cependant, je ne parviens à obtenir ce résultat correct que dans 75 % des cas.

Parfois, j'obtiens 23.0 (GTX 970), 17.01 (GTX 750) ou tout simplement -nan et 0,0 (pour les trois appareils). Ce qui est curieux, c'est que les résultats incorrects respectifs semblent être propres aux appareils ; je n'arrive pas, par exemple, à obtenir 23.0 sur le CPU Intel ou la GTX 750.

Je suis déconcerté car si j'ai commis une erreur algorithmique ou mathématique, cette erreur devrait être cohérente ; au lieu de cela, je n'obtiens des résultats incorrects que de manière aléatoire.

Qu'est-ce que je fais de travers ?

Ce que j'ai essayé

• J'ai vérifié que les données introduites dans les noyaux sont correctes.
• J'ai essayé d'initialiser les deux __local à 0.0, mais cela rend tous les résultats erronés (mais franchement, je ne sais pas vraiment comment l'initialiser correctement).
• J'ai écrit un programme de test qui n'exécute que ce noyau afin d'exclure toute condition de course interagissant avec le reste de mon programme, mais le bogue se produit toujours.

Autres points à noter

• J'utilise le wrapper C++ récupéré directement à partir de la page Github .
• Pour utiliser le wrapper, j'ai défini CL_HPP_MINIMUM_OPENCL_VERSION 120 y CL_HPP_TARGET_OPENCL_VERSION 120 .
• Je compile les noyaux avec l'option -cl-std=CL1.2 drapeau.
• Tous cl::Buffer sont créés à l'aide de seulement les CL_MEM_READ_WRITE drapeau.
• Je teste cela sur Ubuntu 16.04, Ubuntu 14.04 et Debian 8.
• J'ai testé cette méthode sur des processeurs Intel avec la fonction Intel OpenCL Runtime 16.1 pour Ubuntu installé. Le moteur d'exécution indique qu'il prend en charge jusqu'à OpenCL 1.2.
• Je l'ai testé sur les Nvidia GTX 760 et 970. Nvidia ne supporte que l'OpenCL 1.2.
• Les trois plates-formes présentent le même problème à une fréquence variable.

Answer 1

Cela semble compliqué. Il y a plusieurs choses à aborder et elles ne tiendront pas dans les commentaires, donc je vais poster tout cela comme une réponse même si cela ne résout pas (encore) votre problème.

---

Je suis perplexe, car si j'ai fait un calcul algorithmique ou mathématique, je n'ai pas le temps de le faire. mathématique, l'erreur devrait être cohérente ; au lieu de cela, j'obtiens des des résultats incorrects de manière aléatoire.

Un tel comportement est un indicateur typique de conditions de course.

---

J'ai essayé d'initialiser les deux mémoires __locales à 0.0, mais cela a pour effet de fausser tous les résultats (mais franchement, je ne sais pas trop comment faire). tous les résultats sont erronés (mais franchement, je ne sais pas vraiment comment l'initialiser correctement)

En fait, c'est une bonne chose. Nous avons enfin une certaine cohérence.

---

Initialisation de la mémoire locale

L'initialisation de la mémoire locale peut être effectuée en utilisant les éléments de travail, par exemple si vous avez un groupe de travail 1D de 16 éléments et que votre mémoire locale consiste en 16 flottants, il suffit de faire ceci :

local float* ptr = ...          // your pointer to local memory
int idx = get_local_id(0);      // get the index for the current work-item
ptr[idx] = 0.f;                 // init with value 0
barrier(CLK_LOCAL_MEM_FENCE);   // synchronize local memory access within workgroup

Si votre mémoire locale est plus grande, par exemple 64 flottants, vous devrez utiliser une boucle où chaque élément de travail initialise 4 valeurs, du moins c'est la façon la plus efficace. Cependant, personne ne vous empêchera d'utiliser chaque élément de travail pour initialiser chaque valeur de la mémoire locale, même si cela n'a aucun sens puisque vous l'initialisez en fait plusieurs fois.

---

Vos changements

En algorithme original semble avoir été spécialement conçu pour utiliser des tuiles quadratiques.

__local float Asub[TS][TS];
__local float Bsub[TS][TS];

De plus, la taille de la mémoire locale correspond à la taille du groupe de travail, dans leur exemple 32x32. Lorsque je regarde les paramètres de votre noyau pour la mémoire locale, je vois que vous utilisez des paramètres qui sont définis comme M et N dans l'algorithme original. Cela ne semble pas correct.

Mise à jour 1

Puisque vous n'avez pas indiqué si l'algorithme original fonctionne pour vous, voici ce que vous devez faire pour trouver votre erreur :

• Créez un ensemble de données de test. Veillez à n'utiliser que des tailles de données réellement prises en charge par l'algorithme original (par exemple, la taille minimale, les multiples de x, etc.) Utilisez également de grands ensembles de données, car certaines erreurs n'apparaissent que si plusieurs groupes de travail sont envoyés.
• Utilisez l'algorithme original, non modifié, avec vos ensembles de données de test et vérifiez les résultats.
• Modifiez l'algorithme de manière à utiliser une mémoire locale dynamique au lieu d'une mémoire locale de taille fixe, mais assurez-vous qu'elle a la même taille que l'approche de taille fixe. C'est ce que vous avez essayé, mais je pense que cela a échoué en raison de ce que j'ai décrit dans la section "Vos modifications".