Программа CUDA вызывает сбой драйвера nvidia

Моя программа CUDA для расчета Монте-Карло Пи вызывает сбой моего драйвера nvidia, когда я превышаю около 500 испытаний и 256 полных блоков. Кажется, это происходит в функции ядра monteCarlo. Любая помощь приветствуется.

#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>
#include <curand.h>
#include <curand_kernel.h>


#define NUM_THREAD 256
#define NUM_BLOCK 256



///////////////////////////////////////////////////////////////////////////////////////////
///////////////////////////////////////////////////////////////////////////////////////////

// Function to sum an array
__global__ void reduce0(float *g_odata) {
extern __shared__ int sdata[];

// each thread loads one element from global to shared mem
unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
sdata[tid] = g_odata[i];
__syncthreads();

// do reduction in shared mem
for (unsigned int s=1; s < blockDim.x; s *= 2) { // step = s x 2
    if (tid % (2*s) == 0) { // only threadIDs divisible by the step participate
        sdata[tid] += sdata[tid + s];
    }
    __syncthreads();
}

// write result for this block to global mem
if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}

///////////////////////////////////////////////////////////////////////////////////////////
///////////////////////////////////////////////////////////////////////////////////////////
__global__ void monteCarlo(float *g_odata, int  trials, curandState *states){
//  unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int incircle, k;
    float x, y, z;
    incircle = 0;

    curand_init(1234, i, 0, &states[i]);

    for(k = 0; k < trials; k++){
        x = curand_uniform(&states[i]);
        y = curand_uniform(&states[i]);
        z =(x*x + y*y);
        if (z <= 1.0f) incircle++;
    }
    __syncthreads();
    g_odata[i] = incircle;
}
///////////////////////////////////////////////////////////////////////////////////////////
///////////////////////////////////////////////////////////////////////////////////////////
int main() {

    float* solution = (float*)calloc(100, sizeof(float));
    float *sumDev, *sumHost, total;
    const char *error;
    int trials; 
    curandState *devStates;

    trials = 500;
    total = trials*NUM_THREAD*NUM_BLOCK;

    dim3 dimGrid(NUM_BLOCK,1,1); // Grid dimensions
    dim3 dimBlock(NUM_THREAD,1,1); // Block dimensions
    size_t size = NUM_BLOCK*NUM_THREAD*sizeof(float); //Array memory size
    sumHost = (float*)calloc(NUM_BLOCK*NUM_THREAD, sizeof(float));

    cudaMalloc((void **) &sumDev, size); // Allocate array on device
    error = cudaGetErrorString(cudaGetLastError());
    printf("%s\n", error);


    cudaMalloc((void **) &devStates, (NUM_THREAD*NUM_BLOCK)*sizeof(curandState));
    error = cudaGetErrorString(cudaGetLastError());
    printf("%s\n", error);


    // Do calculation on device by calling CUDA kernel
    monteCarlo <<<dimGrid, dimBlock>>> (sumDev, trials, devStates);
    error = cudaGetErrorString(cudaGetLastError());
    printf("%s\n", error);

        // call reduction function to sum
    reduce0 <<<dimGrid, dimBlock, (NUM_THREAD*sizeof(float))>>> (sumDev);
    error = cudaGetErrorString(cudaGetLastError());
    printf("%s\n", error);

    dim3 dimGrid1(1,1,1);
    dim3 dimBlock1(256,1,1);
    reduce0 <<<dimGrid1, dimBlock1, (NUM_THREAD*sizeof(float))>>> (sumDev);
    error = cudaGetErrorString(cudaGetLastError());
    printf("%s\n", error);

    // Retrieve result from device and store it in host array
    cudaMemcpy(sumHost, sumDev, sizeof(float), cudaMemcpyDeviceToHost);
    error = cudaGetErrorString(cudaGetLastError());
    printf("%s\n", error);


    *solution = 4*(sumHost[0]/total);
    printf("%.*f\n", 1000, *solution);
    free (solution);
    free(sumHost);
    cudaFree(sumDev);
    cudaFree(devStates);
    //*solution = NULL;
    return 0;
}

person zetatr    schedule 31.05.2011    source источник


Ответы (2)


Если меньшее количество испытаний работает правильно, и если вы работаете в MS Windows без драйвера NVIDIA Tesla Compute Cluster (TCC) и/или используемый вами графический процессор подключен к дисплею, то вы, вероятно, превышаете «сторожевой таймер» операционной системы. "тайм-аут. Если ядро ​​слишком долго занимает устройство отображения (или любой графический процессор в Windows без TCC), ОС уничтожит ядро, чтобы система не стала неинтерактивной.

Решение состоит в том, чтобы работать на графическом процессоре, не подключенном к дисплею, и, если вы работаете в Windows, используйте драйвер TCC. В противном случае вам нужно будет уменьшить количество испытаний в вашем ядре и запускать ядро ​​несколько раз, чтобы вычислить необходимое количество испытаний.

РЕДАКТИРОВАТЬ: Согласно документам CUDA 4.0 curand (стр. 15, «Примечания к производительности»), вы можете повысить производительность, скопировав состояние генератора в локальное хранилище внутри вашего ядра, а затем сохранив состояние обратно (если оно вам понадобится снова), когда вы закончите:

curandState state = states[i];

for(k = 0; k < trials; k++){
    x = curand_uniform(&state);
    y = curand_uniform(&state);
    z =(x*x + y*y);
    if (z <= 1.0f) incircle++;
}

Далее упоминается, что установка стоит дорого, и предлагается переместить curand_init в отдельное ядро. Это может помочь снизить стоимость вашего ядра MC, чтобы вы не столкнулись со сторожевым таймером.

Я рекомендую прочитать этот раздел документов, там есть несколько полезных рекомендаций.

person harrism    schedule 31.05.2011
comment
Я запускаю Windows с моим графическим процессором, подключенным к дисплею. Я все еще удивлен, что ядро ​​​​заняло так много времени. Могут ли быть причиной вызовы curand_init и curand_uniform? - person zetatr; 31.05.2011
comment
Должно быть легко узнать - замените вызовы curand_uniform на 1.0f и закомментируйте curand_init. Кстати, вам не нужен этот __syncthreads(). - person harrism; 31.05.2011
comment
Спасибо, что сообщили мне о синхронизации. Кроме того, curand_uniform, по-видимому, заставляет ядро ​​работать значительно дольше. Это тоже позор, так как я даже не получаю хорошей сходимости с текущим количеством испытаний. Запуск большего количества ядер позволил бы мне повысить точность, но программа заняла бы гораздо больше времени для неудовлетворительного количества правильных цифр. - person zetatr; 31.05.2011
comment
Я добавил несколько советов по производительности из документации в свой ответ - держу пари, вы можете сократить время, это не должно быть дорогим ядром - curand_uniform всего несколько флопов, и если вы сохраните состояние в локальной переменной, он будет храниться в реестре. Я предполагаю, что настоящие расходы связаны с curand_init(), который компилятор, вероятно, исключил из мертвого кода, когда вы закомментировали curand_uniform(), из-за чего показалось, что curand_uniform был дорогим. Переместите curand_init в отдельное ядро ​​и состояние в локальную переменную, и вам должно быть намного лучше. Вы можете захотеть отдельное состояние для x и y, хотя... - person harrism; 31.05.2011
comment
Спасибо! Эти советы очень помогли. Помещение curand_init в отдельное ядро ​​позволило мне увеличить количество испытаний на пару порядков. Также я создаю отдельный массив состояний для y с другим начальным значением и вызовом curand_init. Это немного увеличило время выполнения, но дало мне как минимум 1 дополнительную цифру по сравнению с тем, что было раньше. Хотя кажется, что этот метод Монте-Карло все еще очень медленно сходится, поскольку у меня есть только 4 правильные цифры с более чем 1,3 миллиардами испытаний. - person zetatr; 31.05.2011
comment
Я не эксперт по MC, но экспериментировали ли вы с разными генераторами, такими как один из генераторов Quadirandom? Если вы нашли мой ответ полезным, пожалуйста, примите его. - person harrism; 31.05.2011
comment
Я просто хотел бы добавить, что при работе с Linux, в дополнение к работе на графическом процессоре, не подключенном к дисплею, X-сервер также должен быть выключен. В моем случае я использовал lightdm в Ubuntu, поэтому понадобилась следующая команда: sudo service lightdm stop - person Adam27X; 23.01.2014
comment
Я не считаю, что отключение X-сервера является обязательным требованием, если вы используете CUDA на графическом процессоре без дисплея (например, Tesla). - person harrism; 24.01.2014

Для тех из вас, у кого есть графический процессор geforce, который не поддерживает драйвер TCC, есть другое решение, основанное на:

http://msdn.microsoft.com/en-us/library/windows/hardware/ff569918(v=vs.85).aspx

  1. запустить regedit,
  2. перейдите к HKEY_LOCAL_MACHINE\System\CurrentControlSet\Control\GraphicsDrivers
  3. создайте новый ключ DWORD с именем TdrLevel, установите значение 0,
  4. перезагрузить ПК.

Теперь ваши долго работающие ядра не должны прерываться. Этот ответ основан на:

Изменение реестра для увеличения времени ожидания графического процессора, Windows 7

Я просто подумал, что было бы полезно предоставить решение и здесь.

person Michal Hosala    schedule 23.06.2014
comment
Зависает ли система/графика, если к этому графическому процессору подключен дисплей? - person Serge Rogatch; 19.08.2016
comment
@SergeRogatch да, полагаю. - person Michal Hosala; 23.08.2016