taille maximale prise en charge pour la bibliothèque cub

Question

Est-ce que quelqu'un sait quelle est la taille maximale prise en charge pour cub :: scan? J'ai reçu un vidage de base pour des tailles d'entrée supérieures à 500 millions. Je voulais vous assurer que je ne fais rien de mal ...

Voici mon code:

#define CUB_STDERR 
#include <stdio.h> 
#include "cub/util_allocator.cuh" 
#include "cub/device/device_scan.cuh" 
#include <sys/time.h> 
using namespace cub; 

bool     g_verbose = false; // Whether to display input/output to console 
CachingDeviceAllocator g_allocator(true); // Caching allocator for device memory 
typedef int mytype; 

/** 
* Solve inclusive-scan problem 
*/ 

static void solve(mytype *h_in, mytype *h_cpu, int n) 
{ 
    mytype inclusive = 0; 
    for (int i = 0; i < n; ++i) { 
     inclusive += h_in[i]; 
     h_cpu[i] = inclusive; 
    } 
} 
static int compare(mytype *h_cpu, mytype *h_o, int n) 
{ 
    for (int i = 0; i < n; i++) { 
     if (h_cpu[i] != h_o[i]) { 
     return i + 1; 
     } 
    } 
    return 0; 
} 

/** 
* Main 
*/ 
int main(int argc, char** argv) 
{ 
    cudaSetDevice(0); 
    struct timeval start, end; 
    int num_items = 1073741824; 
    const int repetitions = 5; 
    mytype *h_in, *h_out, *h_cpu; 
    const int size = num_items * sizeof(mytype); 
    // Allocate host arrays 
    h_in = (mytype *)malloc(size); 
    h_out = (mytype *)malloc(size); 
    h_cpu = (mytype *)malloc(size); 


    // Initialize problem and solution 
    for (int i = 0; i < num_items; i++) { 
     h_in[i] = i; 
     h_out[i] = 0; 
     h_cpu[i] = 0; 
    } 

    solve(h_in, h_cpu, num_items); 

    // Allocate problem device arrays 
    mytype *d_in = NULL; 
    CubDebugExit(g_allocator.DeviceAllocate((void**)&d_in, sizeof(mytype) * num_items)); 

    // Initialize device input 
    CubDebugExit(cudaMemcpy(d_in, h_in, sizeof(mytype) * num_items, cudaMemcpyHostToDevice)); 

    // Allocate device output array 
    mytype *d_out = NULL; 
    CubDebugExit(g_allocator.DeviceAllocate((void**)&d_out, sizeof(mytype) * num_items)); 

    // Allocate temporary storage 
    void   *d_temp_storage = NULL; 
    size_t   temp_storage_bytes = 0; 


    CubDebugExit(DeviceScan::InclusiveSum(d_temp_storage, temp_storage_bytes, d_in, d_out, num_items)); 
    CubDebugExit(g_allocator.DeviceAllocate(&d_temp_storage, temp_storage_bytes)); 

    // Run 
    gettimeofday(&start, NULL); 
    for (long i = 0; i < repetitions; i++) 
     DeviceScan::InclusiveSum(d_temp_storage, temp_storage_bytes, d_in, d_out, num_items); 
    cudaThreadSynchronize(); 
    gettimeofday(&end, NULL); 
    double ctime = end.tv_sec + end.tv_usec/1000000.0 - start.tv_sec - start.tv_usec/1000000.0; 

    cudaMemcpy(h_out, d_out, sizeof(mytype) * num_items, cudaMemcpyDeviceToHost); 
    int cmp = compare(h_cpu, h_out, num_items); 
    printf("%d\t", num_items); 
    if (!cmp) 
     printf("\t%7.4fs \n", ctime); 
    printf("\n"); 
    if (h_in) delete[] h_in; 
    if (h_out) delete[] h_out; 
    if (h_cpu) delete[] h_cpu; 
    if (d_in) CubDebugExit(g_allocator.DeviceFree(d_in)); 
    if (d_out) CubDebugExit(g_allocator.DeviceFree(d_out)); 
    if (d_temp_storage) CubDebugExit(g_allocator.DeviceFree(d_temp_storage)); 

    printf("\n\n"); 

    return 0; 
} 

Answer 1

Le problème est ici:

const int size = num_items * sizeof(mytype); 

Et il peut être fixé en changeant à:

const size_t size = num_items * sizeof(mytype); 

la valeur de num_items dans le code est plus de 1 milliard. Lorsque nous multiplions cela par sizeof(mytype), nous le multiplions par 4, de sorte que le résultat est supérieur à 4 milliards. Cette valeur ne peut pas être stockée dans une variable int. Si vous essayez de l'utiliser de toute façon, votre code hôte ultérieur fera de mauvaises choses. Ce problème (le core dump) n'a en fait rien à voir avec CUDA. Le code serait core dump si vous supprimiez tous les éléments CUB.

Lorsque je modifie la ligne de code ci-dessus, et compilent pour le GPU correct (par exemple -arch=sm_35 dans mon cas, ou -arch=sm_52 pour un GPU Titan X), puis-je obtenir la bonne réponse (et pas de défaut seg/core dump) .

En général, le point de départ correct lorsque la poursuite d'une erreur de type décharge de défaut seg/noyau, est de reconnaître que cette erreur provient de code hôte et vous devez tenter de localiser la ligne exacte du code source qui génère cette Erreur. Cela peut être fait trivialement/fastidieusement en mettant beaucoup d'instructions printf dans votre code, jusqu'à ce que vous identifiez la ligne de votre code après laquelle vous ne voyez aucune sortie printf, ou en utilisant un débogueur de code hôte, comme gdb sur linux. Notez également que ce code tel qu'il est écrit nécessitera un peu plus de 12 Go de mémoire sur l'hôte et un peu plus de 8 Go de mémoire sur le GPU, de sorte qu'il fonctionnera correctement dans de tels paramètres.

Pour référence, voici le code fixe (en fonction de ce OP affiché here):

#define CUB_STDERR 
#include <stdio.h> 
#include "cub/util_allocator.cuh" 
#include "cub/device/device_scan.cuh" 
#include <sys/time.h> 
using namespace cub; 

bool     g_verbose = false; // Whether to display input/output to console 
CachingDeviceAllocator g_allocator(true); // Caching allocator for device memory 
typedef int mytype; 

/** 
* Solve inclusive-scan problem 
*/ 

static void solve(mytype *h_in, mytype *h_cpu, int n) 
{ 
    mytype inclusive = 0; 
    for (int i = 0; i < n; ++i) { 
     inclusive += h_in[i]; 
     h_cpu[i] = inclusive; 
    } 
} 
static int compare(mytype *h_cpu, mytype *h_o, int n) 
{ 
    for (int i = 0; i < n; i++) { 
     if (h_cpu[i] != h_o[i]) { 
     return i + 1; 
     } 
    } 
    return 0; 
} 

/** 
* Main 
*/ 
int main(int argc, char** argv) 
{ 
    cudaSetDevice(0); 
    struct timeval start, end; 
    int num_items = 1073741824; 
    const int repetitions = 5; 
    mytype *h_in, *h_out, *h_cpu; 
    const size_t size = num_items * sizeof(mytype); 
    // Allocate host arrays 
    h_in = (mytype *)malloc(size); 
    h_out = (mytype *)malloc(size); 
    h_cpu = (mytype *)malloc(size); 


    // Initialize problem and solution 
    for (int i = 0; i < num_items; i++) { 
     h_in[i] = i; 
     h_out[i] = 0; 
     h_cpu[i] = 0; 
    } 

    solve(h_in, h_cpu, num_items); 

    // Allocate problem device arrays 
    mytype *d_in = NULL; 
    CubDebugExit(g_allocator.DeviceAllocate((void**)&d_in, sizeof(mytype) * num_items)); 

    // Initialize device input 
    CubDebugExit(cudaMemcpy(d_in, h_in, sizeof(mytype) * num_items, cudaMemcpyHostToDevice)); 

    // Allocate device output array 
    mytype *d_out = NULL; 
    CubDebugExit(g_allocator.DeviceAllocate((void**)&d_out, sizeof(mytype) * num_items)); 

    // Allocate temporary storage 
    void   *d_temp_storage = NULL; 
    size_t   temp_storage_bytes = 0; 


    CubDebugExit(DeviceScan::InclusiveSum(d_temp_storage, temp_storage_bytes, d_in, d_out, num_items)); 
    CubDebugExit(g_allocator.DeviceAllocate(&d_temp_storage, temp_storage_bytes)); 

    // Run 
    gettimeofday(&start, NULL); 
    for (long i = 0; i < repetitions; i++) 
     DeviceScan::InclusiveSum(d_temp_storage, temp_storage_bytes, d_in, d_out, num_items); 
    cudaThreadSynchronize(); 
    gettimeofday(&end, NULL); 
    double ctime = end.tv_sec + end.tv_usec/1000000.0 - start.tv_sec - start.tv_usec/1000000.0; 

    cudaMemcpy(h_out, d_out, sizeof(mytype) * num_items, cudaMemcpyDeviceToHost); 
    int cmp = compare(h_cpu, h_out, num_items); 
    printf("%d\t", num_items); 
    if (!cmp) 
     printf("\t%7.4fs \n", ctime); 
    printf("\n"); 
    if (h_in) delete[] h_in; 
    if (h_out) delete[] h_out; 
    if (h_cpu) delete[] h_cpu; 
    if (d_in) CubDebugExit(g_allocator.DeviceFree(d_in)); 
    if (d_out) CubDebugExit(g_allocator.DeviceFree(d_out)); 
    if (d_temp_storage) CubDebugExit(g_allocator.DeviceFree(d_temp_storage)); 

    printf("\n\n"); 

    return 0; 
}