Кажется, предел CUDA достигнут, но что это за предел? - PullRequest
Новая
       3

Кажется, предел CUDA достигнут, но что это за предел?

1 голос
/ 02 августа 2011

У меня есть программа CUDA, которая, кажется, выходит за пределы какого-то ресурса, но я не могу понять, что это за ресурс. Вот функция ядра: 

__global__ void DoCheck(float2* points, int* segmentToPolylineIndexMap, 
                        int segmentCount, int* output)
{
    int segmentIndex = threadIdx.x + blockIdx.x * blockDim.x;
    int pointCount = segmentCount + 1;

    if(segmentIndex >= segmentCount)
        return;

    int polylineIndex = segmentToPolylineIndexMap[segmentIndex];
    int result = 0;
    if(polylineIndex >= 0)
    {
        float2 p1 = points[segmentIndex];
        float2 p2 = points[segmentIndex+1];
        float2 A = p2;
        float2 a;
        a.x = p2.x - p1.x;
        a.y = p2.y - p1.y;

        for(int i = segmentIndex+2; i < segmentCount; i++)
        {
            int currentPolylineIndex = segmentToPolylineIndexMap[i];

            // if not a different segment within out polyline and
            // not a fake segment
            bool isLegit = (currentPolylineIndex != polylineIndex && 
                currentPolylineIndex >= 0);      

            float2 p3 = points[i];
            float2 p4 = points[i+1];
            float2 B = p4;
            float2 b;
            b.x = p4.x - p3.x;
            b.y = p4.y - p3.y;

            float2 c;
            c.x = B.x - A.x;
            c.y = B.y - A.y;

            float2 b_perp;
            b_perp.x = -b.y;
            b_perp.y = b.x;

            float numerator = dot(b_perp, c);
            float denominator = dot(b_perp, a);
            bool isParallel = (denominator == 0.0);

            float quotient = numerator / denominator;
            float2 intersectionPoint;
            intersectionPoint.x = quotient * a.x + A.x;
            intersectionPoint.y = quotient * a.y + A.y;

            result = result | (isLegit && !isParallel && 
                intersectionPoint.x > min(p1.x, p2.x) && 
                intersectionPoint.x > min(p3.x, p4.x) && 
                intersectionPoint.x < max(p1.x, p2.x) && 
                intersectionPoint.x < max(p3.x, p4.x) && 
                intersectionPoint.y > min(p1.y, p2.y) && 
                intersectionPoint.y > min(p3.y, p4.y) && 
                intersectionPoint.y < max(p1.y, p2.y) && 
                intersectionPoint.y < max(p3.y, p4.y));
        }
    }

    output[segmentIndex] = result;
}

Вот вызов для выполнения функции ядра: 

DoCheck<<<702, 32>>>(
    (float2*)devicePoints, 
    deviceSegmentsToPolylineIndexMap, 
    numSegments, 
    deviceOutput);

Размеры параметров следующие:

  • devicePoints = 22 464 float2s = 179 712 байт
  • deviceSegmentsToPolylineIndexMap = 22 463 дюймов = 89 852 байта
  • numSegments = 1 int = 4 байта
  • deviceOutput = 22 463 дюймов = 89 852 байта

Когда я запускаю это ядро, оно падает на видеокарту. Может показаться, что я нахожусь в каком-то пределе, потому что, если я запускаю ядро, используя DoCheck<<<300, 32>>>(...);, оно работает. Просто чтобы быть понятным, параметры одинаковы, просто количество блоков отличается.

Есть идеи, почему один вылетает из видео драйвера, а другой нет? Похоже, что тот, который потерпел неудачу, все еще в пределах лимита карты по количеству блоков.

Обновление Больше информации о конфигурации моей системы:

  • Видеокарта: nVidia 8800GT
  • CUDA Версия: 1.1
  • ОС: Windows Server 2008 R2

Я также попробовал это на ноутбуке со следующей конфигурацией, но получил те же результаты:

  • Видеокарта: nVidia Quadro FX 880M
  • CUDA Версия: 1.2
  • ОС: Windows 7, 64-разрядная

1 Ответ

6 голосов
/ 02 августа 2011

Ресурс, который истощается - время. На всех текущих платформах CUDA драйвер дисплея содержит сторожевой таймер, который убивает любое ядро, выполнение которого занимает более нескольких секунд. Запуск кода на карте, на которой запущен дисплей, подпадает под это ограничение.

На платформах Windows WDDM, которые вы используете, есть три возможных решения / обходных пути:

  1. Получите карту Telsa и используйте драйвер TCC, который полностью устраняет проблему
  2. Попробуйте изменить настройки реестра, чтобы увеличить лимит таймера (для получения дополнительной информации обратитесь к разделу реестра TdrDelay, но я не пользователь Windows и не могу быть более конкретным)
  3. Измените код ядра так, чтобы он был «входящим», и обрабатывайте параллельную загрузку данных в нескольких запусках ядра, а не в одном. Затраты на запуск ядра не так уж велики, и обработка рабочей нагрузки в течение нескольких запусков ядра часто довольно проста в зависимости от используемого вами алгоритма.
...