С Новым годом! Форум программистов, компьютерный форум, киберфорум
OpenCL
Войти
Регистрация
Восстановить пароль
Карта форума Темы раздела Блоги Сообщество Поиск Заказать работу  
 
 
Рейтинг 4.56/9: Рейтинг темы: голосов - 9, средняя оценка - 4.56
194 / 29 / 5
Регистрация: 11.04.2015
Сообщений: 726
1

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

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

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

Не по теме:

Ну, тут правда то ли я тупой, то ли лыжи не едут. Я ДВА ДНЯ потратил на отладку, чтобы узнать, что проблема ни разу не в моих алгоритмах - просто 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;
}

Если просто запустить эту программу и сравнить значения дистанции, то мы увидим
Код
CPU dist    = 3.359239459231472269e-01
NV dist     = 3.359239460197458449e-01
, что совпадают лишь первые 8 знаков.
Главная причина этого в том, что значение переменной p2 оказывается таким:
Код
CPU p2      = 5.272612557671862987e-08
NV p2       = 5.272612559188061227e-08
и здесь совпадают 9 знаков.
Однако, если мы после основного расчёта (внутри kernel) дистанции продублируем его, разделив расчёт переменной p2 на три отдельных операции, то, О ЧУДО, вывод окажется следующим:
Код
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.
Отсюда и сам мой вопрос - а что, собственно, происходит? Почему, я даже не знаю, как это правильно назвать, видеокарта не хочет с первого раза нормально считать? Почему корректный расчёт происходит только при избыточном дублировании?
На всякий случай, прикладываю полный консольный вывод приведённой выше программы:
Код
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
Эксперт
94731 / 64177 / 26122
Регистрация: 12.04.2006
Сообщений: 116,782
10.04.2024, 00:29
Ответы с готовыми решениями:

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

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

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

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

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

45
194 / 29 / 5
Регистрация: 11.04.2015
Сообщений: 726
19.04.2024, 01:10  [ТС] 21
Author24 — интернет-сервис помощи студентам
Цитата Сообщение от snake32 Посмотреть сообщение
Можно тут по-подробнее?
Да, конечно, расскажу. Делалось это с использованием MatLab, естественно
Как строится расчёт синуса для точности float (если не CORDIC)? Через интервальные аппроксимации параболами (см. рисунок). Например, точность float (ошибка <1.0E-08) обеспечивается всего лишь ~150 интервалами (если разбиваем от 0 до pi/2), при этом, точности 1.0E-12 достигнуть не удаётся и при числе интервалов в 5000. Более того - точность не растёт при увеличении числа интервалов свыше 3000, а падает. Очевидно, что точности double таким образом достичь не удастся.
Для сходящихся функций (синус, косинус) прекрасно работает ряд Тейлора. Точно так же делим участок [0, pi/2] на N интервалов и раскладываем функцию синуса в ряд Тейлора в точке центра каждого интервала, после чего записываем коэффициенты получившегося полинома (с заранее заданной степенью i[SUB]max[/SUB]) в таблицу:
N k0 (*x0) ki (*xi) kimax (*ximax)
0 2.361868486512384E-02 ... 1.651448768486486E-07
1 ... ... ...
Мне удалось добиться ошибки в <3.0E-16 (по данным MatLab) на 125 интервалах с полиномом 6 степени.
Отдельно стоит сказать о том, как я нахожу в таблице нужные коэффициенты полинома. Поскольку участок равномерно разбит на N интервалов, можно использовать простую линейную функцию вида y = k*x, где y - номер интервала, а x - аргумент функции sin(x). Отмечу, однако, что мой "бенчмарк" C++ сказал мне, что "MatLab не прав", и выдал максимальную ошибку в 1.2E-14, что, в целом, меня устроило.
С тангенсом я не стал разбираться и бахнул tan(x) = sin(x) / cos(x). Дорого, но я использую эту функцию довольно редко.
Функции же арксинуса и арккосинуса расходятся, поэтому Тейлор тут мдееееее.. Да тут, как оказалось, что угодно будет мдееееее. Оказывается, арк-функции вычисляют через арктангенс, который сходится по формулам, использующим корень. Но разложить функцию арктангенса через ряд Тейлора, как я понял, невозможно, поэтому пользуются другими "аппроксиматорами". Я решил использовать аппроксимацию Паде: степень полинома числителя - 3, знаменателя - 5, число интервалов - 159. Эта функцию определена для любого x, кроме inf и NaN (нужно для inf - бахайте туда x=квинтиллион). Порядок точности получился тем же, что и для синуса с косинусом. Оговорки есть только с интервальным разбиением - здесь не одно равномерное разбиение, а три (x: 0-10, 10-20, 20-40) и пять определённых вручную "широких" интервалов, неравномерно уходящих в "бесконечность".
Ссылок нет, поскольку когда я это всё делал, большей теорией, чем та, что я сейчас написал, я не обладал. Неделя экспериментов, и программная реализация тригонометрии для double готова. В шейдеры, кстати, я загружал таблицы через SSBO; хотел захардкодить, но компилятор оПтИмИзИрУеТ константы, если они кажется ему слишком большими. Ох уж эти 10 Кб информации. Спасибо, умный компилятор.

Не по теме:

Кстати я также реализовал CORDIC для double в шейдерах (только для синуса). Работало это чудо в 11 (ужасть) раз медленнее, чем моя "табличная" тригонометрия. Я не ожидал, что операции смещения битов в целых числах дороже, чем возведения в степень, умножения и сложения.


Цитата Сообщение от snake32 Посмотреть сообщение
да вроде сводится к нахождению прямой-пересечения двух плоскостей
Да, но так мы найдём две точки, а точкой пересечения является лишь одна из них (это ж дуги). Нужна проверка. Проверка простая - найти 6 дистанций и сравнить их попарно. Именно здесь моя коса находит себе камень - какая-то дистанция оказывается точной лишь до 3 знака и руинит проверку.
Миниатюры
Непредсказуемые изменения точности расчётов внутри kernel  
1
3460 / 1648 / 236
Регистрация: 26.02.2009
Сообщений: 8,051
Записей в блоге: 5
19.04.2024, 01:31 22
Продолжаю эксперименты с точностью)
Получаю интересные результаты, если коротко: зря мы гоним на точность OpenCL,
или, брать результат работы NativeCPU код за эталон - так себе идея

У меня есть ф-ия, которая генерит вершины спирали на поверхности сферы.
Когда я увеличил количество вершин за оборот до 36000, то есть вершины генерятся с шагом в 0,01 градус
и начало спирали (первая вершина) в центре(на нулевом расстоянии) то увидел следующее:
GPU OpenCL 2.0 gfx1030 15.98Gb
| Max delta (km): 0.0000949369370937347
| k: 3 // индекс вершины на спирали где максимальная погрешность
| ___CPU(km): 0 // значение расстояния вычисленное на CPU
| OpenCL(km): 0.0000949369370937347 // значение расстояния вычисленное на OpenCL
| Δlongitude: 6.97205848609883E-10 // разница по долготе между измеряемыми вершинами, градусы
| Δlatitude: 7.49421623424951E-7 // разница по широте между измеряемыми вершинами, градусы
--------------------------------------
CPU OpenCL 1.2 AMD Ryzen 9 5950X 16-Core Processor 31.91Gb
| Max delta (km): 0.0000949369370937347
| k: 4
| ___CPU(km): 0
| OpenCL(km): 0.0000949369370937347
| Δlongitude: 1.23948495911463E-9
| Δlatitude: 9.99228731757285E-7
--------------------------------------
Погрешность стала неприлично большой: 1e-5, что в 1000 раз хуже чем в прошлый раз.
Видно, что CPU рапартует нам о нулевом расстоянии, тогда как разница между координатами точек
точно не нулевая (от 1e-7 до 1e-10)
Тогда я решил распечать первые 10 расстояний посчитаных на CPU:
| distCPU[0] = 0
| distCPU[1] = 0
| distCPU[2] = 0
| distCPU[3] = 0
| distCPU[4] = 0
| distCPU[5] = 0.000134261104008121
| distCPU[6] = 0.00018987387418747
| distCPU[7] = 0.00018987387418747
| distCPU[8] = 0.00018987387418747
| distCPU[9] = 0.000232547053622355
Кликните здесь для просмотра всего текста
GPU OpenCL 2.0 gfx1030 15.98Gb
distCL[0] = 0
distCL[1] = 0
distCL[2] = 0
distCL[3] = 0.0000949369370937347
distCL[4] = 0.0000949369370937347
distCL[5] = 0.000134261104008121
distCL[6] = 0.000164435598561319
distCL[7] = 0.000189873874187469
distCL[8] = 0.000212285444917212
distCL[9] = 0.000251179525784205

CPU OpenCL 1.2 AMD Ryzen 9 5950X 16-Core Processor 31.91Gb
distCL[0] = 0
distCL[1] = 0
distCL[2] = 0
distCL[3] = 0
distCL[4] = 0.0000949369370937347
distCL[5] = 0.000134261104008121
distCL[6] = 0.000134261104008121
distCL[7] = 0.000189873874187469
distCL[8] = 0.000212285444917212
distCL[9] = 0.000251179525784205

Короче, можно смело резать долготу < 1.25E-9 и широту < 1.0e-6
Но это тоже не сильно поможет так как напимер distCPU 6,7,8 абсолютно идентичны, тогда как на OpenCL - нет
1
194 / 29 / 5
Регистрация: 11.04.2015
Сообщений: 726
19.04.2024, 15:07  [ТС] 23
Цитата Сообщение от snake32 Посмотреть сообщение
зря мы гоним на точность OpenCL,
или, брать результат работы NativeCPU код за эталон - так себе идея
Ну, смотрите - как я разрабатывал этот алгоритм. Сначала всё было на бумаге. Потом я написал реализацию на C++. Чтобы её проверить и отладить, я составил несколько тестовых наборов данных и вручную рассчитал результаты для каждого из них (ну, как вручную - с помощью систем символьной математики), чтобы точно понимать, что, где и как должно "идти" в любой момент времени.
Как обычно, чуда не произошло и на этот раз, и моя C++-реализация обрабатывает тестовые наборы не только правильно, но и с той же точностью на каждом шаге, какую дала символьная математика (1e-16).
Именно поэтому, по крайней мере, в рамках данной проблемы, я и принимаю результаты, получаемые на CPU, за эталон, с которым и сравниваю данные, идущие с GPU. Да и давайте честно - библиотеке <math.h> сто лет в обед; не доверять ей - не доверять ничему.
Далее. Как я уже говорил, вычислительный шейдер OpenGL (4.3) выдаёт мне те же данные, что и CPU с оговоркой на точность моей переопределённой тригонометрии. Но OpenCL настолько далеко даже от вычислительного шейдера (спецификация к которому не менялась с 2014 года), что я уверен, что дело в OpenCL. Возможно, дело в OpenCL 3.0, который у меня выставляется автоматически, но я пока не понял, как указать программе конкретную версию. Возможно, это просто ошибка в драйверах, но даже на примере вашего .exeшника мы видим неадекватное снижение точности даже по сравнению с GPU-собратом RX6800XT. Опять же, упомяну, что если я ставлю float в OpenGL, то точности не хватает только при определённых входных данных, а если в OpenCL - оно в принципе считает неверно.
Цитата Сообщение от snake32 Посмотреть сообщение
Короче, можно смело резать долготу < 1.25E-9 и широту < 1.0e-6
А это можно, да Разобраться бы со скачками точности, а там и ограничить можно.
0
824 / 241 / 47
Регистрация: 24.01.2013
Сообщений: 747
19.04.2024, 17:20 24
Цитата Сообщение от Ромуальд_7 Посмотреть сообщение
вычислительный шейдер OpenGL (4.3) выдаёт мне те же данные, что и CPU с оговоркой на точность моей переопределённой тригонометрии
а если в программу OpenCL вставить ваши самописные ряды? какая точность?

З.Ы. и там еще такой нюанс, что математический сопроцессор в CPU считает точнее чем double, там 19-20 десятичных знаков. Потом при записи в память результат округляется до 15-16 знаков.
А видяха если и поддерживает double то в точности по стандарту 15-16 знаков.

В общем, я бы все таки поставил драйвер OpenCL для процессора и подсунул ему девайс-ЦПУ.
Велика вероятность, что дело не в OpenCL, а в видеокарте, тем более, что это Nvidia, там дядюшка Хуанг "срезает углы" где только можно.
1
3460 / 1648 / 236
Регистрация: 26.02.2009
Сообщений: 8,051
Записей в блоге: 5
19.04.2024, 17:40 25
Цитата Сообщение от Ромуальд_7 Посмотреть сообщение
что я уверен, что дело в OpenCL
Если будет время попробуйте посчитать расстояние между точки(center) и точками cpt 4 и 5 на CPU QT
потому как у меня CPU возвращает одинаковое расстояние 0.0000949369370937348 км
center.lng = 37.613094187369
cpt[4].lng = 37.6130941886085
cpt[5].lng = 37.6130941893057

center.lat = 55.7496131585696
cpt[4].lat = 55.7496141577983

cpt[5].lat = 55.7496144076053

// result
distCPU[4] = 0.0000949369370937348
distCL[4] = 0.0000949369370937347

distCPU[5] = 0.0000949369370937348
distCL[5] = 0.000134261104008121
1
194 / 29 / 5
Регистрация: 11.04.2015
Сообщений: 726
20.04.2024, 01:49  [ТС] 26
Цитата Сообщение от snake32 Посмотреть сообщение
Если будет время
Итак Табличка (все расстояния приведены к метрам)
center-cpt[4] center-cpt[5]
CPU (Qt, <math.h>) 1.111111078498496080e-01 1.388888813060294714e-01
GPU (OpenCL 3.0) 1.111111085974727103e-01 1.388888811400676992e-01
GPU (OpenCL 3.0, -cl-opt-disable) 1.111111085571846874e-01 1.388888813060294714e-01
GPU (CS, OpenGL 4.3) 1.1111110798571737e-01 1.3888888127965804e-01
snake32 0.949369370937348e-01 0.949369370937348e-01
QGIS 1.112267452903073e-01 1.39037441788863e-01
Что мы видим - прежде всего то, что QGIS не согласен ни со мной, ни с Вами, но мои результаты ближе.
Теперь о главном - первый замер показывает, что вычислительный шейдер справился в десять раз лучше со своей работой. Отключение же оптимизаций не помогло OpenCL'ю перестать врать и никак не приблизило его даже к уровню шейдера (как я и говорил - столь же неправильные, но другие, значения).
Второй замер оказался интереснее (перепроверено дважды). Вычислительный шейдер снова сильно обходит OpenCL-реализацию, однако лишь до тех пор, пока не ставишь флаг отключения оптимизаций - с флагом совпадают 19 (!все!) знаков, что, помятуя о точности double, определяемой 16 знаками максимум, говорит нам о том, что видеокарта применила точно такой же механизм, что применил процессор (ведь абсолютно одинаковый мусор с 17 по 19 знаки может идти только в этом случае).
И вот об этом "красивом" случае я и говорю - если вычислительный шейдер ведёт себя последовательно, выдавая мне "9.5" правильных знаков из 16, то kernel делает, что захочет - на одних входных данных на тебе 8 точных знаков, и ничего ты с этим не поделаешь, а на других - на тебе 9 знаков, но если отключишь оптимизацию, то, так уж и быть, посчитаю точно. А почему для первого набора ты не хочешь считать точно с отключенной оптимизацией? Принцип "хочу / не хочу"? Ну, компьютеры не так работают (вроде бы). Драйверы NVidia шалят, или в OpenCL 3.0 напортачили?

Не по теме:

Меня, как программиста-учёного, глубоко расстраивают ситуации, в которых ты, даже используя дорогущие профессиональные решения (OpenCL 3.0.15 + CUDA 12.4 + RTX4080), не имеешь физической возможности собрать программу так, чтобы она работала правильно. Конечно, не исключаю, что именно игровые драйверы RTX4080 вносят неоценимый вклад во все мои проблемы.


Как бы то ни было, в скорейшем времени я запущу тестовую программу на A100, и мы сможем узнать больше о виновниках торжества.
Цитата Сообщение от _Develop Посмотреть сообщение
а если в программу OpenCL вставить ваши самописные ряды? какая точность?
На самом деле, сообщение от snake32 пришлось как нельзя кстати, с точки зрения Вашего вопроса. В таблице выше присутствует ячейка, которая (на мой взгляд) полностью исключает влияние встроенной в OpenCL тригонометрии на происходящее. Есть много фактов того, что в вакууме (и при редких обстоятельствах) встроенная тригонометрия работает неотличимо от <math.h>'овской и не может являться причиной бед. Поэтому я не вижу особого смысла в тесте моей тригонометрии - она объективно хуже, хотя компилятор, действительно, может прийти в заблуждение и перестать оптимизировать тригонометрию (если он это делает). Может, попробую после тестов на A100.
Цитата Сообщение от _Develop Посмотреть сообщение
математический сопроцессор в CPU считает точнее чем double, там 19-20 десятичных знаков. Потом при записи в память результат округляется до 15-16 знаков.
А видяха если и поддерживает double то в точности по стандарту 15-16 знаков.
Это да, но и снова мы приходим к ячейке таблицы, в которой результаты от CPU и GPU совпали по 19 знакам. А должны были по 15-16, если бы они считали по-разному. Мистерия.
Цитата Сообщение от _Develop Посмотреть сообщение
В общем, я бы все таки поставил драйвер OpenCL для процессора и подсунул ему девайс-ЦПУ.
Да, тоже займусь скоро. Хоть меня и пугает перспектива "поломки" чего-нибудь, о которой говорил snake32.
Цитата Сообщение от _Develop Посмотреть сообщение
Велика вероятность, что дело не в OpenCL
Ну, вот даже если это так, я не буду готов снять ответственность с разработчиков - Khronos Group - которым в голову пришла гениальная идея об уменьшении контроля разработчика над процессом компиляции kernel'а. С точки зрения научных расчётов, на данный момент OpenCL 3.0 неприменим, поскольку есть лишь один флаг сборки, способный увеличить (привести к уровню, на котором она должна быть по умолчанию) точность, который хочет - работает, не хочет - не работает, а скорость режет. Как я понял, на данный момент не существует инструмента, позволяющего компилировать kernel отдельно от вызова cl::program::build(). Есть очень старые инструменты компиляции OpenCL-кода в PTX-формат, но они не развиваются со времён OpenCL 1.2.
1
824 / 241 / 47
Регистрация: 24.01.2013
Сообщений: 747
20.04.2024, 17:10 27
Цитата Сообщение от Ромуальд_7 Посмотреть сообщение
Да, тоже займусь скоро. Хоть меня и пугает перспектива "поломки" чего-нибудь, о которой говорил snake32.
Да там ничего страшного, установка интеловского драйвера OpenCL затирает доступ к видеокартам NV или AMD и в списке девайсов останется только процессор и встройка от Интел (iGPU). Повторная установка драйвера видеокарты все исправит.

Добавлено через 1 час 7 минут
Цитата Сообщение от Ромуальд_7 Посмотреть сообщение
флаг отключения оптимизаций - с флагом совпадают 19 (!все!) знаков
а что по быстродействию с флагом и без него?
может отключение флага оптимизации вообще приводит к тому что расчет идет на процессоре?
потому то и совпадают все 19 знаков,
иначе такого просто физически не может быть
1
194 / 29 / 5
Регистрация: 11.04.2015
Сообщений: 726
20.04.2024, 18:38  [ТС] 28
Цитата Сообщение от _Develop Посмотреть сообщение
а что по быстродействию с флагом и без него?
В данном случае разница такая, что замерам не поддаётся (<std::chrono> выводит наносекундную разницу, и верить этому, или нет - я не знаю). Моя большая программа, однако, при входных данных, не вызывающих сбоев, показывает разницу втрое. Но я не исследовал вопрос по-нормальному.
Цитата Сообщение от _Develop Посмотреть сообщение
вообще приводит к тому что расчет идет на процессоре?
Ну, это было бы странно, поскольку OpenCL пишет, что используется RTX4080, во время исполнения программы "дёргается" график "GPU utilization" в мониторе ресурсов, немного заполняется её память (по тому же монитору), да и процессор библиотека не видит, ведь я ещё не поставил соответствующие драйверы. Очень сомневаюсь, одним словом.
Цитата Сообщение от _Develop Посмотреть сообщение
иначе такого просто физически не может быть
Да ну почему же. Если в игровых GPU инструкции x64 эмулируются, то что мешает производителю эмулировать и x87?
0
824 / 241 / 47
Регистрация: 24.01.2013
Сообщений: 747
20.04.2024, 19:29 29
Цитата Сообщение от Ромуальд_7 Посмотреть сообщение
да и процессор библиотека не видит, ведь я ещё не поставил соответствующие драйверы.
неа, процессор будет виден, просто драйвер от Интела лучше работает (использует инструкции AVX и др.), соответственно быстродействие с ним в несколько раз лучше. В теме snake32 тут рядом посмотрите тесты, там в его программе 2 процессора можно выбрать в списке, и работают по разному они. Т.е. второй появится после установки драйвера OpenCL от Интел. Разница в скорости расчета раз в 5-6.

Цитата Сообщение от Ромуальд_7 Посмотреть сообщение
Да ну почему же. Если в игровых GPU инструкции x64 эмулируются, то что мешает производителю эмулировать и x87?
эмуляция если и будет то с точностью 15-16 знаков, да и какая эмуляция даст совпадение по всем знакам?
сильно сомневаюсь
0
194 / 29 / 5
Регистрация: 11.04.2015
Сообщений: 726
20.04.2024, 22:20  [ТС] 30
Цитата Сообщение от _Develop Посмотреть сообщение
неа, процессор будет виден,
Ну, слушайте, не виден
Что я делаю в коде:
C++ (Qt)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
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;
}
qDebug() << "num_devices:" << all_devices.size();
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);
И что я получаю на выходе:
Код
Using platform:  "NVIDIA CUDA"
num_devices: 1
Using device: "NVIDIA GeForce RTX 4080"
Using CL version: "OpenCL 3.0 CUDA"
При попытке выставить устройство, отличное от 0, я получаю падение программы, да и не удивительно - ведь она видит лишь одно пригодное устройство.
Поэтому я не сомневаюсь, что расчёт идёт исключительно на GPU. Да и колонка №1 в таблице говорит о том, что даже с -cl-opt-disable регулярно идут бредовые данные, то есть, на CPU программа не переходит.
Цитата Сообщение от _Develop Посмотреть сообщение
эмуляция если и будет то с точностью 15-16 знаков, да и какая эмуляция даст совпадение по всем знакам?
Что я выяснил из интернетов - если вычислители отправляют два числа, у которых имеет место полное совпадение в битовом представлении (все 64 бита одинаковы), то средство вывода их на экран в десятичном виде будет давать одно и то же, даже мусор, теоретически, будет одинаковым. Я вывел на экран 30 значений после запятой, и они также совпали. Предполагаю, что GPU просто выдала абсолютно такой же набор битов. Также я не думаю, что это ошибка или баг.
0
824 / 241 / 47
Регистрация: 24.01.2013
Сообщений: 747
20.04.2024, 23:56 31
Цитата Сообщение от Ромуальд_7 Посмотреть сообщение
Ну, слушайте, не виден
ну не знаю, у меня был виден изначально. Потом я установил драйвер OpenCL от Интела и в списке девайсов появился еще один процессор.
Теперь в списке 3 девайса - видеокарта и два процессора, при том что эти процессоры считают по разному, как писал выше.
0
3460 / 1648 / 236
Регистрация: 26.02.2009
Сообщений: 8,051
Записей в блоге: 5
22.04.2024, 14:56 32
Цитата Сообщение от _Develop Посмотреть сообщение
ну не знаю, у меня был виден изначально.
У вас похоже что-то уже стояло или версия драйвера OpenCL от интел другой версии...
или винда для цп-intel сама ставит эксклюзивно (хотя когда у меня был i7-6700K тоже ничего не было без драйвера),
но для амд без установки драйвера процессор не участвует точно.

Ромуальд_7, короче я не знаю как у вас такая точность получается...
я уже и векторную форму проверил и в ТГ-группе Delphi/Lazarus спросил везде +/- одно и тоже
и получается что не более 0,5 метра точность (5e-4 км на около нулевых расстояний), не важно где считать (CPU/OpenCL), да там и считать особо нечего:
два синуса, три косинуса и арккосинус.
А в векторной форме вообще 3 умножения и арккосинус, думаю, это самый быстрый способ, хотя это надо ещё проверить...
Может позже сделаю бенч, только на Delphi перенесу, чтобы компилятор был другой, на всякий случай.
1
194 / 29 / 5
Регистрация: 11.04.2015
Сообщений: 726
22.04.2024, 19:56  [ТС] 33
snake32, Ну, тут уж что уж
В любом случае, спасибо!
0
3460 / 1648 / 236
Регистрация: 26.02.2009
Сообщений: 8,051
Записей в блоге: 5
23.04.2024, 10:11 34
Цитата Сообщение от Ромуальд_7 Посмотреть сообщение
В любом случае, спасибо!
Да не за что. Самому интересно
Ромуальд_7, а у вас случайно не 32-битный режим?
В 32-битном обычно для вычислений используется сопроцессор x87, где есть режим вычислений 10-байт (extended),
а в x64 - максимум только векторные 8-байтные регистры (double)
0
194 / 29 / 5
Регистрация: 11.04.2015
Сообщений: 726
23.04.2024, 10:16  [ТС] 35
snake32, нене, точно не 32-битный. Да а даже и если так, что бы изменилось? (Просто интересно)
А хотя, вот как работает метод .build()? Как он, опять же, компилирует?
0
824 / 241 / 47
Регистрация: 24.01.2013
Сообщений: 747
23.04.2024, 12:00 36
Ромуальд_7, это то что я писал, сопроцессор x87 считает (внутри себя, т.е. если нет промежуточных сохранений) с точностью Extended (80-bit), а x64 это на расширениях AVX, SSE - тут только 64 bit.
0
194 / 29 / 5
Регистрация: 11.04.2015
Сообщений: 726
23.04.2024, 12:46  [ТС] 37
Цитата Сообщение от _Develop Посмотреть сообщение
это то что я писал
Да я ещё тогда понял
Вопрос-то главный в том, что компилятор kernel'а делает, а не что я указываю для gcc.
0
824 / 241 / 47
Регистрация: 24.01.2013
Сообщений: 747
23.04.2024, 13:01 38
Ромуальд_7, так kernel компилирует драйвер, а gcc это код "обвязки", так же как с вычислительным шейдером.
0
194 / 29 / 5
Регистрация: 11.04.2015
Сообщений: 726
23.04.2024, 13:11  [ТС] 39
Цитата Сообщение от _Develop Посмотреть сообщение
kernel компилирует драйвер
Вооот, но я никак не могу найти подтверждения этому.
Смотрите, в спецификации есть часть:
Код
5.8.6.1. Preprocessor Options
These options control the OpenCL C/C++ preprocessor which is run on each program source before actual compilation.
These options are ignored for programs created with IL.

-D name
Predefine name as a macro, with definition 1.

-D name=definition
The contents of definition are tokenized and processed as if they appeared during translation phase three in a #define directive.
In particular, the definition will be truncated by embedded newline characters.

-D options are processed in the order they are given in the options argument to clBuildProgram or clCompileProgram.
Note that a space is required between the -D option and the symbol it defines, otherwise behavior is implementation-defined.

-I dir
Add the directory dir to the list of directories to be searched for header files. dir can optionally be enclosed in double quotes.

This option is not portable due to its dependency on host file system and host operating system.
It is supported for backwards compatibility with previous OpenCL versions.
Developers are encouraged to create and use explicit header objects by means of clCompileProgram followed by clLinkProgram.
Которая говорит о том, что я мог бы передать параметр для nvcc, вроде вышеупомянутого --fmad=false, но когда я передаю параметр через -D он просто игнорируется. Я не понял, что такое IL, но, похоже, ядро компилируется именно через него.
0
824 / 241 / 47
Регистрация: 24.01.2013
Сообщений: 747
23.04.2024, 13:20 40
Ромуальд_7, с этим не могу помочь, не разбирался с этими параметрами, увы.

Добавлено через 4 минуты
там препроцессор для kernel'а свой, отдельный от С++.
Может есть возможность в самом kernel'е задать дефайн?
0
23.04.2024, 13:20
IT_Exp
Эксперт
87844 / 49110 / 22898
Регистрация: 17.06.2006
Сообщений: 92,604
23.04.2024, 13:20
Помогаю со студенческими работами здесь

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

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

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

Изменения внутри Combobox
Необходимо реализовать следующее: в combobox 5 полей, при изменении какого-либо, измененное...

Изменения внутри файла стилей средствами JS
Доброго времени суток! Подскажите пожалуйста, как можно внести изменения в CSS-файл для...


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

Или воспользуйтесь поиском по форуму:
40
Ответ Создать тему
КиберФорум - форум программистов, компьютерный форум, программирование
Powered by vBulletin
Copyright ©2000 - 2025, CyberForum.ru