Skip to content

Commit

Permalink
Test on all possible CLContext's
Browse files Browse the repository at this point in the history
  • Loading branch information
@SunSerega
SunSerega committed Oct 23, 2023
1 parent 07188da commit b37e4e0
Show file tree
Hide file tree
Showing 12 changed files with 689 additions and 460 deletions.
295 changes: 205 additions & 90 deletions LastPack.log
View file

Large diffs are not rendered by default.

8 changes: 4 additions & 4 deletions Log/Test.log
View file
Original file line number Diff line number Diff line change
Expand Up @@ -11,9 +11,11 @@ Tester: Compiling "Tests/DebugPCU/OpenCLABC.pas"
Tester: Compiling: OK
Tester: Compiling "Tests/DebugPCU/OpenGLABC.pas"
Tester: Compiling: OK
Tester: Compiling "Tests/Exec/CL/ToString.pas"
Tester: Compiling "Tests/CLContextGen.pas"
Tester: Compiling: OK
Tester: Compiling "Tests/Exec/CLABC/gen def context.pas"
Tester: Running Tests/CLContextGen.exe
Tester: Finished running Tests/CLContextGen.exe
Tester: Compiling "Tests/Exec/CL/ToString.pas"
Tester: Compiling: OK
Tester: Compiling "Tests/Exec/CLABC/01#Wrap/01#CLPlatform/1.pas"
Tester: Compiling: OK
Expand Down Expand Up @@ -221,8 +223,6 @@ Tester: Compiling "Samples/OpenGLABC/Точки на поле/Точки.pas"
Tester: Compiling: OK
Tester: Executing Test[Tests/Exec/CL/ToString]
Tester: Done executing
Tester: Executing Test[Tests/Exec/CLABC/gen def context]
Tester: Done executing
Tester: Executing Test[Tests/Exec/CLABC/01#Wrap/01#CLPlatform/1]
Tester: Done executing
Tester: Executing Test[Tests/Exec/CLABC/01#Wrap/02#CLDevice/1]
Expand Down
143 changes: 69 additions & 74 deletions Modules.Packed/OpenCLABC.pas
View file
Original file line number Diff line number Diff line change
Expand Up @@ -3261,7 +3261,9 @@ CLArrayProperties = class

end;

public static procedure GenerateAndCheckDefault(test_size: integer := 1024*24; test_max_seconds: real := 0.5);
public function TestSanity(test_size: integer := 1024*24): TimeSpan?;

public static function GenerateAndCheckAllPossible(test_size: integer := 1024*24; test_max_seconds: real := 0.5): List<ValueTuple<CLContext,TimeSpan>>;

private static function LoadTestContext: CLContext;

Expand Down Expand Up @@ -35370,86 +35372,84 @@ function CombineUseAsyncQueueN7<TInp1,TInp2,TInp3,TInp4,TInp5,TInp6,TInp7>(use:

{$region CLContext}

static procedure CLContext.GenerateAndCheckDefault(test_size: integer; test_max_seconds: real);
type
CLContextTestData = static class

private const test_prog = 'kernel void k(global int* v) { v[get_global_id(0)]++; }';

public static function MakeTestQueue(test_size: integer): CommandQueue<boolean>;
begin
// Без кеширования, чтобы у всех
// запусков .TestSanity были равные условия
// что касается внутренних оптимизаций очереди
var rng := new Random;
var test_arr := ArrGen(test_size, i->rng.Next);
var Q_Arr := HFQ(c->new CLArray<integer>(c, test_size), false).Multiusable;
Result :=
HFQ(c->CLProgramCode.Create(test_prog, c)['k'], false).MakeCCQ
.ThenExec1(test_size, Q_Arr.MakeCCQ
.ThenWriteArray(test_arr)
) +
Q_Arr.MakeCCQ
.ThenGetArray
.ThenConvert(test_res->
test_res.Zip(test_arr, (a,b)->a=b+1).All(r->r), false
);
end;

end;

function CLContext.TestSanity(test_size: integer): TimeSpan?;
begin
if Interlocked.CompareExchange(default_was_inited, 1, 0)<>0 then
raise new System.InvalidOperationException($'%CLContext:GenerateAndCheckDefault:NotFirst%');
var Q := CLContextTestData.MakeTestQueue(test_size);
var sw := Stopwatch.StartNew;

Result := nil;
if not self.SyncInvoke(Q) then exit;

var test_prog := 'kernel void k(global int* v) { v[get_global_id(0)]++; }';
Result := sw.Elapsed;
end;

static function CLContext.GenerateAndCheckAllPossible(test_size: integer; test_max_seconds: real): List<ValueTuple<CLContext,TimeSpan>>;
begin
var test_max_time := TimeSpan.FromSeconds(test_max_seconds);
Result := new List<ValueTuple<CLContext, TimeSpan>>;

var Q_Test: CommandQueue<boolean>;
var P_Arr := new ParameterQueue<CLArray<integer>>('A');
var P_Prog := new ParameterQueue<CLProgramCode>('Prog');
begin
var rng := new Random;
var test_arr := ArrGen(test_size, i->rng.Next);
Q_Test :=
CLKernelCCQ.Create(P_Prog.ThenConvert(p->p['k'], false, true))
.ThenExec1(test_size, P_Arr.MakeCCQ
.ThenWriteArray(test_arr)
) +
P_Arr.MakeCCQ
.ThenGetArray
.ThenConvert(test_res->
test_res.Zip(test_arr, (a,b)->a=b+1).All(r->r), false
);
end;

var c := CLPlatform.All.SelectMany(pl->
foreach var pl in CLPlatform.All do
begin
var dvcs := CLDevice.GetAllFor(pl, clDeviceType.DEVICE_TYPE_ALL).ToList;

System.Threading.Tasks.Parallel.For(0,dvcs.Count, i->
var keep := new boolean[dvcs.Count];
var thrs := ArrGen(dvcs.Count, i->new Thread(()->
begin
var del := true;

var c := new CLContext(dvcs[i]);
var S_Arr := P_Arr.NewSetter(new CLArray<integer>(c, test_size));
var S_Prog := P_Prog.NewSetter(new CLProgramCode(test_prog, c));

var thr := new Thread(()->
begin
del := not c.SyncInvoke(Q_Test, S_Arr, S_Prog);
end);
thr.IsBackground := true;
thr.Start;
keep[i] := nil <> CLContext.Create(dvcs[i]).TestSanity(test_size)
end));
foreach var thr in thrs do thr.IsBackground := true;
foreach var thr in thrs do thr.Start;
Thread.Sleep(test_max_time);
foreach var thr in thrs do thr.Abort;

for var i := 0 to dvcs.Count-1 do
if not keep[i] then dvcs[i] := nil;
dvcs.RemoveAll(d->d=nil);

var chosen := new List<CLDevice>(dvcs.Count);
foreach var choise in |true,false|.CartesianPower(dvcs.Count) do
begin
chosen.Clear;
for var i := 0 to dvcs.Count-1 do
if choise[i] then chosen += dvcs[i];
if chosen.Count=0 then continue;

if not thr.Join(test_max_time) then
thr.Abort;
var c := new CLContext(chosen, chosen[0]);
var time := c.TestSanity(test_size);
if time=nil then continue;

if del then dvcs[i] := nil;
end);
dvcs.RemoveAll(d->d=nil);
Result.Add(ValueTuple.Create(c, time.Value));
end;

Result := if dvcs.Count=0 then
System.Linq.Enumerable.Empty&<(CLContext,TimeSpan)> else
|true,false|.Cartesian(dvcs.Count)
.Select(choise->
begin
Result := new List<CLDevice>(dvcs.Count);
for var i := 0 to dvcs.Count-1 do
if choise[i] then Result += dvcs[i];
end)
.Where(l_dvcs->l_dvcs.Count<>0)
.Select(l_dvcs->
begin
var c := new CLContext(l_dvcs, l_dvcs[0]);

var sw := Stopwatch.StartNew;
c.SyncInvoke(Q_Test.DiscardResult,
P_Arr.NewSetter(new CLArray<integer>(c, test_size)),
P_Prog.NewSetter(new CLProgramCode(test_prog, c))
);
sw.Stop;

Result := (c, sw.Elapsed);
end);
end)
.DefaultIfEmpty((default(CLContext),TimeSpan.Zero))
.MinBy(t->t[1])[0];
end;

Interlocked.CompareExchange(_default, c, nil);
end;

static function CLContext.LoadTestContext: CLContext;
Expand All @@ -35469,11 +35469,6 @@ function CombineUseAsyncQueueN7<TInp1,TInp2,TInp3,TInp4,TInp5,TInp6,TInp7>(use:
var br := new System.IO.BinaryReader(System.IO.File.OpenRead(fname));
try
var pl_name := br.ReadString;
if pl_name.Length=0 then
begin
$'Pregenerated context was empty'.Println;
exit;
end;

var pl := CLPlatform.All.SingleOrDefault(pl->pl.Properties.Name=pl_name);
if pl=nil then
Expand Down
143 changes: 69 additions & 74 deletions Modules/OpenCLABC.pas
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1428,7 +1428,9 @@ EventRetainReleaseData = record

end;

public static procedure GenerateAndCheckDefault(test_size: integer := 1024*24; test_max_seconds: real := 0.5);
public function TestSanity(test_size: integer := 1024*24): TimeSpan?;

public static function GenerateAndCheckAllPossible(test_size: integer := 1024*24; test_max_seconds: real := 0.5): List<ValueTuple<CLContext,TimeSpan>>;

private static function LoadTestContext: CLContext;

Expand Down Expand Up @@ -9968,86 +9970,84 @@ function HPQ(p: CLContext->(); need_own_thread: boolean) :=

{$region CLContext}

static procedure CLContext.GenerateAndCheckDefault(test_size: integer; test_max_seconds: real);
type
CLContextTestData = static class

private const test_prog = 'kernel void k(global int* v) { v[get_global_id(0)]++; }';

public static function MakeTestQueue(test_size: integer): CommandQueue<boolean>;
begin
// Без кеширования, чтобы у всех
// запусков .TestSanity были равные условия
// что касается внутренних оптимизаций очереди
var rng := new Random;
var test_arr := ArrGen(test_size, i->rng.Next);
{%> var Q_Arr := HFQ(c->new CLArray<integer>(c, test_size), false).Multiusable;%}
{%> Result :=%}
{%> HFQ(c->CLProgramCode.Create(test_prog, c)['k'], false).MakeCCQ%}
{%> .ThenExec1(test_size, Q_Arr.MakeCCQ%}
{%> .ThenWriteArray(test_arr)%}
{%> ) +%}
{%> Q_Arr.MakeCCQ%}
{%> .ThenGetArray%}
{%> .ThenConvert(test_res->%}
{%> test_res.Zip(test_arr, (a,b)->a=b+1).All(r->r), false%}
{%> )!!}test_arr := test_arr; Result := nil{%};
end;

end;

function CLContext.TestSanity(test_size: integer): TimeSpan?;
begin
if Interlocked.CompareExchange(default_was_inited, 1, 0)<>0 then
raise new System.InvalidOperationException($'%CLContext:GenerateAndCheckDefault:NotFirst%');
var Q := CLContextTestData.MakeTestQueue(test_size);
var sw := Stopwatch.StartNew;

Result := nil;
if not self.SyncInvoke(Q) then exit;

var test_prog := 'kernel void k(global int* v) { v[get_global_id(0)]++; }';
Result := sw.Elapsed;
end;

static function CLContext.GenerateAndCheckAllPossible(test_size: integer; test_max_seconds: real): List<ValueTuple<CLContext,TimeSpan>>;
begin
var test_max_time := TimeSpan.FromSeconds(test_max_seconds);
Result := new List<ValueTuple<CLContext, TimeSpan>>;

var Q_Test: CommandQueue<boolean>;
var P_Arr := new ParameterQueue<CLArray<integer>>('A');
var P_Prog := new ParameterQueue<CLProgramCode>('Prog');
begin
var rng := new Random;
var test_arr := ArrGen(test_size, i->rng.Next);
{%> Q_Test :=%}
{%> CLKernelCCQ.Create(P_Prog.ThenConvert(p->p['k'], false, true))%}
{%> .ThenExec1(test_size, P_Arr.MakeCCQ%}
{%> .ThenWriteArray(test_arr)%}
{%> ) +%}
{%> P_Arr.MakeCCQ%}
{%> .ThenGetArray%}
{%> .ThenConvert(test_res->%}
{%> test_res.Zip(test_arr, (a,b)->a=b+1).All(r->r), false%}
{%> )!!}test_arr := test_arr{%};
end;

var c := CLPlatform.All.SelectMany(pl->
foreach var pl in CLPlatform.All do
begin
var dvcs := CLDevice.GetAllFor(pl, clDeviceType.DEVICE_TYPE_ALL).ToList;

System.Threading.Tasks.Parallel.For(0,dvcs.Count, i->
var keep := new boolean[dvcs.Count];
var thrs := ArrGen(dvcs.Count, i->new Thread(()->
begin
var del := true;

var c := new CLContext(dvcs[i]);
var S_Arr := P_Arr.NewSetter(new CLArray<integer>(c, test_size));
var S_Prog := P_Prog.NewSetter(new CLProgramCode(test_prog, c));

var thr := new Thread(()->
begin
del := not c.SyncInvoke(Q_Test, S_Arr, S_Prog);
end);
thr.IsBackground := true;
thr.Start;
keep[i] := nil <> CLContext.Create(dvcs[i]).TestSanity(test_size)
end));
foreach var thr in thrs do thr.IsBackground := true;
foreach var thr in thrs do thr.Start;
Thread.Sleep(test_max_time);
foreach var thr in thrs do thr.Abort;

for var i := 0 to dvcs.Count-1 do
if not keep[i] then dvcs[i] := nil;
dvcs.RemoveAll(d->d=nil);

var chosen := new List<CLDevice>(dvcs.Count);
foreach var choise in |true,false|.CartesianPower(dvcs.Count) do
begin
chosen.Clear;
for var i := 0 to dvcs.Count-1 do
if choise[i] then chosen += dvcs[i];
if chosen.Count=0 then continue;

if not thr.Join(test_max_time) then
thr.Abort;
var c := new CLContext(chosen, chosen[0]);
var time := c.TestSanity(test_size);
if time=nil then continue;

if del then dvcs[i] := nil;
end);
dvcs.RemoveAll(d->d=nil);
Result.Add(ValueTuple.Create(c, time.Value));
end;

Result := if dvcs.Count=0 then
System.Linq.Enumerable.Empty&<(CLContext,TimeSpan)> else
|true,false|.Cartesian(dvcs.Count)
.Select(choise->
begin
Result := new List<CLDevice>(dvcs.Count);
for var i := 0 to dvcs.Count-1 do
if choise[i] then Result += dvcs[i];
end)
.Where(l_dvcs->l_dvcs.Count<>0)
.Select(l_dvcs->
begin
var c := new CLContext(l_dvcs, l_dvcs[0]);

var sw := Stopwatch.StartNew;
c.SyncInvoke(Q_Test.DiscardResult,
P_Arr.NewSetter(new CLArray<integer>(c, test_size)),
P_Prog.NewSetter(new CLProgramCode(test_prog, c))
);
sw.Stop;

Result := (c, sw.Elapsed);
end);
end)
.DefaultIfEmpty((default(CLContext),TimeSpan.Zero))
.MinBy(t->t[1])[0];
end;

Interlocked.CompareExchange(_default, c, nil);
end;

static function CLContext.LoadTestContext: CLContext;
Expand All @@ -10067,11 +10067,6 @@ function HPQ(p: CLContext->(); need_own_thread: boolean) :=
var br := new System.IO.BinaryReader(System.IO.File.OpenRead(fname));
try
var pl_name := br.ReadString;
if pl_name.Length=0 then
begin
$'Pregenerated context was empty'.Println;
exit;
end;

var pl := CLPlatform.All.SingleOrDefault(pl->{%>pl.Properties.Name!!}nil{%}=pl_name);
if pl=nil then
Expand Down
2 changes: 0 additions & 2 deletions Packing/Descriptions/OpenCLABC.missing.log
View file
Original file line number Diff line number Diff line change
Expand Up @@ -796,5 +796,3 @@

# Err:Blittable:Source:CLMemory:FillNativeArrayArea

# CLContext:GenerateAndCheckDefault:NotFirst

Loading

0 comments on commit b37e4e0

Please sign in to comment.