<bgdev />free

Вход

Задача НЕ за интервю
7

0 1 2 3 4 ....8 9 10 11 12 ....18 19 20 21 22 ....26 27 28 29 30 ....32 33 34 35 36
#12988 (ツ) gat3way
Създадено на 28.09.2020, видяно: 1600 пъти.
Евлампи
gat3way

Въпреки всичко, човек не трябва да изпада в такъв идеализъм и резултатите върху GPU-та са доста по-консистентни и независими от разни външни условия.

Така е ама както сам виждаш има разни неща дето са неочевидни и скалирането на решение не е гарантирано линейно и има разни скрити цени дето не е много ясно как се амортизират. Меренето и оптимизациите са тънка работа

Е, скалирането е друг проблем, но като цяло проблемът с profile-ването и оптимизациите ако и да звучи странно, е доста по-прост върху GPU-та, то там всичко е по-просто. Като цяло, GPU кода е далеч по-благ просто защото обикновено не може да става и дума за сравнение като комплексност с CPU кода. Съответно много по-лесно се осмисля какво и как да се промени. Това хубаво не е тривиално все пак наистина, щото си има неговите номера, наложени от хардуера, но всичко е (беше поне) доста добре документирано, и AMD и NVidia имаха мноооого чудни документи по въпроса за писането на читави кернели и утилизиране на хардуера доколкото може. Дори, колкото и да е тъпо, към един момент, дизасемблираните кернели за AMD ми се виждаха по-четливи и разбираеми от x86 асемблера, въпреки че това беше отпреди GCN, с оная малоумна VLIW архитектура, с клаузите, слотовете по клаузите и прочее тъпотии. С GCN (а и при nvidia) е още по-просто и четливо.

#12989 (ツ) johnfound
Създадено на 28.09.2020, видяно: 1590 пъти.
|

Да, затова пуснах C кода. Няма как да пусна кода на Го, защото зависи от твърде много други файлове. Като преработя версията на tries за C може да я пусна и нея.

Мисля, че имаш грешка в тeзи два реда:

  start = ts.tv_sec * 1000.0 + ts.tv_nsec / 10000000.0; 
...
  end = ts.tv_sec * 1000.0 + ts.tv_nsec / 10000000.0; 

За да преобразуваш ns в ms трябва да делиш на 1Е6, а не на 1Е7.

Не е голям проблем - реално на моята машина грешката е около 500ms в твоя полза. ;-)

#12990 (ツ) ФейкПрофил
Създадено на 28.09.2020, видяно: 1589 пъти.
johnfound
|

Това е кода на чисто C, който сравнява всеки със всеки стринг. На моят лаптоп това е резултата от първите 10 стринга:


0: 4 2011.848800 ms
1: 4 2011.129300 ms
2: 4 2012.483100 ms
3: 4 2012.424100 ms
4: 40 2911.208800 ms
5: 4 2013.644900 ms
6: 9 2011.395600 ms
7: 6 2019.304200 ms
8: 4 2011.881200 ms
9: 6 2021.873200 ms
10: 58 2915.040600 ms

О, мерси, ще имам някакви опорни стойности на какъв хардуер работиш и какво с какво сравняваме. :-)

При мене същият сорс дава:

0: 4 12943.155175 ms
1: 4 12042.569583 ms
2: 4 12948.384989 ms
3: 4 12050.330475 ms
4: 41 12948.288226 ms
5: 4 12044.346433 ms
6: 9 12948.630645 ms
7: 6 12044.806212 ms
8: 4 12045.221166 ms
9: 6 12949.675423 ms
10: 58 12048.351350 ms

Тоест, твоят компютър е приблизително 6 пъти по-бърз от моят.

Впрочем, мисля, че намерих причината за големите забавяния - последните версии на Intel имат ужасни латентности на някои инструкции, при това неочевидни... :-(

Левенщайна ти е бавен :) аз исках да ускоря моя, но това което намерих като "оптимизиран" алгоритъм се оказа 2пъти по-бавен от библиотечката която ползвах :Д

Единствената оптимизация, която успях да направя е да ползвм u8 вместо char, което смъкна 500мс на стринг и стана ~3450мс

#12991 (ツ) |
Създадено на 28.09.2020, видяно: 1587 пъти.
johnfound
|

Да, затова пуснах C кода. Няма как да пусна кода на Го, защото зависи от твърде много други файлове. Като преработя версията на tries за C може да я пусна и нея.

Мисля, че имаш грешка в тeзи два реда:

  start = ts.tv_sec * 1000.0 + ts.tv_nsec / 10000000.0; 
...
  end = ts.tv_sec * 1000.0 + ts.tv_nsec / 10000000.0; 

За да преобразуваш ns в ms трябва да делиш на 1Е6, а не на 1Е7.

Не е голям проблем - реално на моята машина грешката е около 500ms в твоя полза. ;-)

Ugh, натиснал съм една нула повече. Мерси. :)

#12992 (ツ) Delegate
Последно редактирано на 28.09.2020 от Delegate, видяно: 1584 пъти.

ФейкПрофил, и аз мисля, че му е бавен Левенщ*ейнът*. Тоя C# вариант, при мен е по-бърз от неговия.


 public static int LevenshteinDistance(string value1, string value2)
        {
            if (value1.Length == 0)
            {
                return 0;
            }

            int[] costs = new int[value1.Length];

            // Add indexing for insertion to first row
            for (int i = 0; i < costs.Length; )
            {
                costs[i] = ++i;
            }

            int minSize = value1.Length < value2.Length ? value1.Length : value2.Length;


            for (int i = 0; i < minSize; i++)
            {
                // cost of the first index
                int cost = i;
                int addationCost = i;

                // cache value for inner loop to avoid index lookup and bonds checking, profiled this is quicker
                char value2Char = value2[i];

                for (int j = 0; j < value1.Length; j++)
                {
                    int insertionCost = cost;

                    cost = addationCost;

                    // assigning this here reduces the array reads we do, improvment of the old version
                    addationCost = costs[j];

                    if (value2Char != value1[j])
                    {
                        if (insertionCost < cost)
                        {
                            cost = insertionCost;
                        }

                        if (addationCost < cost)
                        {
                            cost = addationCost;
                        }

                        ++cost;
                    }

                    costs[j] = cost;

                }


            }

            return costs[costs.Length - 1];
        }

#12993 (ツ) |
Последно редактирано на 28.09.2020 от |, видяно: 1576 пъти.
gat3way

Ами спокойно може да съм го окензал някъде, не съм го тествал особено много.

Хмм, това __local в OpenCL като __shared__ в CUDA ли е? Ако да, в levenshtein използваш един и същи буфер за всички тредове. Ако не е, res и off в reduce ще са грешни.

#12995 (ツ) gat3way
Създадено на 28.09.2020, видяно: 1560 пъти.

Оооо даммм, прав си. Требва да е двумерен масив и да се индексира спрямо get_local_id(0). Обаче 64 му идва прекалено много иииии...

ptxas error   : Entry function 'levenshtein' uses too much shared data (0xc804 bytes, 0xc000 max)

Затва, трябва да се ограничи workgroup size-а на 32 вместо 64. Това обаче естествено води до прилично влошаване на производителността....иииии...

execution time = 55.904923 seconds
#12996 (ツ) |
Създадено на 28.09.2020, видяно: 1553 пъти.
gat3way

Оооо даммм, прав си. Требва да е двумерен масив и да се индексира спрямо get_local_id(0). Обаче 64 му идва прекалено много иииии...

ptxas error   : Entry function 'levenshtein' uses too much shared data (0xc804 bytes, 0xc000 max)

Затва, трябва да се ограничи workgroup size-а на 32 вместо 64. Това обаче естествено води до прилично влошаване на производителността....иииии...

execution time = 55.904923 seconds

Има още един бъг, за B трябва да сетваш и използваш temp2, a не temp.

Иначе аз пробвах да направя column да се алокира в “стека” като махнах local, но като че ли стана още по-бавно от твоето.

#12997 (ツ) gat3way
Създадено на 28.09.2020, видяно: 1549 пъти.

Не, не го прави това, двумерен масив в local, индексиран спрямо workitem-а е по-добре. Иначе компилатора си решава къде ще отиде това в "стека" и това спокойно може да свърши в глобалната памет. В opencl-а е много важно да се казва в кой memory space кое ходи, иначе компилатора прави каквото си реши и това често изобщо не е умно. Той може и много умни неща да направи, ако например масивът е малък и има наклонности, може да отиде в регистри, но (поне при AMD навремето), тея неща често ходеха в това дето му викаха "scratch register space" или нещо от сорта което си е глобална памет пар екселанс.

Нещо от сорта:

__kernel void levenshtein(__global const uchar *A, __global const uchar *B, __global uint *C, uint index) 
{
    uint s1len, s2len, x, y, lastdiag, olddiag;
    uint temp, temp2;
    uint i = get_local_id(0);
    __local uint column[32][MAX_STRING];

    int j = get_global_id(0);
    temp = j * (MAX_STRING+sizeof(uint));
    s1len = (A[temp+3] << 24) | (A[temp+2] << 16) | (A[temp+1] << 8) | A[temp];
    temp = index * (MAX_STRING+sizeof(uint));
    s2len = (B[temp+3] << 24) | (B[temp+2] << 16) | (B[temp+1] << 8) | B[temp];


    for (y = 1; y <= s1len; y++)
        column[i][y] = y;
    for (x = 1; x <= s2len; x++) 
    {
        column[i][0] = x;
        for (y = 1, lastdiag = x-1; y <= s1len; y++) 
        {
            olddiag = column[i][y];
            column[i][y] = MIN3(column[i][y] + 1, column[i][y-1] + 1, lastdiag + (A[temp + sizeof(int) + y-1] == B[temp + sizeof(int) + 
x - 1] ? 0 : 1));
            lastdiag = olddiag;
        }
    }

    C[i] = column[i][s1len];
}

И в хост кода size_t local_item_size става 32 (но за reduce, трябва да се въведе друг local_item_size дето пак си е 64).

#13000 (ツ) |
Създадено на 28.09.2020, видяно: 1543 пъти.
gat3way

Не, не го прави това, двумерен масив в local, индексиран спрямо workitem-а е по-добре. Иначе компилатора си решава къде ще отиде това в "стека" и това спокойно може да свърши в глобалната памет. В opencl-а е много важно да се казва в кой memory space кое ходи, иначе компилатора прави каквото си реши и това често изобщо не е умно. Той може и много умни неща да направи, ако например масивът е малък и има наклонности, може да отиде в регистри, но (поне при AMD навремето), тея неща често ходеха в това дето му викаха "scratch register space" или нещо от сорта което си е глобална памет пар екселанс.

Нещо от сорта:

__kernel void levenshtein(__global const uchar *A, __global const uchar *B, __global uint *C, uint index) 
{
    uint s1len, s2len, x, y, lastdiag, olddiag;
    uint temp, temp2;
    uint i = get_local_id(0);
    __local uint column[32][MAX_STRING];

    int j = get_global_id(0);
    temp = j * (MAX_STRING+sizeof(uint));
    s1len = (A[temp+3] << 24) | (A[temp+2] << 16) | (A[temp+1] << 8) | A[temp];
    temp = index * (MAX_STRING+sizeof(uint));
    s2len = (B[temp+3] << 24) | (B[temp+2] << 16) | (B[temp+1] << 8) | B[temp];


    for (y = 1; y <= s1len; y++)
        column[i][y] = y;
    for (x = 1; x <= s2len; x++) 
    {
        column[i][0] = x;
        for (y = 1, lastdiag = x-1; y <= s1len; y++) 
        {
            olddiag = column[i][y];
            column[i][y] = MIN3(column[i][y] + 1, column[i][y-1] + 1, lastdiag + (A[temp + sizeof(int) + y-1] == B[temp + sizeof(int) + 
x - 1] ? 0 : 1));
            lastdiag = olddiag;
        }
    }

    C[i] = column[i][s1len];
}

И в хост кода size_t local_item_size става 32 (но за reduce, трябва да се въведе друг local_item_size дето пак си е 64).

Да, аз пробвах да го направя и uchar, но е по-бавно от свалянето на броя тредове. На теслата е 24 секунди за 1000 стринга.

#13001 (ツ) gat3way
Създадено на 28.09.2020, видяно: 1541 пъти.

Ами аз както казах, доскоро не бях писал кернел от поне две години :)

Но иначе, навремето (и за амд и за нвидия), гранулярността за достъп до всякаква памет беше 32 бита. Като достъпваш uchar, пак четеш 32 бита, но имплицитно се шибват маските и това е по-бавно (отделно и там конфликтите при достъпа на памет когато се пише, които са сходни като идея с false sharing-а при CPU-то). Но да де, когато имаш ограничения на паметта, това може и да излезе по-добрият вариант. То затва ми липсва amdappprofiler-а, беше много хубаво нещо. Със сигурност има нещо и за nvidia сега, тогава нямаше толкова яко, но нвидията много дръпна последните години и сигурно имат вече по-добро. То всъщност май и тогава имаха, но беше само за куда-та май. Абе всъщност не знам.

#13003 (ツ) |
Създадено на 28.09.2020, видяно: 1519 пъти.
gat3way

Ами аз както казах, доскоро не бях писал кернел от поне две години :)

Но иначе, навремето (и за амд и за нвидия), гранулярността за достъп до всякаква памет беше 32 бита. Като достъпваш uchar, пак четеш 32 бита, но имплицитно се шибват маските и това е по-бавно (отделно и там конфликтите при достъпа на памет когато се пише, които са сходни като идея с false sharing-а при CPU-то). Но да де, когато имаш ограничения на паметта, това може и да излезе по-добрият вариант. То затва ми липсва amdappprofiler-а, беше много хубаво нещо. Със сигурност има нещо и за nvidia сега, тогава нямаше толкова яко, но нвидията много дръпна последните години и сигурно имат вече по-добро. То всъщност май и тогава имаха, но беше само за куда-та май. Абе всъщност не знам.

На мен твоя код ми харесва повече от моя, но за съжаление не скейлва добре. 200К стринга ги сравнява за 4688 секунди, докато моя го прави за 1335.

От друга страна, за 1000 стринга твоя му трябват 23.9 секунди, докато на моя 118.3.

Предполагам, че твоите кърнъли са too fine grained и овърхеда за стартирането се акумулира.

#13004 (ツ) gat3way
Създадено на 28.09.2020, видяно: 1517 пъти.
#13010 (ツ) Golden Gega
Създадено на 29.09.2020, видяно: 1502 пъти.
#13012 (ツ) Delegate
Последно редактирано на 29.09.2020 от Delegate, видяно: 1489 пъти.
ФейкПрофил
Delegate

Малко я барнах заявката - чувствително подобрение. Не ползвам сторнати процедури.

Successfully run. Total query runtime: 5 min 41 secs.
100 rows affected.

Абе нещо пак не е така. Задачата е за всеки от 100-те хиляди стринга да намериш най-близкия от другата колекция. С други думи, решението ти трябва да произведе 100_000 реда на изхода.

ПП: задачата много прилича на spellchecker обаче там алгоритмите се издъниха здраво :Д Както казва мастър Линус - когато теорията и практиката се сблъскат, практиката винаги печели.

Еми, на Джони програмката например вади 100 резултата, не 100 000. От друга страна, оригиналното условие на | е :

Колекция А е от 1 милион стринга, колекция Б е от 10 милиона стринга.

За всеки стринг от колекция Б трябва да се намери стринга от колекция А, който е най-близък до него.

Джон, ти защо вадиш само 100 резултата и при това редицата ти от резултати 4,4,4,4,40,4,9 ти е одобренa от | :-)

#13017 (ツ) ФейкПрофил
Създадено на 29.09.2020, видяно: 1483 пъти.

Тази редица е одобрена. Май оригиналното условие просто е наобратно - за всеки от 100те трябва да се намери най-близкия от 100–000

#13029 (ツ) Delegate
Създадено на 29.09.2020, видяно: 1465 пъти.

Някой, ако иска да тества C# версия (exe)

Attached files:
FileSizeUploadedDownloadsMD5 hash
leven.zip3458656 bytes29.09.2020105d591d3fb2c55246916365eb7a955b943
#13031 (ツ) Courvoisier
Създадено на 29.09.2020, видяно: 1461 пъти.
Delegate

Някой, ако иска да тества C# версия (exe)

Дай сорс де, exe да изпълнявам от интернета малко....

#13032 (ツ) synergie
Създадено на 29.09.2020, видяно: 1460 пъти.
Courvoisier
Delegate

Някой, ако иска да тества C# версия (exe)

Дай сорс де, exe да изпълнявам от интернета малко....

Пусни го у некой контейнер.

#13033 (ツ) Delegate
Последно редактирано на 29.09.2020 от Delegate, видяно: 1456 пъти.

А на Джони изпълнихте екзето, а ? rofl

using System;
using System.Collections.Generic;
using System.Text;
using System.IO;

namespace Leven
{
    class Program
    {
        public static int Distance(string value1, string value2)
        {
            if (value1.Length == 0)
            {
                return 0;
            }
            int[] costs = new int[value1.Length];
            // Add indexing for insertion to first row
            for (int i = 0; i < costs.Length; )
            {
                costs[i] = ++i;
            }

            int minSize = value1.Length < value2.Length ? value1.Length : value2.Length;
            for (int i = 0; i < minSize; i++)
            {
                // cost of the first index
                int cost = i;
                int addationCost = i;

                // cache value for inner loop to avoid index lookup and bonds checking, profiled this is quicker
                char value2Char = value2[i];

                for (int j = 0; j < value1.Length; j++)
                {
                    int insertionCost = cost;
                    cost = addationCost;
                    // assigning this here reduces the array reads we do, improvment of the old version
                    addationCost = costs[j];
                    if (value2Char != value1[j])
                    {
                        if (insertionCost < cost)
                        {
                            cost = insertionCost;
                        }
                        if (addationCost < cost)
                        {
                            cost = addationCost;
                        }
                        ++cost;
                    }
                    costs[j] = cost;
                }
            }
            return costs[costs.Length - 1];
        }

        static void Main(string[] args)
        {

            var ListA = File.ReadAllLines("Dataset.txt");
            var ListB = File.ReadAllLines("t2.txt");
            List<string> lineListB = new List<string>(ListB);
            List<string> lineListA = new List<string>(ListA);

            //var watch = System.Diagnostics.Stopwatch.StartNew();

            for (int i = 0; i < lineListB.Count; i++)
            {
                var watch = System.Diagnostics.Stopwatch.StartNew();
                int min_dist = Int32.MaxValue;

                for (int j = 0; j < lineListA.Count; j++)
                {
                    int dist = Distance(lineListB[i], lineListA[j]);
                    if (dist < min_dist)
                    {
                        min_dist = dist;
                    }
                }

                watch.Stop();
                Console.WriteLine(i + " elapsed (ms)/dist:" + watch.ElapsedMilliseconds + "/" + min_dist);

            }


        }

    }
}



0 1 2 3 4 ....8 9 10 11 12 ....18 19 20 21 22 ....26 27 28 29 30 ....32 33 34 35 36

Задача НЕ за интервю
7

AsmBB v3.0 (check-in: a316dab8b98d07d9); SQLite v3.42.0 (check-in: 831d0fb2836b71c9);
©2016..2023 John Found; Licensed under EUPL. Powered by Assembly language Created with Fresh IDE