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

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

10.04.2024, 00:29. Показов 2923. Ответов 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
Native x86
Эксперт Hardware
 Аватар для quwy
6853 / 3787 / 1024
Регистрация: 13.02.2013
Сообщений: 11,861
10.04.2024, 02:38
Очень похожая ситуация и решение.
1
194 / 29 / 5
Регистрация: 11.04.2015
Сообщений: 735
10.04.2024, 18:44  [ТС]
Цитата Сообщение от quwy Посмотреть сообщение
Очень похожая ситуация и решение.
Благодарю Вас за ответ!
Ситуация похожая, я бы даже сказал - та же, однако решение не приведено, к сожалению.
Там люди приходят к выводу, что необходимо компилировать kernel с параметром
C
1
-fmad=false
однако такого параметра НЕ СУЩЕСТВУЕТ и никогда не существовало, судя по всему. Приведу вырезку из спецификации, прикрепив скрин ниже (или см. первый блок страницы 4 OpenCL 3.0 Reference Guide).
Есть параметр
C
1
-cl-opt-disable
, который, цитирую "отключает все оптимизации кода", но он не меняет ровным счётом ничего. Чтобы убедиться, что компилятор вообще слушает, что я говорю, я бахнул параметр
C
1
-cl-fast-relaxed-math
, который должен упрощать математику, и он действительно упростил - теперь во втором выводе p2 я также получаю округлённое до 9 знаков (неправильное) значение, следовательно, опции компилятор применяет. Возможно, дело здесь и не в оптимизациях, раз их отключение ни на что не повлияло.
К сожалению, в обсуждении по приведённой Вами ссылке нигде не показывается момент применения параметра
C
1
-fmad=false
к процедуре компиляции, поэтому я просто-напросто не могу этого повторить, чтобы проверить. Параметры же
C
1
--fmad=false, fmad=false, -fmad-disable
и прочие синонимичные выводят одинаковую ошибку компиляции - "неизвестный параметр".

Не по теме:

Кроме того, я вынужден отметить, что в том обсуждении вопрошающий в целом не совсем понимает, чего хочет, поскольку выводимые его программами значения совпадают по 15 значащим цифрам, и именно такую точность гарантирует тип double. По сути, у автора вопроса не было проблем с точностью в принципе. В моём же случае бобры сгрызают 6-7 знаков из 15.

Миниатюры
Непредсказуемые изменения точности расчётов внутри kernel  
0
Native x86
Эксперт Hardware
 Аватар для quwy
6853 / 3787 / 1024
Регистрация: 13.02.2013
Сообщений: 11,861
11.04.2024, 02:33
Цитата Сообщение от Ромуальд_7 Посмотреть сообщение
однако такого параметра НЕ СУЩЕСТВУЕТ
А чем вы компилируете кернел?
Это странно, потому что поиск по этому выражению дает многовато результатов, как для несуществующего параметра.
0
194 / 29 / 5
Регистрация: 11.04.2015
Сообщений: 735
11.04.2024, 10:23  [ТС]
Цитата Сообщение от quwy
А чем вы компилируете кернел?
Штатными средствами библиотеки
C++
1
2
#include <CL/cl.hpp> // Windows
#include <CL/opencl.hpp> // Linux
, а именно (часть из управляющей программы выше):
C++
1
2
3
4
5
6
cl::Program program(context,sources);
const char * params = "-fmad=false"; // выдаст ошибку — «неизвестный параметр»
    if(program.build(default_device, params) != CL_SUCCESS){
        qDebug() << " Error building: " << QString::fromStdString(program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(default_device));
        return 1;
    }
А поиск выдаёт так много результатов, я полагаю, потому, что речь там идёт о CUDA, а не об OpenCL. Как я уже упоминал, поиск по этому же выражению (и синонимичным) в официальных спецификациях (к версиям 1, 2 и 3) не даёт результатов.
0
 Аватар для snake32
3502 / 1685 / 236
Регистрация: 26.02.2009
Сообщений: 8,368
Записей в блоге: 6
11.04.2024, 14:30
Ромуальд_7, когда я програмил OpenCL 1.0 я вставлял в каждый кернел такую строку:
Code
1
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
В ряд ли вам это поможет)
Могу прогнать ваш готовый *.exe на Radeon 6800XT, интересно, будут ли такие же потери точности...
1
827 / 244 / 47
Регистрация: 24.01.2013
Сообщений: 750
11.04.2024, 15:40
Ромуальд_7, я где то тут писал уже, что столкнулся с тем, что видяха считает тригонометрию с низкой точностью.
Ну это и понятно, для вывода графики точности большой не нужно.

Попробуй ту же программу OpenCL выполнить на процессоре, т.е. девайс подсунуть CPU вместо GPU.
Интересно какие результаты будут...
0
 Аватар для snake32
3502 / 1685 / 236
Регистрация: 26.02.2009
Сообщений: 8,368
Записей в блоге: 6
11.04.2024, 16:14
Цитата Сообщение от _Develop Посмотреть сообщение
что видяха считает тригонометрию с низкой точностью
У меня вроде нормально считало, когда я сравнивал с ЦПУ, но это было давно. Может сегодня уже по другому.
0
194 / 29 / 5
Регистрация: 11.04.2015
Сообщений: 735
11.04.2024, 17:23  [ТС]
snake32, да, эта строка, к сожалению, ничего не изменила.
А можете ли Вы просто взять мой kernel и вставить в одну из Ваших хостовых программ? Дело в том, что я прогаю на Qt, а с точки зрения сборки проекта в один .exe для исполнения на машине без Qt под Windows это — головная боль. Нужно собирать в статику, но я не знаю, что там по OpenCL в статических сборках. На линуксе как-то поудобнее это реализовано.

Не по теме:

Я пока не умею в распространение на мультиплатформе, уж простите :)


Цитата Сообщение от _Develop
столкнулся с тем, что видяха считает тригонометрию с низкой точностью
Так в том и дело, что тригонометрия-то заявлена как double-совместимая, хоть в спецификации и говорится о каком-то «округлении результата до чётного». Кроме того, основная проблема не в неточном значении, а в том, что если конструкция «Б» вычисляется после конструкции «А», то «А» = 9 корректных знаков, а «Б» = 16 знаков; конструкции аналогичны по смыслу и порядку операций. Но если кажущуюся рабочей конструкцию «Б» вставить ВМЕСТО «А», она также начинает давать 9 корректных знаков. Это просто магия — если уж GPU руководствуется одним принципом для всей программы, то откуда такое поведение?
Цитата Сообщение от _Develop
Попробуй ту же программу OpenCL выполнить на процессоре, т.е. девайс подсунуть CPU вместо GPU.
Интересно какие результаты будут...
Чуть позже попробую, да. Однако думается мне, что в зависимости от указанного устройства выбирается компилятор, и для CPU будет выбран g++, который такого не наделает
0
194 / 29 / 5
Регистрация: 11.04.2015
Сообщений: 735
12.04.2024, 16:47  [ТС]
Товарищи, новый симптом!
В представленном мной выше kernel-коде есть строки (69 - 77)
C
1
2
3
4
5
6
7
8
9
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);
реализующие корректный расчёт.
Я закомментировал строки печати p2_h1 и p2_h2, и p2_sum стала считаться неточно.
p2_sum вновь начинает считаться точно, если оставить печать хотя бы одной промежуточной переменной.
В связи с этим новый вопрос - GPU всегда ведёт себя как лентяй - делает плохо, пока никто не смотрит? Нет, ну правда - что это?
0
 Аватар для snake32
3502 / 1685 / 236
Регистрация: 26.02.2009
Сообщений: 8,368
Записей в блоге: 6
17.04.2024, 18:31
Набросал проектик.
Максимальная разница OpenCL vs CPU x86_64 4.50651360534016E-8 километров.
То есть меньше 0,1 мм на поверхности Земли.
Ромуальд_7, куда ещё точнее то?
Код ядра OpenCL:
C++
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
__constant double rzem = 6371.11;
__constant double eps   = 1e-9;
 
__kernel void keDist(   __global double * dist, // (km)
                                __global double2 * coords,//lon lat (deg)
                                            const double2 pt ) // lon lat (deg)
 
{   uint index = get_global_id(0);
    double2 cp = radians( coords[index] );
    double2 st = radians( pt );
    double sin_sr1 = sin( st.y );
    double sin_sr2 = sin( cp.y );
    double cos_sr1 = cos( st.y );
    double cos_sr2 = cos( cp.y );
    double v = clamp( sin_sr1*sin_sr2 + cos_sr1*cos_sr2 * cos( st.x-cp.x ), -1.0, 1.0 );
    dist[index] = fmax( rzem * acos( v ), 0.0 );
}
Код Lazarus x86_64:
Delphi
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
function Distance( constref pt0,pt1:TDVec2 ):double;
    var cp,st:TDVec2;
        sin_sr1,sin_sr2,
      cos_sr1,cos_sr2,v:double;
begin
  cp := radians( pt0 );
  st := radians( pt1 );
 
  sin_sr1 := sin( st.lat );
  sin_sr2 := sin( cp.lat );
  cos_sr1 := cos( st.lat );
  cos_sr2 := cos( cp.lat );
 
  v := clamp( sin_sr1*sin_sr2 + cos_sr1*cos_sr2 * cos( st.lng-cp.lng ), -1.0, 1.0 );
  Result := max( Rearth*arccos( v ), 0.0 );
end;
Цитата Сообщение от _Develop Посмотреть сообщение
Попробуй ту же программу OpenCL выполнить на процессоре, т.е. девайс подсунуть CPU вместо GPU.
Интересно какие результаты будут...
Абсолютно такая же погрешность на любом OpenCL-устройстве.
Миниатюры
Непредсказуемые изменения точности расчётов внутри kernel   Непредсказуемые изменения точности расчётов внутри kernel  
Вложения
Тип файла: 7z Distance.7z (724.9 Кб, 5 просмотров)
2
194 / 29 / 5
Регистрация: 11.04.2015
Сообщений: 735
17.04.2024, 20:37  [ТС]
snake32, к сожалению, не всё так просто. Эта функция в вакууме действительно выдаёт более, чем приемлемые результаты, но дело в том, что это — лишь ядро моей большой программы, которое вызывается миллионы раз за один прогон. Это ядро сверху обёрнуто некоторым количеством проверок в виде дополнительных операций, и в результате при некоторых трудноуловимых, но катастрофичных сценариях точность падает до третьего знака после запятой (а у меня проверка стоит именно по третьему знаку — дальше некуда — будут ложные срабатывания). В какой-то момент kernel просто виснет в цикле while, не пройдя проверку на выход из него; в этом случае разница между GPU и CPU составляет 14 знаков (была дооооолгая отладка). Отмечу, что среднее рассогласование в ответственный момент варьируется от 7 правильных знаков после запятой, до уже упомянутых трёх.
На самом деле, причина даже ясна — в функции расчёта дистанции есть часть .. cos(delta_lambda), и думается мне, что в какой-то момент оптимизированная компилятором функция не видит разницы между cos(0.000000032467178) и cos(0), и пожалуйста — точность 7 знаков, которая далее после пяти операций упадёт до 3.
Мой основно вопрос в том ведь, как заставить kernel работать так, как он может работать. На тестовом примере (приведённом здесь) УДАЁТСЯ заставить GPU правильно всё посчитать, разбив операцию на три элементарных и прописав флаг -cl-opt-disable, но это не работает на большой программе — цифры идут другие, но в равной степени неправильные.
0
 Аватар для snake32
3502 / 1685 / 236
Регистрация: 26.02.2009
Сообщений: 8,368
Записей в блоге: 6
17.04.2024, 20:59
Ромуальд_7, можешь мой экзешник запустить?

Цитата Сообщение от Ромуальд_7 Посмотреть сообщение
в функции расчёта дистанции есть часть .. cos(delta_lambda), и думается мне, что в какой-то момент оптимизированная компилятором функция не видит разницы
Может надо проверять что если точки близки, то сразу возвращать 0?

Кстати, на домашнем ПК GPU Radeon 6800XT чуть хуже точность.
А на 5950Х одинаковое, но оно и понятно драйвер OpenCL один и тот же стоит от Intel
Миниатюры
Непредсказуемые изменения точности расчётов внутри kernel  
1
194 / 29 / 5
Регистрация: 11.04.2015
Сообщений: 735
17.04.2024, 23:38  [ТС]
Цитата Сообщение от snake32 Посмотреть сообщение
можешь мой экзешник запустить?
Тээээкс.. запустил (приложил скриншот).
Что-то мне жёстко не нравится - разница вдвое. Кстати, мой 5900x прога не увидела, а Ваш 5950x - шмогла.
Цитата Сообщение от snake32 Посмотреть сообщение
Может надо проверять что если точки близки, то сразу возвращать 0?
Звучит как план, конечно, но я же пытаюсь в хорошее математическое ПО С этой точки зрения, числа 1,2345678987654'321 и 1,2345678987654'999 отличаются достаточно, чтобы что-то где-то посчиталось неправильно.

Не по теме:

Я бы, откровенно, плясал, если бы GPU начала выдавать хотя бы 10 совпадающих с CPU знаков


Цитата Сообщение от snake32 Посмотреть сообщение
на домашнем ПК GPU Radeon 6800XT чуть хуже точность.
О - такая же, как у меня - 9,013..E-8?
Миниатюры
Непредсказуемые изменения точности расчётов внутри kernel  
0
 Аватар для snake32
3502 / 1685 / 236
Регистрация: 26.02.2009
Сообщений: 8,368
Записей в блоге: 6
18.04.2024, 11:18
Цитата Сообщение от Ромуальд_7 Посмотреть сообщение
в хорошее математическое ПО С этой точки зрения, числа 1,2345678987654'321 и 1,2345678987654'999 отличаются достаточно, чтобы что-то где-то посчиталось неправильно.
Как по мне:
Лучше пусть работает менее точно, но стабильно везде,
чем местами точнее, но не стабильно.

Цитата Сообщение от Ромуальд_7 Посмотреть сообщение
О - такая же, как у меня - 9,013..E-8?
Получается точность на nVidia 4080 в 2 раза хуже чем на AMD RX6800XT (9.01e-8 vs 4.93e-8)
Надо исходные данные изменить, у меня минимум 100м между точками, и в большинстве случаев именно в этом месте максимальная погрешность, кроме Radeon, у него на 3 точке макс погрешность(i=3). Чуть позже сделаю чтобы с нуля тестило.
Цитата Сообщение от Ромуальд_7 Посмотреть сообщение
Кстати, мой 5900x прога не увидела, а Ваш 5950x - шмогла
Потому что нет драйвера OpenCL для CPU.
Я ставил драйвер от Intel 2016 года, но из-за него могут другие OpenCL-устройства отвалится.
Например видеокарта - придётся сверху переставить драйвер ВК.
Если для вас это не проблема можете скачать OpenCL драйвер здесь

Добавлено через 53 секунды
Цитата Сообщение от Ромуальд_7 Посмотреть сообщение
точность падает до третьего знака после запятой
Это ваще жесть
1
827 / 244 / 47
Регистрация: 24.01.2013
Сообщений: 750
18.04.2024, 11:28
Цитата Сообщение от snake32 Посмотреть сообщение
Максимальная разница OpenCL vs CPU x86_64 4.50651360534016E-8 километров.
Ну так 1e-8 это и есть точность float (7-8 знаков), о чем и речь.

Я конкретных значений не сравнивал, считал интерференционную картину в оптике, так на видяхе картинка была "рваная", а на ЦПУ плавная. Похожее было когда рисовал фрактал, сначала в точности float, потом пришлось перейти на double.

Ромуальд_7, может стоит попробовать считать тригонометрию самому с помощью рядов, а не использовать встроенные функции. Заодно убедиться, что проблема именно в этом, и оценить падение быстродействия на этом.
1
 Аватар для snake32
3502 / 1685 / 236
Регистрация: 26.02.2009
Сообщений: 8,368
Записей в блоге: 6
18.04.2024, 11:59
Цитата Сообщение от _Develop Посмотреть сообщение
Ну так 1e-8 это и есть точность float (7-8 знаков), о чем и речь.
Нет, если делать хоть какие-нибудь вычисления на флоатах, ваши 1e-8 легко превратятся в 1e-5.
А если я перепишу алгоритм вычислений расстояний на сфере с double на float то разница будет 2 км (уже проверял, давно), а тут дельта миллиметры, на расстояниях от 100 метров до 20 км.
Цитата Сообщение от _Develop Посмотреть сообщение
может стоит попробовать считать тригонометрию самому с помощью рядов
Это ж циклы использовать - долго. + накопительная погрешность.

Можно отказаться от тригонометрии в пользу векторной математики, но тогда
1) хранить вершины как единичные 3D-вектора (это + 50% памяти, так как вместо 2 double'ов надо юзать 3 на вершину)
2) придётся постоянно конвертировать из 3D-декартовой в сферическую и обратно
3) всё рано для расчёта расстояния потребуется arccos, ну хотя бы он будет 1

По-хорошему надо и такую серию тестов на точность и скорость делать.
1
194 / 29 / 5
Регистрация: 11.04.2015
Сообщений: 735
18.04.2024, 13:28  [ТС]
Цитата Сообщение от snake32 Посмотреть сообщение
Лучше пусть работает менее точно, но стабильно везде,
Вот я ж согласен - поэтому и решил пересесть с вычислительных шейдеров. Но приколов с реализациями драйверов хватает и на OpenCL, смотрю я.
Цитата Сообщение от snake32 Посмотреть сообщение
Чуть позже сделаю чтобы с нуля тестило.
Спасибо - лишние проверки это здорово
Цитата Сообщение от snake32 Посмотреть сообщение
но из-за него могут другие OpenCL-устройства отвалится.
Понял. Тогда мне пока и так хватает трудностей, я думаю. Но спасибо - буду знать.
Цитата Сообщение от snake32 Посмотреть сообщение
Это ваще жесть
Вот ото ж
Цитата Сообщение от _Develop Посмотреть сообщение
сначала в точности float, потом пришлось перейти на double.
Тут, кстати, тоже интересная история. Решил я попробовать бахнуть везде не double, а float - там, где раньше я был недоволен третьим знаком после запятой, теперь вообще другие числа идут; прям настолько другие значения, что об их точности в сравнении с CPU даже речи идти не может. Я ничего не понимаю. В вычислительных шейдерах OpenGL на float'e я, хотя бы, что-то вменяемое получал (в точно таком же алгоритме), а тут просто комшар какой-то. И хоть с флагом -cl-opt-disable собирай, хоть с чем - без разницы.
Цитата Сообщение от _Develop Посмотреть сообщение
может стоит попробовать считать тригонометрию самому с помощью рядов, а не использовать встроенные функции.
На ВШ я так и делал - ведь double там не определён для тригонометрии. И было там бы всё хорошо, если бы не ограничение по времени их исполнения, зашитое в драйверах. Вот прям везде палок навставляют.
Цитата Сообщение от _Develop Посмотреть сообщение
Заодно убедиться, что проблема именно в этом
Ну вот на примере, который я привёл, видно, что GPU может делать нормально, но не хочет, когда программа разрастается. Проблема явно не в точности а в чём-то ещё. Возможно, так NVidia пересаживает всех на CUDA - не знаю.
Цитата Сообщение от snake32 Посмотреть сообщение
Нет, если делать хоть какие-нибудь вычисления на флоатах, ваши 1e-8 легко превратятся в 1e-5.
Вот как раз - как я писал в этом же сообщении - в моём случае значения просто уезжают с "крайне неточных" до "что это за число??? Откуда это?".
Цитата Сообщение от snake32 Посмотреть сообщение
Это ж циклы использовать - долго. + накопительная погрешность.
Кстати в итоге я сделал таблицы (большиииииие) аппроксимации по Тейлору для синуса и косинуса и по Паде для арктангенса. Точность получалась до 14 знака, а по скорости - обходил math.h на 17% (на CPU). Я был очень горд собой, пока не обнаружил, что шейдерами-то я не воспользуюсь нифига xD
Цитата Сообщение от snake32 Посмотреть сообщение
Можно отказаться от тригонометрии
Вот к сожалению, нельзя. Я и тут спрашивал в своё время, и искал информацию и книжки смотрел - без шансов. Вся геодезия на тригонометрии. Как пример - задача о нахождении точки пересечения двух дуг большого круга на сфере (не говоря об эллипсоиде). Вот никак без тригонометрии. Кстати, Ваши три шага в этом алгоритме и применяются - проще нельзя. Именно тут точности и не хватает для прохождения проверок.
0
827 / 244 / 47
Регистрация: 24.01.2013
Сообщений: 750
18.04.2024, 14:18
Цитата Сообщение от Ромуальд_7 Посмотреть сообщение
Вот прям везде палок навставляют.
Это точно, в теории масса способов, а кинься что то делать - сплошные грабли везде.
0
 Аватар для snake32
3502 / 1685 / 236
Регистрация: 26.02.2009
Сообщений: 8,368
Записей в блоге: 6
18.04.2024, 17:34
Цитата Сообщение от Ромуальд_7 Посмотреть сообщение
Кстати в итоге я сделал таблицы (большиииииие) аппроксимации по Тейлору для синуса и косинуса и по Паде для арктангенса. Точность получалась до 14 знака, а по скорости - обходил math.h на 17% (на CPU)
Можно тут по-подробнее? (размер таблиц, код или ссылку откуда формулу брали)
Цитата Сообщение от Ромуальд_7 Посмотреть сообщение
Как пример - задача о нахождении точки пересечения двух дуг большого круга на сфере
да вроде сводится к нахождению прямой-пересечения двух плоскостей
0
Надоела реклама? Зарегистрируйтесь и она исчезнет полностью.
inter-admin
Эксперт
29715 / 6470 / 2152
Регистрация: 06.03.2009
Сообщений: 28,500
Блог
18.04.2024, 17:34
Помогаю со студенческими работами здесь

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

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

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

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

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


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

Или воспользуйтесь поиском по форуму:
20
Ответ Создать тему
Новые блоги и статьи
Новый ноутбук
volvo 07.12.2025
Всем привет. По скидке в "черную пятницу" взял себе новый ноутбук Lenovo ThinkBook 16 G7 на Амазоне: Ryzen 5 7533HS 64 Gb DDR5 1Tb NVMe 16" Full HD Display Win11 Pro
Музыка, написанная Искусственным Интеллектом
volvo 04.12.2025
Всем привет. Некоторое время назад меня заинтересовало, что уже умеет ИИ в плане написания музыки для песен, и, собственно, исполнения этих самых песен. Стихов у нас много, уже вышли 4 книги, еще 3. . .
От async/await к виртуальным потокам в Python
IndentationError 23.11.2025
Армин Ронахер поставил под сомнение async/ await. Создатель Flask заявляет: цветные функции - провал, виртуальные потоки - решение. Не threading-динозавры, а новое поколение лёгких потоков. Откат?. . .
Поиск "дружественных имён" СОМ портов
Argus19 22.11.2025
Поиск "дружественных имён" СОМ портов На странице: https:/ / norseev. ru/ 2018/ 01/ 04/ comportlist_windows/ нашёл схожую тему. Там приведён код на С++, который показывает только имена СОМ портов, типа,. . .
Сколько Государство потратило денег на меня, обеспечивая инсулином.
Programma_Boinc 20.11.2025
Сколько Государство потратило денег на меня, обеспечивая инсулином. Вот решила сделать интересный приблизительный подсчет, сколько государство потратило на меня денег на покупку инсулинов. . . .
Ломающие изменения в C#.NStar Alpha
Etyuhibosecyu 20.11.2025
Уже можно не только тестировать, но и пользоваться C#. NStar - писать оконные приложения, содержащие надписи, кнопки, текстовые поля и даже изображения, например, моя игра "Три в ряд" написана на этом. . .
Мысли в слух
kumehtar 18.11.2025
Кстати, совсем недавно имел разговор на тему медитаций с людьми. И обнаружил, что они вообще не понимают что такое медитация и зачем она нужна. Самые базовые вещи. Для них это - когда просто люди. . .
Создание Single Page Application на фреймах
krapotkin 16.11.2025
Статья исключительно для начинающих. Подходы оригинальностью не блещут. В век Веб все очень привыкли к дизайну Single-Page-Application . Быстренько разберем подход "на фреймах". Мы делаем одну. . .
Фото: Daniel Greenwood
kumehtar 13.11.2025
Расскажи мне о Мире, бродяга
kumehtar 12.11.2025
— Расскажи мне о Мире, бродяга, Ты же видел моря и метели. Как сменялись короны и стяги, Как эпохи стрелою летели. - Этот мир — это крылья и горы, Снег и пламя, любовь и тревоги, И бескрайние. . .
КиберФорум - форум программистов, компьютерный форум, программирование
Powered by vBulletin
Copyright ©2000 - 2025, CyberForum.ru