argument configuration non valide pour fil bloquer supérieure à 16bit

Question

Ce code fonctionne très bien:

#include <stdio.h> 
#define N 1000 // <-- Works for values < 2^16 

__global__ 
void add(int *a, int *b) { 
    int i = blockIdx.x; 
    if (i<N) { 
     b[i] = 2*a[i]; 
    } 
} 
int main() { 
    int max_value[2]; 
    int ha[N], hb[N]; 
    int *da, *db; 
    cudaMalloc((void **)&da, N*sizeof(int)); 
    cudaMalloc((void **)&db, N*sizeof(int)); 
    for (int i = 0; i<N; ++i) { 
     ha[i] = i; 
    } 
    cudaMemcpy(da, ha, N*sizeof(int), cudaMemcpyHostToDevice); 
    add<<<N, 1>>>(da, db); 
    cudaMemcpy(hb, db, N*sizeof(int), cudaMemcpyDeviceToHost); 
    max_value[0] = hb[0]; 
    int i; 
    for (i = 0; i < N; i++) { 
     if (hb[i] > max_value[0]) { 
      max_value[0] = hb[i]; 
      max_value[1] = i; 
     } 
    } 
    cudaFree(da); 
    cudaFree(db); 
    printf("Max number %d, from value:%d \n", max_value[0], max_value[1]); 
    getchar(); 
    return 0; 
} 

Mais quand changer le numéro N (éléments du tableau) de 1000 à> (2) -1 le programme se bloque.

Je pensais qu'il était un débordement sur l'hôte, donc je me suis déplacé la déclaration de tableau de ha et hb-BSS segment et changé N-1000000.

#include <stdio.h> 
#define N 1000000 // <---- 

__global__ 
void add(int *a, int *b) { 
    int i = blockIdx.x; 
    if (i<N) { 
     b[i] = 2*a[i]; 
    } 
} 
static int ha[N]; // <---- 
static int hb[N]; // <---- 
int main() { 
    int max_value[2]; 
    // int ha[N], hb[N]; 
    int *da, *db; 
    cudaMalloc((void **)&da, N*sizeof(int)); 
    cudaMalloc((void **)&db, N*sizeof(int)); 
    for (int i = 0; i<N; ++i) { 
     ha[i] = i; 
    } 
    cudaMemcpy(da, ha, N*sizeof(int), cudaMemcpyHostToDevice); 
    add<<<N, 1>>>(da, db); 
    cudaMemcpy(hb, db, N*sizeof(int), cudaMemcpyDeviceToHost); 
    max_value[0] = hb[0]; 
    int i; 
    for (i = 0; i < N; i++) { 
     if (hb[i] > max_value[0]) { 
      max_value[0] = hb[i]; 
      max_value[1] = i; 
     } 
    } 
    cudaFree(da); 
    cudaFree(db); 
    printf("Max number %d, from value:%d \n", max_value[0], max_value[1]); 
    getchar(); 
    return 0; 
} 

Maintenant, je ne suis pas une erreur, mais le tableau hbest vide.
Quel est le problème avec mon code?
Comment puis-je allouer de grandes baies à un périphérique et obtenir un résultat valide?

JOUR: J'ai inséré le code de vérification d'erreur, l'erreur
je reçois est ->« arguments Configuration invalide ».
Le code mis à jour est le suivant:

#include <stdio.h> 
#include <time.h> 
#include <math.h> 
#include <thrust/system_error.h> 
#include <thrust/system/cuda/error.h> 
#include <sstream> 
const int N = 70000; 

#define checkCudaErrors(error) {\ 
    if (error != cudaSuccess) {\ 
     printf("CUDA Error - %s:%d: '%s'\n",__FILE__,__LINE__,cudaGetErrorString(error));\ 
     exit(1);\ 
     }\ 
}\ 

__global__ 
void add(int *a, int *b) { 
    int i = blockIdx.x; 
    if (i<N) { 
     b[i] = 2*a[i]; 
    } 
} 
static int ha[N]; 
static int hb[N]; 
int main() { 
    // int ha[N], hb[N]; 
    int max_value[2]; 

    int deviceCount = 0; 
    cudaGetDeviceCount(&deviceCount); 
    cudaError_t err=cudaDeviceReset(); 
    if(err!=cudaSuccess){printf("%s in %s at line %d\n",cudaGetErrorString(err),__FILE__,__LINE__);} 
    printf("Device count: %d \n", deviceCount); 

    for (int i = 0; i<N; ++i) { ha[i] = i; } 
    int *da, *db; 
    checkCudaErrors(cudaMalloc((void **)&da, N*sizeof(int))); 
    checkCudaErrors(cudaMalloc((void **)&db, N*sizeof(int))); 
    checkCudaErrors(cudaMemcpy(da, ha, N*sizeof(int), cudaMemcpyHostToDevice)); 
    add<<<N, 1>>>(da, db); // <--- Invalid configuration error 
    checkCudaErrors(cudaMemcpy(hb, db, N*sizeof(int), cudaMemcpyDeviceToHost)); 
    max_value[0] = hb[0]; 
    int i; 
    for (i = 0; i < N; i++) { 
     if (hb[i] > max_value[0]) { 
      max_value[0] = hb[i]; 
      max_value[1] = i; 
     } 
    } 
    cudaError_t error = cudaGetLastError();  
    if(error != cudaSuccess) { 
     printf("CUDA error: %s\n", cudaGetErrorString(error)); 
     getchar(); 
     exit(-1); 
    } 
    getchar(); 
    return 0; 
} 

Le dispositif est une GeForce GTX 470 et je suis en utilisant la compilation
nvcc de new.cu de foo

Answer 1

Votre appareil (GTX 470) est un utilitaire cc2.0 e (capacité de calcul). L'erreur d'argument de configuration non valide est due au fait que pour les périphériques cc2.0, le nombre de blocs pour une grille 1D est limité à 65535. Cette information est détectable dans le programming guide ("Dimension maximale x de une grille de blocs de thread ") ou en exécutant l'exemple de code CUDA deviceQuery. Donc, votre choix de N ici est trop grand:

add<<<N, 1>>>(da, db); 
    ^

La solution habituelle pour cela avec un dispositif de cc2.0 est de créer une grille de threadblocks qui est multidimensionnelle, ce qui permet un plus grand nombre de threadblocks. Les paramètres de lancement du noyau peuvent en fait être des variables dim3 qui permettent la spécification de grilles multidimensionnelles (de threads) ou de threads multidimensionnels (de threads).

Pour ce faire correctement, vous devrez également modifier votre code de noyau pour créer un ID de thread unique globalement approprié à partir des variables multidimensionnelles qui s'offrent à vous.

L'exemple chiffré suivant donne un ensemble minimal possible des changements pour démontrer le concept, et semble fonctionner correctement pour moi:

$ cat t363.cu 
#include <stdio.h> 
#include <time.h> 
#include <math.h> 
#include <thrust/system_error.h> 
#include <thrust/system/cuda/error.h> 
#include <sstream> 
const int N = 70000; 

#define checkCudaErrors(error) {\ 
    if (error != cudaSuccess) {\ 
     printf("CUDA Error - %s:%d: '%s'\n",__FILE__,__LINE__,cudaGetErrorString(error));\ 
     exit(1);\ 
     }\ 
}\ 

__global__ 
void add(int *a, int *b) { 
    int i = blockIdx.x + blockIdx.y*gridDim.x; 
    if (i<N) { 
     b[i] = 2*a[i]; 
    } 
} 
static int ha[N]; 
static int hb[N]; 
int main() { 
    int max_value[2]; 

    int deviceCount = 0; 
    cudaGetDeviceCount(&deviceCount); 
    cudaError_t err=cudaDeviceReset(); 
    if(err!=cudaSuccess){printf("%s in %s at line %d\n",cudaGetErrorString(err),__FILE__,__LINE__);} 
    printf("Device count: %d \n", deviceCount); 

    for (int i = 0; i<N; ++i) { ha[i] = i; } 
    int *da, *db; 
    checkCudaErrors(cudaMalloc((void **)&da, N*sizeof(int))); 
    checkCudaErrors(cudaMalloc((void **)&db, N*sizeof(int))); 
    checkCudaErrors(cudaMemcpy(da, ha, N*sizeof(int), cudaMemcpyHostToDevice)); 
    dim3 mygrid(N/10, 10); 
    add<<<mygrid, 1>>>(da, db); 
    checkCudaErrors(cudaMemcpy(hb, db, N*sizeof(int), cudaMemcpyDeviceToHost)); 
    max_value[0] = hb[0]; 
    int i; 
    for (i = 0; i < N; i++) { 
     if (hb[i] > max_value[0]) { 
      max_value[0] = hb[i]; 
      max_value[1] = i; 
     } 
    } 
    printf("max_value[0] = %d, max_value[1] = %d\n", max_value[0], max_value[1]); 
    cudaError_t error = cudaGetLastError(); 
    if(error != cudaSuccess) { 
     printf("CUDA error: %s\n", cudaGetErrorString(error)); 
     getchar(); 
     exit(-1); 
    } 
    return 0; 
} 
$ nvcc -arch=sm_20 -o t363 t363.cu 
nvcc warning : The 'compute_20', 'sm_20', and 'sm_21' architectures are deprecated, and may be removed in a future release (Use -Wno-deprecated-gpu-targets to suppress warning). 
$ ./t363 
Device count: 4 
max_value[0] = 139998, max_value[1] = 69999 
$ 

Notes:

Si vous avez exécuté votre code d'origine sur un CC3 .0 ou supérieur, il ne devrait pas lancer cette erreur. Les nouveaux appareils CUDA ont augmenté la limite de la grille 1D à 2^31-1. Mais si vous vouliez dépasser ce nombre de blocs (autour de 2B), vous deviez à nouveau passer à une grille multidimensionnelle. Les périphériques cc2.0 sont obsolètes dans CUDA 8 et leur prise en charge est supprimée de la prochaine version de CUDA 9.