Quelle est la granularité du CUDA Memory Checker?

J'ai un noyau qui fonctionne dans les versions de débogage mais échoue dans les versions de version. Je suspecte que je lis ou écris hors des limites, mais le vérificateur de mémoire de CUDA ne montre aucune erreur. Alors, je l'ai fait un test où je lecture et d'écriture hors des limites avec le noyau suivant:Quelle est la granularité du CUDA Memory Checker?

__global__ void addKernel(int *c, const int *a, const int *b) 
{ 
    int x[1]; 
    for (int i(0); i < 100; ++i) { 
    x[i] = i; 
    } 
    int t(0); 
    for (int i(0); i < 100; ++i) { 
    t += x[i]; 
    } 
    c[0] = t; 
}

Le CUDA Memory Checker n'a pas pris le hors d'écritures hors limites et lit. J'ai été capable d'obtenir une écriture hors limites pour être détectée en augmentant la boucle à 1.000.000 rounds, ce qui a amené le noyau à faire une pause à i = 502.586 (2MB hors limites).

Est-ce la granularité attendue au cours de laquelle le vérificateur mémoire CUDA fonctionne et est-ce que je peux faire pour l'obtenir pour détecter de petites hors des limites écrit (de l'ordre de quelques octets?)

Source

2013-02-15 Roger Dahl

Peut-être que le compilateur fait une de ses optimisations incroyables? Le résultat est sommatif de i de 0 à 100. Que se passe-t-il, pass 100 est un paramètre du noyau (inconnu au moment de la compilation)? – pQB

Semblait improbable. Je l'ai essayé de toute façon, et ça n'a rien changé :) –

Je pense que vous Nous sommes juste en train de nous faire trébucher dans cet exemple par l'optimisation, comme cela a déjà été suggéré.

Voici mon cas de test:

$ cat t1130.cu 
#include <stdio.h> 

__global__ void addKernel(int *c) 
{ 
    int x[1]; 
    for (int i(0); i < 100; ++i) { 
    x[i] = i; 
#ifdef FORCE 
    printf("%d ", i); 
#endif 
    } 
    int t(0); 
    for (int i(0); i < 100; ++i) { 
    t += x[i]; 
    } 
    c[0] = t; 
} 

int main(){ 

    int *d_c; 
    cudaMalloc(&d_c, sizeof(int)); 
    addKernel<<<1,1>>>(d_c); 
    cudaDeviceSynchronize(); 
} 
$ nvcc -o t1130 t1130.cu 
$ cuda-memcheck ./t1130 
========= CUDA-MEMCHECK 
========= ERROR SUMMARY: 0 errors 
$ nvcc -DFORCE -o t1130 t1130.cu 
$ cuda-memcheck ./t1130 
========= CUDA-MEMCHECK 
========= Invalid __local__ write of size 4 
=========  at 0x00000168 in addKernel(int*) 
=========  by thread (0,0,0) in block (0,0,0) 
=========  Address 0x00fffd10 is out of bounds 
=========  Saved host backtrace up to driver entry point at kernel launch time 
=========  Host Frame:/lib64/libcuda.so.1 (cuLaunchKernel + 0x2cd) [0x15865d] 
=========  Host Frame:./t1130 [0x16ca1] 
=========  Host Frame:./t1130 [0x314b3] 
=========  Host Frame:./t1130 [0x27a1] 
=========  Host Frame:./t1130 [0x269c] 
=========  Host Frame:./t1130 [0x26b6] 
=========  Host Frame:./t1130 [0x2600] 
=========  Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x21d65] 
=========  Host Frame:./t1130 [0x2489] 
========= 
0 1 2 3 ========= Program hit cudaErrorLaunchFailure (error 4) due to "unspecified launch failure" on CUDA API call to cudaDeviceSynchronize. 
=========  Saved host backtrace up to driver entry point at error 
=========  Host Frame:/lib64/libcuda.so.1 [0x2f31b3] 
=========  Host Frame:./t1130 [0x354a6] 
=========  Host Frame:./t1130 [0x2605] 
=========  Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x21d65] 
=========  Host Frame:./t1130 [0x2489] 
========= 
========= ERROR SUMMARY: 2 errors 
$

Donc, si nous compilons votre code en l'état, il fonctionne sans erreur sous cuda-memcheck. Mais si nous forçons le compilateur à générer la boucle réelle, l'accès invalide est détecté et signalé.

En expliquant également votre commentaire, si nous utilisons -G au lieu de -DFORCE dans le cas ci-dessus, cuda-memcheck détectera également et signalera l'erreur, puisque les optimisations sont désactivées.

Source

2016-04-01 21:27:34

Quelle est la granularité du CUDA Memory Checker?

Répondre

Questions connexes