Форум программистов, компьютерный форум, киберфорум
OpenCL
Войти
Регистрация
Восстановить пароль
Блоги Сообщество Поиск  
 
 
Рейтинг 4.73/11: Рейтинг темы: голосов - 11, средняя оценка - 4.73
194 / 29 / 5
Регистрация: 11.04.2015
Сообщений: 735

Непредсказуемые изменения точности расчётов внутри kernel

10.04.2024, 00:29. Показов 3735. Ответов 45
Метки нет (Все метки)

Студворк — интернет-сервис помощи студентам
Всем доброго времени суток!
Не так много времени прошло, но я снова вынужден обратиться за помощью к сообществу.

Не по теме:

Ну, тут правда то ли я тупой, то ли лыжи не едут. Я ДВА ДНЯ потратил на отладку, чтобы узнать, что проблема ни разу не в моих алгоритмах - просто OpenCL'ю иногда хочется посчитать 9 знаков после запятой, а не 15. Прошу прощения.


Смотрите, есть простая программа по вычислению длины отрезка дуги на поверхности сферы:
C
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
double distanceD(double point1[2], double point2[2])
{
    double R = 6371110.0;
    double phi1 = point1[0];
    double lambda1 = point1[1];
 
    double phi2 = point2[0];
    double lambda2 = point2[1];
 
    double delta_lambda = lambda2 - lambda1;
 
    double p1 = sin(delta_lambda) * cos(phi2);
    double p2 = cos(phi1) * sin(phi2) - sin(phi1) * cos(phi2) * cos(delta_lambda);
    double q = sin(phi1) * sin(phi2) + cos(phi1) * cos(phi2) * cos(delta_lambda);
    double res = abs(atan2(sqrt(p1 * p1 + p2 * p2), q) * R);
 
    return res;
}
Когда я проводил тесты OpenCL-порта моей крупной программы, я обнаружил, что в некоторых случаях значение функции distanceD, полученное на GPU отличается от CPU-варианта на 7 знаков после запятой (то есть, я получаю 9 корректных знаков после запятой, после чего идёт мусор). То есть, это точнее, чем float, но и далеко от double. В ходе некоторых экспериментов мне удалось заставить GPU выдать абсолютно точное значение, но воспользоваться этим "изобретением" мне не удаётся.
Поясняю. Вот код управляющей программы (Qt 5, CUDA 12.4, OpenCL 3.0 (auto)):
Кликните здесь для просмотра всего текста
C++ (Qt)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
#include <QCoreApplication>
#include <QDebug>
#include <QFile>
#include <QtMath>
#include <CL/cl.hpp>
 
double distanceD(double point1[2], double point2[2])
{
    double R = 6371110.0;
    double phi1 = point1[0];
    double lambda1 = point1[1];
 
    double phi2 = point2[0];
    double lambda2 = point2[1];
 
    double delta_lambda = lambda2 - lambda1;
 
    double p1 = sin(delta_lambda) * cos(phi2);
    double p2 = cos(phi1) * sin(phi2) - sin(phi1) * cos(phi2) * cos(delta_lambda);
    double q = sin(phi1) * sin(phi2) + cos(phi1) * cos(phi2) * cos(delta_lambda);
    double res = abs(atan2(sqrt(p1 * p1 + p2 * p2), q) * R);
 
    qDebug() << Qt::scientific << qSetRealNumberPrecision(18) << "CPU p2      =" << p2;
    qDebug() << Qt::scientific << qSetRealNumberPrecision(18) << "CPU res     =" << res;
 
    double p2_h1 = cos(phi1) * sin(phi2);
    double p2_h2 = sin(phi1) * cos(phi2) * cos(delta_lambda);
    double p2_sum = p2_h1 - p2_h2;
 
    qDebug() << Qt::scientific << qSetRealNumberPrecision(18) << "CPU p2_h1   =" << p2_h1 ;
    qDebug() << Qt::scientific << qSetRealNumberPrecision(18) << "CPU p2_h2   =" << p2_h2 ;
    qDebug() << Qt::scientific << qSetRealNumberPrecision(18) << "CPU p2_sum  =" << p2_sum;
 
    return res;
}
 
 
int main(int argc, char *argv[])
{
    std::vector<cl::Platform> all_platforms;
    cl::Platform::get(&all_platforms);
    if(all_platforms.size() == 0){
        qDebug() << "No platforms found. Check OpenCL installation.";
        return 1;
    }
    cl::Platform default_platform=all_platforms[0];
    qDebug() << "Using platform: " << QString::fromStdString(default_platform.getInfo<CL_PLATFORM_NAME>());
 
    std::vector<cl::Device> all_devices;
    default_platform.getDevices(CL_DEVICE_TYPE_ALL, &all_devices);
    if(all_devices.size() == 0){
        qDebug() << "No devices found. Check OpenCL installation or make sure at least one compatible GPU is connected.";
        return 1;
    }
    cl::Device default_device=all_devices[0];
    qDebug() << "Using device: " << QString::fromStdString(default_device.getInfo<CL_DEVICE_NAME>());
 
    qDebug() << "Using CL version: " << QString::fromStdString(default_device.getInfo<CL_DEVICE_VERSION>());
 
    cl::Context context(default_device);
 
    cl::Program::Sources sources;
 
    QFile kernelCode1(":/CalcDist.cl");
    kernelCode1.open(QFile::ReadOnly | QFile::Text);
    std::string codeString = kernelCode1.readAll().toStdString();
    unsigned long long codeStringSize = codeString.length();
 
    sources.push_back({codeString.c_str(), codeStringSize});
 
    cl::Program program(context,sources);
    if(program.build({default_device})!=CL_SUCCESS){
        qDebug() << " Error building: " << QString::fromStdString(program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(default_device));
        return 1;
    }
 
    cl::CommandQueue queue(context,default_device);
 
    cl::compatibility::make_kernel<> calcDist(cl::Kernel(program, "mainProgram"));
    cl::EnqueueArgs eargs(queue, cl::NullRange, cl::NDRange(1), cl::NullRange);
 
    calcDist(eargs).wait();
 
    double p1  [2] = {0.891949793450334649, 0.513485940430910115};
    double p2  [2] = {0.891949846176460226, 0.513485940430959964};
 
    double dist = distanceD(p1, p2);
    qDebug() << Qt::scientific << qSetRealNumberPrecision(18) << "CPU dist    =" << dist;
 
    return 0;
}

и код файла CalcDist.cl:
Кликните здесь для просмотра всего текста
C
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
double distanceD(double2 point1, double2 point2);
 
void __kernel mainProgram()
{
    long long index = get_global_id(0);
 
    if (index == 0)
    {
        double2 p1   = (double2)(0.891949793450334649, 0.513485940430910115);
        double2 p2   = (double2)(0.891949846176460226, 0.513485940430959964);
 
        double dist = distanceD(p1, p2);
 
        printf("NV dist     = %.18e\n", dist);
    }
}
 
 
double distanceD(double2 point1, double2 point2)
{
    double R = 6371110.0;
 
    double phi1 = point1[0];
    double lambda1 = point1[1];
 
    double phi2 = point2[0];
    double lambda2 = point2[1];
 
    double delta_lambda = lambda2 - lambda1;
 
    // переменные, устраняющие дублирование тригонометрических расчётов
    double cosphi1 = cos(phi1);
    double cosphi2 = cos(phi2);
    double sinphi1 = sin(phi1);
    double sinphi2 = sin(phi2);
    double cosdeltalambda = cos(delta_lambda);
    //
 
    double p1 = sin(delta_lambda) * cosphi2;
 
    // одношаговый расчёт
        double p2 = cosphi1 * sinphi2 - sinphi1 * cosphi2 * cosdeltalambda;
    // многошаговый расчёт
        // double p2_h1 = cosphi1 * sinphi2;
        // double p2_h2 = sinphi1 * cosphi2 * cosdeltalambda;
        // double p2 = p2_h1 - p2_h2;
    //
 
    // одношаговый расчёт
        double q = sinphi1 * sinphi2 + cosphi1 * cosphi2 * cosdeltalambda;
    // многошаговый расчёт
        // double q_h1 = sinphi1 * sinphi2;
        // double q_h2 = cosphi1 * cosphi2 * cosdeltalambda;
        // double q = q_h1 + q_h2;
    //
 
    // одношаговый расчёт
        double res = fabs(atan2(sqrt(p1 * p1 + p2 * p2), q) * R);
    // многошаговый расчёт
        // double res_spqr = sqrt(p1 * p1 + p2 * p2);
        // double res_atan2 = atan2(res_spqr, q);
        // double res_atan2_R = res_atan2  * R;
        // double res = fabs(res_atan2_R);
    //
 
    printf("NV p2       = %.18e\n", p2 );
    printf("NV res      = %.18e\n", res);
 
    double p2_h1 = cosphi1 * sinphi2;
    double p2_h2 = sinphi1 * cosphi2 * cosdeltalambda;
    double p2_sum = p2_h1 - p2_h2;
    double new_res = fabs(atan2(sqrt(p1 * p1 + p2_sum * p2_sum), q) * R);
 
    printf("NV p2_h1    = %.18e\n", p2_h1 );
    printf("NV p2_h2    = %.18e\n", p2_h2 );
    printf("NV p2_sum   = %.18e\n", p2_sum);
    printf("NV new_res  = %.18e\n", new_res);
 
    return res;
}

Если просто запустить эту программу и сравнить значения дистанции, то мы увидим
Code
1
2
CPU dist    = 3.359239459231472269e-01
NV dist     = 3.359239460197458449e-01
, что совпадают лишь первые 8 знаков.
Главная причина этого в том, что значение переменной p2 оказывается таким:
Code
1
2
CPU p2      = 5.272612557671862987e-08
NV p2       = 5.272612559188061227e-08
и здесь совпадают 9 знаков.
Однако, если мы после основного расчёта (внутри kernel) дистанции продублируем его, разделив расчёт переменной p2 на три отдельных операции, то, О ЧУДО, вывод окажется следующим:
Code
1
2
3
4
5
6
CPU p2_h1   = 4.886896713611054155e-01
CPU p2_h2   = 4.886896186349798388e-01
CPU p2_sum  = 5.272612557671862987e-08
NV p2_h1    = 4.886896713611054155e-01
NV p2_h2    = 4.886896186349798388e-01
NV p2_sum   = 5.272612557671862987e-08
для CPU ничего не изменилось (резонно), но что произошло на GPU???

Не по теме:

Какие-то приколы из мира Arduino, чесслово


Я думал, что нашёл проблему, и заменил старые одношаговые секции кода новыми - многошаговыми. КАК же я был удивлён, когда на выходе снова получил восемь несчастных корректных знаков после запятой из 15.
Отсюда и сам мой вопрос - а что, собственно, происходит? Почему, я даже не знаю, как это правильно назвать, видеокарта не хочет с первого раза нормально считать? Почему корректный расчёт происходит только при избыточном дублировании?
На всякий случай, прикладываю полный консольный вывод приведённой выше программы:
Code
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
Using platform:  "NVIDIA CUDA"
Using device:  "NVIDIA GeForce RTX 4080"
Using CL version:  "OpenCL 3.0 CUDA"
CPU p2      = 5.272612557671862987e-08
CPU res     = 3.359239459231472269e-01
CPU p2_h1   = 4.886896713611054155e-01
CPU p2_h2   = 4.886896186349798388e-01
CPU p2_sum  = 5.272612557671862987e-08
CPU dist    = 3.359239459231472269e-01
NV p2       = 5.272612559188061227e-08
NV res      = 3.359239460197458449e-01
NV p2_h1    = 4.886896713611054155e-01
NV p2_h2    = 4.886896186349798388e-01
NV p2_sum   = 5.272612557671862987e-08
NV new_res  = 3.359239459231471714e-01
NV dist     = 3.359239460197458449e-01
Заранее выражаю ОГРОМНУЮ благодарность за помощь!

Не по теме:

За помощь с перевоспитанием сумасбродной видеокарты, судя по всему

0
Programming
Эксперт
39485 / 9562 / 3019
Регистрация: 12.04.2006
Сообщений: 41,671
Блог
10.04.2024, 00:29
Ответы с готовыми решениями:

Повышение точности расчетов
Здрям! Подскажите, пожалуйста, как заставить эту заразу, во-первых, воспринимать отрицательные величины, во-вторых считать с точностью до...

Повышение точности расчетов в Matlab
Повышение точности расчетов в Matlab : http://www.advanpix.com/ Бесплатный Toolbox Приведу пример: N = 7; ...

Уменьшение точности расчетов (округление)
Доброго времени суток. Есть пример x=3/5000=0,0006 Нужно чтобы ответ был 0. Именно ответ, а не формат ответа. Чтобы маткад дальше...

45
 Аватар для snake32
3590 / 1720 / 236
Регистрация: 26.02.2009
Сообщений: 8,713
Записей в блоге: 5
23.04.2024, 14:09
Студворк — интернет-сервис помощи студентам
Цитата Сообщение от _Develop Посмотреть сообщение
а x64 это на расширениях AVX, SSE - тут только 64 bit
Тут кстати пишут что можно вызвать х87 ф-ии на х64. Но это надо на асме походу самому писать.

Цитата Сообщение от Ромуальд_7 Посмотреть сообщение
Я не понял, что такое IL, но, похоже, ядро компилируется именно через него.
Intermediate Language (IL) - промежуточный язык
https://openwall.info/wiki/john/development/AMD-IL
2
 Аватар для snake32
3590 / 1720 / 236
Регистрация: 26.02.2009
Сообщений: 8,713
Записей в блоге: 5
23.04.2024, 23:41
Цитата Сообщение от Ромуальд_7 Посмотреть сообщение
Итак Табличка (все расстояния приведены к метрам)
Не даёт мне покоя : У меня новые данные к табличке:
Code
1
2
3
4
Embarcadero® Delphi 11 Version 28.0.48361.3236
 
Win32    1.11125588355752e-01    1.38901044890629e-01
Win64    0.0                     1.34261104008121e-01
Пока все вычисления не сделал в 1 строку - точность хромала (как в прошлый раз)
Delphi
1
2
3
4
5
6
7
8
9
10
11
12
13
14
function Distance( const pt0,pt1:TDVec2 ):double;
    var cp,st:TDVec2;
        //v:double;
begin
  cp := radians( pt0 );
  st := radians( pt1 );
 
  //v := ;
 
  //clamp( v, 0.0, 1.0 );
  Result := arccos( sin( st.y )*sin( cp.y ) + cos( st.y )*cos( cp.y ) * cos( st.x-cp.x ) ) * Rearth;
  {if Result < epsDist then
    Result := 0.0;}
end;
0
194 / 29 / 5
Регистрация: 11.04.2015
Сообщений: 735
24.04.2024, 00:48  [ТС]
Цитата Сообщение от snake32 Посмотреть сообщение
Intermediate Language (IL) - промежуточный язык
О, спасибо! Но понять, о чём там речь, конечно, тяжело Тем не менее, всё указывает на то, что я запускаю kernel не так, поэтому почему мои опции компиляции игнорируются - всё ещё загадка.
Цитата Сообщение от snake32 Посмотреть сообщение
Не даёт мне покоя
А что именно не даёт покоя? С QGIS'ом, который использует GDAL для расчётов, спорить, я полагаю, бесполезно, а в остальном - наши данные сходятся в рамках использования различных формул для вычисления длины ортодромии. Я-то использую частный случай формулы Винсенти для эллипсоида, у которого большая и малая полуоси совпадают и равны среднему радиусу Земли; формулы брал отсюда. Это уникальный случай - википедия не только не наврала, но и стала-таки энциклопедией, когда других источников я в своё время не нашёл.
Кстати, интересно, что делфи либо применяет какие-то лишние оптимизации к варианту x64, либо влияние x87 действительно очень велико. Вариант, посчитанный для Win64, получается удивительно грубым для такой маленькой формулы.
Цитата Сообщение от snake32 Посмотреть сообщение
Пока все вычисления не сделал в 1 строку - точность хромала (как в прошлый раз)
А это снова какая-то мистика. На GPU получилось ровно наоборот. Мерзкие компиляторсы
0
 Аватар для snake32
3590 / 1720 / 236
Регистрация: 26.02.2009
Сообщений: 8,713
Записей в блоге: 5
24.04.2024, 11:40
Цитата Сообщение от Ромуальд_7 Посмотреть сообщение
А что именно не даёт покоя?
Более точная точность
Цитата Сообщение от Ромуальд_7 Посмотреть сообщение
С QGIS'ом, который использует GDAL для расчётов, спорить, я полагаю, бесполезно, а в остальном - наши данные сходятся
Да, похоже QGIS использует более точную(double-double?) математику + геоид с учётом разных радиусов от центра Земли у полюсов и экватора.
В текущей формуле сферы мы лишь можем подобрать чуть больший радиус на 55 параллели, например, я встречал такой радиус 6378137 м
0
827 / 244 / 47
Регистрация: 24.01.2013
Сообщений: 750
24.04.2024, 19:05
Короче не поленился и просчитал сам и кажется понял почему такой разброс результатов.

Сначала посчитал в MS Visual C++, результат как у snake32 = 0,13426110400812105.
Надо отдать должное компилятору - никакой разницы как записывать функцию нет, в одну строку или с промежуточными переменными, и даже debug/release на результат не влияют.

Потом взял обычный виндовый калькулятор и просчитал в нем (там вроде 32 цифры после запятой, но надо их проверить?).
Результат = 0,11111110745692463905484496187387.
Не знаю на сколько он точен, но в процессе ручного счета увидел что арккосинус берется от числа 0,99999999999999984792607679235729,
а оно отличается от единицы всего в 16-й цифре, т.е. для типа double это на грани фола.
И последующий разброс результатов уже зависит от компилятора и девайся, кто как это разрулит.

В общем для достижения большей точности нужен тип данных шире чем double иначе никак.
1
 Аватар для snake32
3590 / 1720 / 236
Регистрация: 26.02.2009
Сообщений: 8,713
Записей в блоге: 5
25.04.2024, 13:08
Нашёл длинную арифметику использующую несколько ячеек флоата/дабла под одну цифру
https://www.davidhbailey.com/dhbpapers/qd.pdf
Прочитал по диагонали... надо пробовать...

Добавлено через 3 минуты
https://github.com/scibuilder/... /qd_real.h
0
Надоела реклама? Зарегистрируйтесь и она исчезнет полностью.
inter-admin
Эксперт
29715 / 6470 / 2152
Регистрация: 06.03.2009
Сообщений: 28,500
Блог
25.04.2024, 13:08

Преобразовать функцию для повышения точности расчетов
Есть функция f = sqrt(1 + x) - 1. Для x, близких к нулю (порядка 10 в -15 степени) не хватает точности типа double. Относительная ошибка...

Построить график изменения вероятности нахождения в состояниях со временем (по результатам расчетов)
Здравствуйте, помогите пожалуйста доработать программу по заданию. Нужно доделать граф состояний. Построить график изменения...

Электротехника: График изменения тока протекающего через конденсатор на переменном E. Из расчетов переходных процессов
Добрый вечер, подскажите пожалуйста как в маткаде построить графики: Где &quot;71,565&quot; и &quot;108,435&quot; в градусах. Можно ли сделать...

Непредсказуемые синие экраны
Здравствуйте. Уже на протяжении года меня преследуют эти ошибки, но проявляются они абсолютно в разные моменты. На моей памяти было...

Stripos - непредсказуемые результаты
Получаю значения фун-ей stripos и далее использую его в фун-и substr_replace. Все замечательно работает. Но при вызове var_dump...


Искать еще темы с ответами

Или воспользуйтесь поиском по форуму:
46
Ответ Создать тему
Новые блоги и статьи
Саморегулирующийся социальный контракт для сервера cross-section.
Hrethgir 14.08.2026
С кодом конечно таких глубоких размышлений пока не было, впрочем я уже привык к алгоритмизации. Суть предмета записи: снова в диалоге с нейросетью (я взял пока себе ник для учётки админа - Rector). . . .
Часы электронные
Uhbif79 12.08.2026
Выкладываю программу часов. Программа позволяет: 1. Использовать системное время и дату, 2. Есть возможность вводить время и дату вручную. 3. Реализованы 2 будильника: начало и конец рабочего дня. . . .
Часы с будильником на основе класса QLCDNumber
Uhbif79 12.08.2026
Всем добрый день, выкладываю программу часов с будильником на основе класса QLCDNumber. Здесь я пробовал самостоятельно создавал классы, впервые столкнулся с видимостью переменной одного класса из. . .
Установка MinGW GCC 16.2 и CMake
8Observer8 10.08.2026
VK Видео: https:/ / vkvideo. ru/ video-240781534_456239017 YouTube: eY5-5PyI9NM Текстовая версия
Неделя из жизни имитационной модели склада: мои кривые руки растут, откуда надо
anaschu 10.08.2026
Неделя из жизни имитационной модели склада: как я почти написал неправильную логику и что с этим делать Работаю сейчас над учебно-рабочим проектом: строю в AnyLogic имитационную модель процессов. . .
Калькулятор для расчета родства
russiannick 07.08.2026
1. Задача: Создать калькулятор для расчета родства. Родственных связей существует 8 ступеней, такие как: p - отец P - мать q - муж Q - жена b - брат B - сестра s - сын S - дочь
Мир по моей воле
kumehtar 07.08.2026
Когда-то кажется, что всё просто. Ты весь такой светлый. Причиняешь добро. Борешься за справедливость в этом тёмном мире. Потом начинаешь замечать одну неприятную вещь. Почти каждый хороший. . .
Кредитный калькулятор
Maks 05.08.2026
Решение задачи по прикладной информатике средствами 1С. Задача: Напишите приложение-калькулятор, которое помогает рассчитывать параметры кредита для аннуитетного и дифференцированного видов. . .
КиберФорум - форум программистов, компьютерный форум, программирование
Powered by vBulletin
Copyright ©2000 - 2026, CyberForum.ru