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
Е, скалирането е друг проблем, но като цяло проблемът с profile-ването и оптимизациите ако и да звучи странно, е доста по-прост върху GPU-та, то там всичко е по-просто. Като цяло, GPU кода е далеч по-благ просто защото обикновено не може да става и дума за сравнение като комплексност с CPU кода. Съответно много по-лесно се осмисля какво и как да се промени. Това хубаво не е тривиално все пак наистина, щото си има неговите номера, наложени от хардуера, но всичко е (беше поне) доста добре документирано, и AMD и NVidia имаха мноооого чудни документи по въпроса за писането на читави кернели и утилизиране на хардуера доколкото може. Дори, колкото и да е тъпо, към един момент, дизасемблираните кернели за AMD ми се виждаха по-четливи и разбираеми от x86 асемблера, въпреки че това беше отпреди GCN, с оная малоумна VLIW архитектура, с клаузите, слотовете по клаузите и прочее тъпотии. С GCN (а и при nvidia) е още по-просто и четливо.
Мисля, че имаш грешка в т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 в твоя полза. 
Левенщайна ти е бавен :) аз исках да ускоря моя, но това което намерих като "оптимизиран" алгоритъм се оказа 2пъти по-бавен от библиотечката която ползвах :Д
Единствената оптимизация, която успях да направя е да ползвм u8 вместо char, което смъкна 500мс на стринг и стана ~3450мс
ФейкПрофил, и аз мисля, че му е бавен Левенщ*ейнът*. Тоя 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];
}
Хмм, това __local в OpenCL като __shared__ в CUDA ли е? Ако да, в levenshtein използваш един и същи буфер за всички тредове. Ако не е, res и off в reduce ще са грешни.
Оооо даммм, прав си. Требва да е двумерен масив и да се индексира спрямо 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, но като че ли стана още по-бавно от твоето.
Не, не го прави това, двумерен масив в 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 стринга.
Ами аз както казах, доскоро не бях писал кернел от поне две години :)
Но иначе, навремето (и за амд и за нвидия), гранулярността за достъп до всякаква памет беше 32 бита. Като достъпваш uchar, пак четеш 32 бита, но имплицитно се шибват маските и това е по-бавно (отделно и там конфликтите при достъпа на памет когато се пише, които са сходни като идея с false sharing-а при CPU-то). Но да де, когато имаш ограничения на паметта, това може и да излезе по-добрият вариант. То затва ми липсва amdappprofiler-а, беше много хубаво нещо. Със сигурност има нещо и за nvidia сега, тогава нямаше толкова яко, но нвидията много дръпна последните години и сигурно имат вече по-добро. То всъщност май и тогава имаха, но беше само за куда-та май. Абе всъщност не знам.
На мен твоя код ми харесва повече от моя, но за съжаление не скейлва добре. 200К стринга ги сравнява за 4688 секунди, докато моя го прави за 1335.
От друга страна, за 1000 стринга твоя му трябват 23.9 секунди, докато на моя 118.3.
Предполагам, че твоите кърнъли са too fine grained и овърхеда за стартирането се акумулира.
Чакай сега че съм на друга вълна, спуквам се от смях, тоя яко е огаврил алтернативките-десни тука.
Хахаха, тоя пич плюс цивренето на Рамбо дават добър старт на деня, а и спря да вали 
Еми, на Джони програмката например вади 100 резултата, не 100 000. От друга страна, оригиналното условие на | е :
Джон, ти защо вадиш само 100 резултата и при това редицата ти от резултати 4,4,4,4,40,4,9 ти е одобренa от | 
Тази редица е одобрена. Май оригиналното условие просто е наобратно - за всеки от 100те трябва да се намери най-близкия от 100–000
Някой, ако иска да тества C# версия (exe)
| File | Size | Uploaded | Downloads | MD5 hash |
|---|---|---|---|---|
| leven.zip | 3458656 bytes | 29.09.2020 | 145 | d591d3fb2c55246916365eb7a955b943 |
Дай сорс де, exe да изпълнявам от интернета малко....
А на Джони изпълнихте екзето, а ?
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