Блог

У меня есть программа 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 amp;amp; 
                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 amp;amp; !isParallel amp;amp; 
                intersectionPoint.x > min(p1.x, p2.x) amp;amp; 
                intersectionPoint.x > min(p3.x, p4.x) amp;amp; 
                intersectionPoint.x < max(p1.x, p2.x) amp;amp; 
                intersectionPoint.x < max(p3.x, p4.x) amp;amp; 
                intersectionPoint.y > min(p1.y, p2.y) amp;amp; 
                intersectionPoint.y > min(p3.y, p4.y) amp;amp; 
                intersectionPoint.y < max(p1.y, p2.y) amp;amp; 
                intersectionPoint.y < max(p3.y, p4.y));
        }
    }

    output[segmentIndex] = resu<
}
  

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

 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-разрядная

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

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

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