Le lancement du noyau CUDA échoue lors de l'utilisation de divers décalages dans les données d'entrée.

Question

Mon code donne un message d'erreur et j'essaie d'en trouver la cause. Pour faciliter la recherche du problème, j'ai supprimé le code qui n'a apparemment aucun rapport avec le message d'erreur. Si vous pouvez me dire pourquoi le code simple suivant produit un message d'erreur, je pense que je devrais être en mesure de corriger mon code original :

#include "cuComplex.h"
#include <cutil.h>

__device__ void compute_energy(void *data, int isample, int nsamples) {
  cuDoubleComplex * const nminusarray          = (cuDoubleComplex*)data;
  cuDoubleComplex * const f                    = (cuDoubleComplex*)(nminusarray+101);
  double          * const abs_est_errorrow_all = (double*)(f+3);
  double          * const rel_est_errorrow_all = (double*)(abs_est_errorrow_all+nsamples*51);
  int             * const iid_all              = (int*)(rel_est_errorrow_all+nsamples*51);
  int             * const iiu_all              = (int*)(iid_all+nsamples*21);
  int             * const piv_all              = (int*)(iiu_all+nsamples*21);
  cuDoubleComplex * const energyrow_all        = (cuDoubleComplex*)(piv_all+nsamples*12);
  cuDoubleComplex * const refinedenergyrow_all = (cuDoubleComplex*)(energyrow_all+nsamples*51);
  cuDoubleComplex * const btplus_all           = (cuDoubleComplex*)(refinedenergyrow_all+nsamples*51);

  cuDoubleComplex * const btplus           = btplus_all+isample*21021;

  btplus[0] = make_cuDoubleComplex(0.0, 0.0);
}

__global__ void computeLamHeight(void *data, int nlambda) {
  compute_energy(data, blockIdx.x, nlambda);
}

int main(int argc, char *argv[]) {
  void *device_data;

  CUT_DEVICE_INIT(argc, argv);
  CUDA_SAFE_CALL(cudaMalloc(&device_data, 184465640));
  computeLamHeight<<<dim3(101, 1, 1), dim3(512, 1, 1), 45000>>>(device_data, 101);
  CUDA_SAFE_CALL(cudaThreadSynchronize());
}

J'utilise une GeForce GTX 480 et je compile le code comme suit :

nvcc -L /soft/cuda-sdk/4.0.17/C/lib -I /soft/cuda-sdk/4.0.17/C/common/inc -lcutil_x86_64 -arch sm_13 -O3 -Xopencc "-Wall" Main.cu

La sortie est :

Using device 0: GeForce GTX 480
Cuda error in file 'Main.cu' in line 31 : unspecified launch failure.

EDIT : J'ai maintenant simplifié davantage le code. Le code simplifié suivant produit toujours le message d'erreur :

#include <cutil.h>

__global__ void compute_energy(void *data) {
  *(double*)((int*)data+101) = 0.0;
}

int main(int argc, char *argv[]) {
  void *device_data;

  CUT_DEVICE_INIT(argc, argv);
  CUDA_SAFE_CALL(cudaMalloc(&device_data, 101*sizeof(int)+sizeof(double)));
  compute_energy<<<dim3(1, 1, 1), dim3(1, 1, 1)>>>(device_data);
  CUDA_SAFE_CALL(cudaThreadSynchronize());
}

Il est maintenant facile de voir que le décalage devrait être valide. J'ai essayé d'exécuter cuda-memcheck et il dit ce qui suit :

========= CUDA-MEMCHECK
Using device 0: GeForce GTX 480
Cuda error in file 'Main.cu' in line 13 : unspecified launch failure.
========= Invalid __global__ write of size 8
=========     at 0x00000020 in compute_energy
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0x200200194 is misaligned
=========
========= ERROR SUMMARY: 1 error

J'ai essayé de faire des recherches sur Internet pour trouver ce que signifie le fait que l'adresse soit mal alignée, mais je n'ai pas trouvé d'explication. Quel est le problème ?

Answer 1

Il était très difficile d'analyser votre code original avec toutes ces constantes magiques, mais votre cas de reproduction mis à jour rend le problème immédiatement évident. L'architecture du GPU exige que tous les pointeurs soient alignés sur les limites des mots. Votre noyau contient un accès au pointeur qui n'est pas correctement aligné sur les mots. Les doubles sont un type de 64 bits, et votre adressage n'est pas aligné sur une frontière de 64 bits. Ceci :

*(double*)((int*)data+100) = 0.0; // 50th double

ou ceci :

*(double*)((int*)data+102) = 0.0; // 51st double

sont toutes deux légales. Ceci :

*(double*)((int*)data+101) = 0.0; // not aligned to a 64 bit boundary

ne l'est pas.

Answer 2

L'erreur indique un accès à la mémoire hors limites, veuillez vérifier la valeur de l'offset.