From 9b178235e54d19d4cbbe3319f081020b6ce76851 Mon Sep 17 00:00:00 2001 From: NHALX Date: Sun, 15 Dec 2013 19:29:13 +0100 Subject: [PATCH] Fixed bug in lattice linext mapping. - Added multi metric wrapper to ll_formulas - Added CPU version of build combo. - Misc code reorganization. --- projects/VisualStudio/llio.vcxproj | 11 +- projects/VisualStudio/llio.vcxproj.filters | 43 ++-- src/league/build_combo/build_combo.c | 185 +++++++----------- src/league/build_combo/build_combo.h | 14 +- src/league/build_combo/item_filter.c | 75 +++++++ src/league/build_combo/item_filter.h | 14 ++ src/league/build_combo/kernel/k_build_combo.c | 110 +++++++++++ src/league/build_combo/kernel/k_build_combo.h | 18 ++ .../unit_test/unittest_buildcombo.c | 36 ++-- src/league/build_path/db_input.c | 21 -- .../kernel/{metric_ADPS.c => metric_area.c} | 21 +- .../kernel/{metric_ADPS.h => metric_area.h} | 9 +- src/league/build_path/opencl_bind.c | 129 +----------- src/league/build_path/opencl_bind.h | 2 +- .../build_path/unit_test/unittest_buildpath.c | 11 +- .../build_path/unit_test/unittest_kernel.cl | 2 +- .../build_path/unit_test/unittest_opencl.c | 1 + src/league/database/db_search.c | 26 +++ src/league/database/db_search.h | 7 + src/league/ll_formulas.h | 13 +- src/league/main.c | 69 +++++++ src/league/unit_test/unittest_league.c | 3 +- src/opencl_host/function.c | 114 +++++++++++ src/opencl_host/function.h | 16 +- src/opencl_host/host.c | 3 + src/poset/lattice.c | 3 +- 26 files changed, 625 insertions(+), 331 deletions(-) create mode 100644 src/league/build_combo/item_filter.c create mode 100644 src/league/build_combo/item_filter.h create mode 100644 src/league/build_combo/kernel/k_build_combo.c create mode 100644 src/league/build_combo/kernel/k_build_combo.h rename src/league/build_path/kernel/{metric_ADPS.c => metric_area.c} (91%) rename src/league/build_path/kernel/{metric_ADPS.h => metric_area.h} (68%) create mode 100644 src/league/database/db_search.c create mode 100644 src/league/database/db_search.h create mode 100644 src/league/main.c diff --git a/projects/VisualStudio/llio.vcxproj b/projects/VisualStudio/llio.vcxproj index 9e8d7f4..92eaee4 100644 --- a/projects/VisualStudio/llio.vcxproj +++ b/projects/VisualStudio/llio.vcxproj @@ -14,13 +14,16 @@ + + - + + @@ -45,15 +48,19 @@ + + - + + + diff --git a/projects/VisualStudio/llio.vcxproj.filters b/projects/VisualStudio/llio.vcxproj.filters index f022d7f..54b86c7 100644 --- a/projects/VisualStudio/llio.vcxproj.filters +++ b/projects/VisualStudio/llio.vcxproj.filters @@ -66,9 +66,6 @@ league\build_path - - league\build_path - league\build_path @@ -81,9 +78,6 @@ combinatorics - - league\build_combo - combinatorics @@ -102,6 +96,24 @@ league + + league\build_combo + + + league + + + league\build_combo + + + league\database + + + league\build_combo + + + league\build_path + @@ -146,9 +158,6 @@ league\build_path - - league\build_path - league\build_path @@ -176,11 +185,23 @@ opencl_host + + league + + + league\build_combo + league\build_combo - - league + + league\database + + + league\build_combo + + + league\build_path diff --git a/src/league/build_combo/build_combo.c b/src/league/build_combo/build_combo.c index ff37585..306bc16 100644 --- a/src/league/build_combo/build_combo.c +++ b/src/league/build_combo/build_combo.c @@ -1,128 +1,89 @@ -#ifndef __OPENCL_VERSION__ -#include "../../opencl_host/dummy.h" -#include -#undef DEBUG_PRINTF -#define DEBUG_PRINTF +#ifndef _MSC_VER +#include #else -#define DEBUG_PRINTF +#define alloca _alloca #endif +#include -#include "types.h" -#include "league/database/db_layout.h" -#include "league/ll_formulas.h" -#include "league/common.h" +#include "league/build_combo/build_combo.h" #include "combinatorics/combination.h" +#include "combinatorics/factorial.h" +#include "league/unit_test/find_max.h" // TODO: remove this +#include "league/build_combo/kernel/k_build_combo.h" -static -int -metric(__constant item_t db_items[], - llf_criteria *cfg, - setmax_t *combo, - uint_t combo_n, - __local setmax_t *unique, - float *output) +void build_comboGPU(opencl_function *func, llf_criteria *cfg, const item_t *items, size_t item_n, size_t inventory_n, setmax_t *combo) { - size_t i; - - VECTOR(stats) = VECTOR_ZERO_INIT; - - for (i = 0; i < combo_n; ++i) - { - itemid_t item = combo[i]; - itemid_t passive = db_items[item].passive_id; - - stats_add(db_items, &stats, item, ++unique[passive] == 1); - } - - *output = llf_dmgtotal(cfg, &stats); - return 0; -} + ulong_t ncombo = NK_MULTISET(item_n, inventory_n); + opencl_workset workset; + opencl_kernel_arg *output; + result_t *results, local_best, best = { 0, 0 }; -////////////////////////////////////////////////////////////// + workset = opencl_workcfg(func->ctx, ncombo, build_combo__allocnfo__(func, item_n)); + output = build_combo__bind__(func, inventory_n, items, item_n, cfg, workset.pass_size, workset.local_size); + results = output->buf_data; + for (ulong_t i = 0, offset = 0; i < workset.iterations; ++i) + { + opencl_run(func, 1, TRUE, offset, &workset); + local_best = FindMax(results, output->buf_size / sizeof *results); + + if (best.metric < local_best.metric) + best = local_best; -result_t -BuildCombo(__constant item_t db_items[], - llf_criteria *cfg, - ulong_t rank, - uint_t db_len, - uint_t combo_len, - __local setmax_t *unique) -{ - setmax_t combo[NK_MULTISET_MAX_K]; - result_t dps; - - combo_unrank(rank, db_len, combo_len, combo); - dps.index = rank; - metric(db_items, cfg, combo, combo_len, unique, &dps.metric); - return dps; + offset += workset.pass_size; + } + + combo_unrank(best.index, item_n, inventory_n, combo); } -#ifndef __OPENCL_VERSION__ -#include "opencl_host/function.h" -opencl_allocinfo -build_combo__allocnfo__(opencl_function *func, size_t items_n) -{ - cl_ulong klmem; - opencl_allocinfo nfo = { 0 }; - - NOFAIL(clGetKernelWorkGroupInfo(func->kernel, func->ctx->device, - CL_KERNEL_LOCAL_MEM_SIZE, sizeof klmem, &klmem, 0)); - /*TODO: fill these out - nfo.fixed.constant += sizeof (item_t)* items_n; - - nfo.scale_workgroup.local += sizeof (ideal_t)* items_n; - nfo.scale_workgroup.local += sizeof (result_t); - nfo.scale_workgroup.local += klmem; - - nfo.scale_reduce.global = sizeof (result_t);*/ - return nfo; -} - -opencl_kernel_arg * -build_combo__bind__(opencl_function *X, - uint_t combo_n, - item_t *items, uint_t items_n, - llf_criteria *cfg, - - size_t pass_size, - size_t local_size) +void build_comboCPU(opencl_function *func, llf_criteria *cfg, const item_t *items, size_t item_n, size_t inventory_n, setmax_t *combo) { - size_t outlen; - - ka_ignore(X); - ka_mconst(X, "db_items", 0, items, sizeof *items * items_n); - ka_value(X, "cfg_input", cfg, sizeof *cfg); - ka_value(X, "db_len", &items_n, sizeof items_n); - ka_value(X, "combo_len", &combo_n, sizeof combo_n); - ka_mlocal(X, "pasv_scratch", sizeof (setmax_t) * items_n * local_size); - ka_mlocal(X, "scratch", sizeof (result_t) * local_size); - - outlen = pass_size / local_size; - return ka_mglobal(X, "output", A_OUT, CL_MEM_WRITE_ONLY, 0, sizeof(result_t)* outlen); + ulong_t ncombo = NK_MULTISET(item_n, inventory_n); + ulong_t pass_size; + ulong_t iterations; + ulong_t offset = 0; + result_t result = { 0,0 }; + + pass_size = 1000000; // TODO: is this suffeciently large to + + if (pass_size > ncombo) + pass_size = (size_t)ncombo; + + iterations = ncombo / pass_size; + + for (ulong_t j = 0; j < iterations; ++j) + { + int i; + result_t local_best; + setmax_t *scratch; + result_t r; + + #pragma omp parallel default(shared) private(i,r,local_best,scratch) + { + local_best = (result_t){ 0, 0 }; + scratch = alloca(sizeof *scratch * item_n); + + #pragma omp for schedule(static) nowait + for (i = 0; i < pass_size; ++i) + { + r = build_combo_m(items, cfg, offset + i, item_n, inventory_n, scratch); + if (local_best.metric < r.metric) + local_best = r; + } + + #pragma omp critical + { + if (result.metric < local_best.metric) + result = local_best; + } + } + + offset += pass_size; + printf("CPU: progress: %llu/%llu (%f)\n", j+1, iterations, (float)(j+1) / (float)iterations); + } + + combo_unrank(result.index, item_n, inventory_n, combo); } - -#else -#include "opencl_host/kernel/reduce.cl" -#include "combinatorics/combination.c" -#include "combinatorics/factorial.c" - -__kernel void -build_combo( - ulong_t start_offset, - __constant item_t db_items[], - llf_criteria cfg, - uint_t db_len, - uint_t combo_len, - __local setmax_t *unique, - __local result_t *scratch, - __global result_t *result) -{ - result_t dps = BuildCombo(db_items, &cfg, start_offset + get_global_id(0), db_len, combo_len, unique); - Reduce(dps, scratch, result); -} - -#endif diff --git a/src/league/build_combo/build_combo.h b/src/league/build_combo/build_combo.h index df9cf7c..d4366eb 100644 --- a/src/league/build_combo/build_combo.h +++ b/src/league/build_combo/build_combo.h @@ -1,15 +1,11 @@ #ifndef _LEAGUE_BUILD_COMBO_H_ #define _LEAGUE_BUILD_COMBO_H_ - -#ifndef __OPENCL_VERSION__ - -#include "opencl_host/function.h" -#include "league/database/db_layout.h" -#include "league/ll_formulas.h" #include "league/common.h" +#include "league/ll_formulas.h" +#include "opencl_host/function.h" +#include "combinatorics/factorial.h" -opencl_allocinfo build_combo__allocnfo__(opencl_function *func, size_t items_n); -opencl_kernel_arg *build_combo__bind__(opencl_function *X, uint_t combo_n, const item_t *items, uint_t items_n, llf_criteria *cfg, size_t pass_size, size_t local_size); -#endif +void build_comboGPU(opencl_function *func, llf_criteria *cfg, const item_t *items, size_t item_n, size_t inventory_n, setmax_t *combo); +void build_comboCPU(opencl_function *func, llf_criteria *cfg, const item_t *items, size_t item_n, size_t inventory_n, setmax_t *combo); #endif diff --git a/src/league/build_combo/item_filter.c b/src/league/build_combo/item_filter.c new file mode 100644 index 0000000..5f611d2 --- /dev/null +++ b/src/league/build_combo/item_filter.c @@ -0,0 +1,75 @@ +#include +#include +#include + +#include "league/database/database.h" +#include "league/ll_formulas.h" +#include "league/common.h" +#include "league/build_combo/item_filter.h" +#include "combinatorics/combination.h" + + + +static bool_t +Test(const item_t *item, uint_t flags) +{ + VECTOR(stats) = VECTOR_ZERO_INIT; + VECTOR_ADD(stats, item->stats); + VECTOR_ADD(stats, item->passive); + + if ((flags & FC_SUSTAIN) && ( + F_HPRegen(stats) || + F_LIFESTEAL_PERCENT(stats) || + F_LIFESTEAL_FLAT(stats) || + F_SPELLVAMP(stats) + )) + return TRUE; + + if ((flags & FC_PHYSICAL_DAMAGE) && ( + F_AD(stats) || + F_CRIT_CHANCE(stats) || + F_CRIT_BONUS(stats) || + F_ATTACK_SPEED(stats) || + F_ARMORPEN_FLAT(stats) || + F_ARMORPEN_PERCENT(stats) || + F_HP2AD(stats) + )) + return TRUE; + + if (flags & FC_MAGIC_DAMAGE) + { + assert(0); + } + + if ((flags & FC_SURVIVABILITY) && ( + F_HP(stats) + )) + return TRUE; + + return FALSE; +} + +item_t* +item_filter(uint_t flags, size_t *count) +{ + item_t *filtered; + *count = 0; + for (size_t i = 0; i < DB_LEN; ++i) + { + if (Test(&db_items[i], flags)) + (*count)++; + } + + filtered = malloc((*count) * sizeof *filtered); + assert(filtered); + if (!filtered) + return 0; + + for (size_t i = 0, j = 0; i < DB_LEN; ++i) + { + if (Test(&db_items[i], flags)) + memcpy(&filtered[j++], &db_items[i], sizeof *filtered); + } + + return filtered; +} diff --git a/src/league/build_combo/item_filter.h b/src/league/build_combo/item_filter.h new file mode 100644 index 0000000..9ae672c --- /dev/null +++ b/src/league/build_combo/item_filter.h @@ -0,0 +1,14 @@ +#ifndef _BUILD_COMBO_ITEM_FILTER_H_ +#define _BUILD_COMBO_ITEM_FILTER_H_ + +#include "types.h" +#include "league/database/db_layout.h" + +#define FC_SUSTAIN 0x00000001 +#define FC_PHYSICAL_DAMAGE 0x00000010 +#define FC_MAGIC_DAMAGE 0x00000100 +#define FC_SURVIVABILITY 0x00001000 + +item_t*item_filter(uint_t flags, size_t *n); + +#endif diff --git a/src/league/build_combo/kernel/k_build_combo.c b/src/league/build_combo/kernel/k_build_combo.c new file mode 100644 index 0000000..dc4af37 --- /dev/null +++ b/src/league/build_combo/kernel/k_build_combo.c @@ -0,0 +1,110 @@ +#ifndef __OPENCL_VERSION__ +#include "opencl_host/dummy.h" +#include +#undef DEBUG_PRINTF +#define DEBUG_PRINTF +#else +#define DEBUG_PRINTF +#endif + +#include "types.h" +#include "league/database/db_layout.h" +#include "league/ll_formulas.h" +#include "league/common.h" +#include "combinatorics/combination.h" + + +result_t +build_combo_m(__constant item_t db_items[], + llf_criteria *cfg, + ulong_t rank, + uint_t db_len, + uint_t combo_len, + __local setmax_t *unique) +{ + setmax_t combo[NK_MULTISET_MAX_K]; + result_t dps; + VECTOR(stats) = VECTOR_ZERO_INIT; + + combo_unrank(rank, db_len, combo_len, combo); + + for (size_t i = 0; i < combo_len; ++i) + { + itemid_t item = combo[i]; + itemid_t passive = db_items[item].passive_id; + + stats_add(db_items, &stats, item, ++unique[passive] == 1); + } + + dps.index = rank; + dps.metric = llf_metric(cfg, &stats); + return dps; +} + + +#ifndef __OPENCL_VERSION__ +#include "opencl_host/function.h" +opencl_allocinfo +build_combo__allocnfo__(opencl_function *func, size_t items_n) +{ + cl_ulong klmem; + opencl_allocinfo nfo = { 0 }; + + NOFAIL(clGetKernelWorkGroupInfo(func->kernel, func->ctx->device, + CL_KERNEL_LOCAL_MEM_SIZE, sizeof klmem, &klmem, 0)); + + nfo.fixed.constant += sizeof (item_t)* items_n; + + nfo.scale_workgroup.local += sizeof (setmax_t)* items_n; + nfo.scale_workgroup.local += sizeof (result_t); + nfo.scale_workgroup.local += klmem; + + nfo.scale_reduce.global = sizeof (result_t); + return nfo; +} + + +opencl_kernel_arg * +build_combo__bind__(opencl_function *X, + uint_t combo_n, + item_t *items, uint_t items_n, + llf_criteria *cfg, + + size_t pass_size, + size_t local_size) +{ + size_t outlen; + + ka_ignore(X); + ka_mconst(X, "db_items", 0, items, sizeof *items * items_n); + ka_value(X, "cfg_input", cfg, sizeof *cfg); + ka_value(X, "db_len", &items_n, sizeof items_n); + ka_value(X, "combo_len", &combo_n, sizeof combo_n); + ka_mlocal(X, "pasv_scratch", sizeof (setmax_t) * items_n * local_size); + ka_mlocal(X, "scratch", sizeof (result_t) * local_size); + + outlen = pass_size / local_size; + return ka_mglobal(X, "output", A_OUT, CL_MEM_WRITE_ONLY, 0, sizeof(result_t)* outlen); +} + +#else +#include "opencl_host/kernel/reduce.cl" +#include "combinatorics/combination.c" +#include "combinatorics/factorial.c" + +__kernel void +build_combo( + ulong_t start_offset, + __constant item_t db_items[], + llf_criteria cfg, + uint_t db_len, + uint_t combo_len, + __local setmax_t *unique, + __local result_t *scratch, + __global result_t *result) +{ + result_t dps = build_combo_m(db_items, &cfg, start_offset + get_global_id(0), db_len, combo_len, unique); + Reduce(dps, scratch, result); +} + +#endif diff --git a/src/league/build_combo/kernel/k_build_combo.h b/src/league/build_combo/kernel/k_build_combo.h new file mode 100644 index 0000000..fa75d23 --- /dev/null +++ b/src/league/build_combo/kernel/k_build_combo.h @@ -0,0 +1,18 @@ +#ifndef _LEAGUE_BUILD_COMBO_KERNEL_H_ +#define _LEAGUE_BUILD_COMBO_KERNEL_H_ + +#ifndef __OPENCL_VERSION__ + +#include "opencl_host/function.h" +#include "league/database/db_layout.h" +#include "league/ll_formulas.h" +#include "league/common.h" +#include "combinatorics/factorial.h" + +opencl_allocinfo build_combo__allocnfo__(opencl_function *func, size_t items_n); +opencl_kernel_arg *build_combo__bind__(opencl_function *X, uint_t combo_n, const item_t *items, uint_t items_n, llf_criteria *cfg, size_t pass_size, size_t local_size); + +result_t build_combo_m(const item_t db_items[], llf_criteria *cfg, ulong_t rank, uint_t db_len, uint_t combo_len, setmax_t *unique); +#endif + +#endif diff --git a/src/league/build_combo/unit_test/unittest_buildcombo.c b/src/league/build_combo/unit_test/unittest_buildcombo.c index 1b780d8..e6bc10a 100644 --- a/src/league/build_combo/unit_test/unittest_buildcombo.c +++ b/src/league/build_combo/unit_test/unittest_buildcombo.c @@ -1,13 +1,15 @@ #include -#include "league/build_combo/build_combo.h" +#include "league/build_combo/kernel/k_build_combo.h" #include "combinatorics/combination.h" #include "combinatorics/factorial.h" #include "league/unit_test/find_max.h" +#include "league/build_combo/build_combo.h" #define UTBC_DEF "-DUNIT_TEST -ID:/GitRoot/llio/src/" -#define UTBC_SRC "#include \"league/build_combo/build_combo.c\"" +#define UTBC_SRC "#include \"league/build_combo/kernel/k_build_combo.c\"" #define ITEM_COUNT 3 +#define INVENTORY 2 const item_t db_test[ITEM_COUNT] = { @@ -30,17 +32,15 @@ const item_t db_test[ITEM_COUNT] = {0}} /*B. F. Sword*/, }; + void unittest_buildcombo() { -#define K 2 - setmax_t combo[K]; - opencl_workset workset; + setmax_t combo[INVENTORY]; + opencl_context gpu; opencl_function *func; - opencl_kernel_arg *output; - ulong_t ncombo; - result_t *results, best; + llf_criteria cfg = { 0 }; cfg.time_frame = 3; cfg.ad_ratio = 3.4f; @@ -50,24 +50,22 @@ unittest_buildcombo() cfg.enemy_mr = 100; cfg.build_maxcost = 15000; cfg.build_maxinventory = 6; + cfg.metric_type = METRIC_ALL_IN; opencl_init(&gpu, 1); func = opencl_build(&gpu, "build_combo", UTBC_SRC, UTBC_DEF); - ncombo = NK_MULTISET(ITEM_COUNT, K); - workset.iterations = 1; - workset.local_size = 1; - workset.pass_size = (size_t)ncombo; - workset.total = ncombo; - - output = build_combo__bind__(func, K, db_test, ITEM_COUNT, &cfg, workset.pass_size, workset.local_size); - results = output->buf_data; - opencl_run(func, 1, TRUE, 0, &workset); - best = FindMax(results, output->buf_size / sizeof *results); - combo_unrank(best.index, ITEM_COUNT, K, combo); + build_comboGPU(func, &cfg, db_test, ITEM_COUNT, INVENTORY, combo); assert(combo[0] == 2 && combo[1] == 2); opencl_function_free(func, 1); opencl_free(&gpu); + + combo[0] = 0; + combo[1] = 0; + + build_comboCPU(func, &cfg, db_test, ITEM_COUNT, INVENTORY, combo); + + assert(combo[0] == 2 && combo[1] == 2); } diff --git a/src/league/build_path/db_input.c b/src/league/build_path/db_input.c index b0ac6d6..5488893 100644 --- a/src/league/build_path/db_input.c +++ b/src/league/build_path/db_input.c @@ -11,27 +11,6 @@ #include "league/database/database.h" -static int item_cmp_id(const void * a, const void * b) -{ - return ((item_t*)a)->id - ((item_t*)b)->id; -} - -#define PTR2INDEX(E, B, SIZE) \ - (!E) ? -1 : ((uintptr_t)E - (uintptr_t)B) / SIZE; - - -// TODO: verify db is actually sorted somewhere. -size_t dbi_find(itemid_t item) // ASSUMES: database is sorted -{ - item_t key; - uintptr_t entry; - - key.id = item; - entry = (uintptr_t)bsearch(&key, db_items, DB_LEN, sizeof key, &item_cmp_id); - - assert(entry); - return PTR2INDEX(entry, db_items, sizeof key); -} static size_t FindItemIndexS(char *item) // TODO: make less terrible. diff --git a/src/league/build_path/kernel/metric_ADPS.c b/src/league/build_path/kernel/metric_area.c similarity index 91% rename from src/league/build_path/kernel/metric_ADPS.c rename to src/league/build_path/kernel/metric_area.c index 86e954c..bff5959 100644 --- a/src/league/build_path/kernel/metric_ADPS.c +++ b/src/league/build_path/kernel/metric_area.c @@ -6,7 +6,7 @@ #include "types.h" #include "league/database/db_layout.h" #include "league/ll_formulas.h" -#include "league/build_path/kernel/metric_ADPS.h" +#include "league/build_path/kernel/metric_area.h" result_t RateBuildpath(__global ideal_t *linext, DB, llf_criteria *cfg, __local ideal_t *pasv_scratch, lattice_info *info, count_t nth_extension); //////////////////////////////////////////////////////// @@ -52,7 +52,7 @@ AddItem(DB, __local ideal_t *unique, VECTOR(*stats), size_t item) static int -MetricAreaDPS(DB, +MetricArea(DB, llf_criteria *cfg, itemid_t *dbi, uint_t dbi_n, @@ -78,12 +78,7 @@ MetricAreaDPS(DB, result += (prev_dps * ((float)db_items[dbi[bs]].upgrade_cost)); //total_cost += db_items[dbi[bs]].upgrade_cost; - -#ifdef METRIC_ADPS - prev_dps = llf_dmgtotal(cfg, &stats); -#else - prev_dps = llf_sustain(cfg, &stats); -#endif + prev_dps = llf_metric(cfg, &stats); } *output = result; //*output = (total_cost) ? result / total_cost : 0; @@ -93,7 +88,7 @@ MetricAreaDPS(DB, #ifdef __OPENCL_VERSION__ #include "opencl_host/kernel/reduce.cl" -__kernel void metric_ADPS( +__kernel void metric_area( ulong_t linext_offset, lattice_info info, __global ideal_t *linext, @@ -124,7 +119,7 @@ __kernel void metric_ADPS( #endif opencl_allocinfo -metric_ADPS__allocnfo__(opencl_function *func, size_t items_n) +metric_area__allocnfo__(opencl_function *func, size_t items_n) { cl_ulong klmem; opencl_allocinfo nfo = { 0 }; @@ -143,7 +138,7 @@ metric_ADPS__allocnfo__(opencl_function *func, size_t items_n) } opencl_kernel_arg * -metric_ADPS__bind__(opencl_function *X, bool_t copy_output, +metric_area__bind__(opencl_function *X, bool_t copy_output, lattice_info *info, opencl_kernel_arg *linext, item_t *items, size_t items_n, @@ -168,7 +163,7 @@ metric_ADPS__bind__(opencl_function *X, bool_t copy_output, #undef X -result_t metric_ADPS(__global ideal_t *linext, DB, llf_criteria *cfg, lattice_info *info, count_t nth_extension) +result_t metric_area(__global ideal_t *linext, DB, llf_criteria *cfg, lattice_info *info, count_t nth_extension) { const size_t size = sizeof (ideal_t) * PASV_SCRATCH_LEN(info->linext_width); ideal_t *pasv_scratch = alloca(size); @@ -201,7 +196,7 @@ RateBuildpath(__global ideal_t *linext, DB, llf_criteria *cfg, __local ideal_t * for (i = 0; i < info->linext_width; i++) ids[i] = linext[i]; - error = MetricAreaDPS(db_items, cfg, ids, info->linext_width, pasv_scratch, &metric.metric); + error = MetricArea(db_items, cfg, ids, info->linext_width, pasv_scratch, &metric.metric); metric.index = (error != ERROR_NONE) ? error : nth_extension; } diff --git a/src/league/build_path/kernel/metric_ADPS.h b/src/league/build_path/kernel/metric_area.h similarity index 68% rename from src/league/build_path/kernel/metric_ADPS.h rename to src/league/build_path/kernel/metric_area.h index 613ef79..5de75e1 100644 --- a/src/league/build_path/kernel/metric_ADPS.h +++ b/src/league/build_path/kernel/metric_area.h @@ -3,6 +3,7 @@ #include "types.h" #include "league/common.h" +#include "league/ll_formulas.h" #include "poset/lattice_types.h" #include "poset/kernel/lattice_kernel.h" @@ -19,7 +20,7 @@ #ifndef __OPENCL_VERSION__ opencl_kernel_arg * -metric_ADPS__bind__(opencl_function *x, bool_t copy_output, +metric_area__bind__(opencl_function *x, bool_t copy_output, lattice_info *info, opencl_kernel_arg *linext, item_t *items, size_t items_n, @@ -27,9 +28,9 @@ metric_ADPS__bind__(opencl_function *x, bool_t copy_output, size_t pass_size, size_t local_size); -opencl_allocinfo metric_ADPS__allocnfo__(opencl_function *, size_t items_n); -result_t metric_ADPS(__global ideal_t *linext, DB, llf_criteria *cfg, lattice_info *info, count_t nth_extension); +opencl_allocinfo metric_area__allocnfo__(opencl_function *, size_t items_n); +result_t metric_area(ideal_t *linext, const item_t *items, llf_criteria *cfg, lattice_info *info, count_t nth_extension); #else -__kernel void metric_ADPS(ulong_t linext_offset, lattice_info info, __global ideal_t *linext, DB, llf_criteria cfg, __local ideal_t *pasv_scratch, __local result_t* scratch, __global result_t* result); +__kernel void metric_area(ulong_t linext_offset, lattice_info info, __global ideal_t *linext, DB, llf_criteria cfg, __local ideal_t *pasv_scratch, __local result_t* scratch, __global result_t* result); #endif #endif diff --git a/src/league/build_path/opencl_bind.c b/src/league/build_path/opencl_bind.c index 2d54251..6b00fae 100644 --- a/src/league/build_path/opencl_bind.c +++ b/src/league/build_path/opencl_bind.c @@ -16,126 +16,17 @@ #include "poset/lattice.h" #include "opencl_host/host.h" #include "opencl_bind.h" -#include "kernel/metric_ADPS.h" +#include "kernel/metric_area.h" -#define ROUND_UP(N,M) ((N + M - 1) / M * M) - -size_t floor_pow2(size_t x) -{ - size_t y; - - do { - y = x; - x = x & (x - 1); - } while (x != 0); - - return y; -} - -#define GPU_SATURATION 256 - -opencl_allocinfo -sum_allocinfo(opencl_allocinfo *nfo, size_t nfo_n) -{ - opencl_allocinfo ai = { 0 }; - - for (size_t i = 0; i < nfo_n; ++i) - { - #define MERGE(F) \ - ai.F.constant += nfo[i].F.constant; \ - ai.F.global += nfo[i].F.global; \ - ai.F.local += nfo[i].F.local; - - MERGE(fixed) - MERGE(scale_pass) - MERGE(scale_workgroup) - MERGE(scale_reduce) - #undef MERGE - } - - return ai; -} - -opencl_workset -ConfigureWorkload(opencl_context *ctx, count_t linext_count, opencl_allocinfo nfo) -{ - size_t local_size, pass_size, saturation; - cl_ulong global_limit; - cl_ulong iterations; - opencl_workset work; - //NOFAIL(clGetKernelWorkGroupInfo(kernel, ctx->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof hint, &hint, 0)); - - { - const cl_ulong max_local = ctx->cfg_max_local_storage; - const cl_ulong max_global = ctx->cfg_max_global_storage; - const cl_ulong max_constant = ctx->cfg_max_const_storage; - - #define CONSTRAINT(F,F2,DEFAULT) \ - ((nfo.F2.F) \ - ? floor_pow2((max_##F - nfo.fixed.F) / nfo.F2.F) \ - : DEFAULT) - - #define MINIMUM(CRIT,F2,D) \ - min(CRIT(local,F2,D),min(CRIT(constant,F2,D),CRIT(global,F2,D))) - - local_size = MINIMUM(CONSTRAINT, scale_workgroup, ctx->cfg_max_workgroup_size); - global_limit = MINIMUM(CONSTRAINT, scale_pass, linext_count); - } - - saturation = min(floor_pow2(global_limit / ctx->cfg_compute_units), GPU_SATURATION); - pass_size = (ctx->cfg_compute_units * saturation * local_size); - -#define ALLOC_SIZE(F) ( \ - nfo.scale_workgroup.F*local_size + \ - nfo.scale_pass.F*pass_size + \ - nfo.fixed.F \ - ) - - printf("OpenCL: %s: %llu\n", "const_alloc", ALLOC_SIZE(constant)); - printf("OpenCL: %s: %llu\n", "global_alloc", ALLOC_SIZE(global)); - printf("OpenCL: %s: %llu\n", "local_alloc", ALLOC_SIZE(local)); - - assert(ALLOC_SIZE(constant) < ctx->cfg_max_const_storage); - assert(ALLOC_SIZE(global) < ctx->cfg_max_global_storage); - assert(ALLOC_SIZE(local) < ctx->cfg_max_local_storage); - - if ((count_t)pass_size > linext_count) - pass_size = (size_t)ROUND_UP(linext_count, 2); - - if (local_size > pass_size) - local_size = ROUND_UP(pass_size, 2); - - linext_count = ROUND_UP(linext_count, local_size); - pass_size = ROUND_UP(pass_size, local_size); - iterations = (cl_ulong)ceil((double)linext_count / (double)pass_size); - - printf("OpenCL: n=%llu, pass_size=%u, local_size=%u, iterations=%llu\n", linext_count, pass_size, local_size, iterations); - - assert(iterations * (cl_ulong)pass_size >= linext_count); - assert(local_size % 2 == 0); // NOTE: the reduction algorithm in the kernel needs a power of 2 local size - - work.iterations = iterations; - work.local_size = local_size; - work.pass_size = pass_size; - work.total = linext_count; - - return work; -} -#undef CONSTRAINT -#undef MINIMUM -#undef ALLOC_SIZE - - opencl_workset clbp_bind(opencl_function *f, ideal_lattice *l, item_t *items, size_t items_len, llf_criteria *cfg_input, opencl_kernel_arg **output_ptr) { opencl_kernel_arg *linext, *output; opencl_context *x = f->ctx; - opencl_allocinfo nfo[2] = { 0 }; opencl_workset wset[2]; opencl_workset work; lattice_info info; @@ -144,19 +35,15 @@ clbp_bind(opencl_function *f, ideal_lattice *l, item_t *items, size_t items_len info.max_neighbors = l->max_neighbors; info.linext_count = l->linext_count; - nfo[0] = linext__allocnfo__(&f[0], l); - nfo[1] = metric_ADPS__allocnfo__(&f[1], items_len); - - wset[0] = ConfigureWorkload(x, l->linext_count, nfo[0]); - wset[1] = ConfigureWorkload(x, l->linext_count, nfo[1]); - + wset[0] = opencl_workcfg(x, l->linext_count, linext__allocnfo__(&f[0], l)); + wset[1] = opencl_workcfg(x, l->linext_count, metric_area__allocnfo__(&f[1], items_len)); // TODO: this doesn't really handle the minimum resource requirement right work = (wset[0].local_size < wset[1].local_size) - ? wset[0] - : wset[1]; + ? wset[0] + : wset[1]; linext = linext__bind__(&f[0], FALSE, l, work.pass_size); - output = metric_ADPS__bind__(&f[1], TRUE, + output = metric_area__bind__(&f[1], TRUE, &info, linext, items, items_len, cfg_input, work.pass_size, @@ -172,10 +59,10 @@ clbp_init(opencl_context *ctx) { #define C_DEFINES "-ID:/GitRoot/llio/src/" - char *kernels[CLBP_KERNEL_N] = { "linext", "metric_ADPS" }; + char *kernels[CLBP_KERNEL_N] = { "linext", "metric_area" }; char *files[CLBP_KERNEL_N] = { "D:/GitRoot/llio/src/poset/kernel/lattice_kernel.c", - "D:/GitRoot/llio/src/league/build_path/kernel/metric_ADPS.c", + "D:/GitRoot/llio/src/league/build_path/kernel/metric_area.c", }; return opencl_buildfilev(ctx, kernels, CLBP_KERNEL_N, files, CLBP_KERNEL_N, C_DEFINES); diff --git a/src/league/build_path/opencl_bind.h b/src/league/build_path/opencl_bind.h index 533df0e..1c40488 100644 --- a/src/league/build_path/opencl_bind.h +++ b/src/league/build_path/opencl_bind.h @@ -1,5 +1,5 @@ #include "opencl_host/host.h" -#include "kernel/metric_ADPS.h" +#include "kernel/metric_area.h" #include "poset/kernel/lattice_kernel.h" #include "poset/lattice.h" diff --git a/src/league/build_path/unit_test/unittest_buildpath.c b/src/league/build_path/unit_test/unittest_buildpath.c index 731708c..9cd6629 100644 --- a/src/league/build_path/unit_test/unittest_buildpath.c +++ b/src/league/build_path/unit_test/unittest_buildpath.c @@ -14,11 +14,12 @@ #include "opencl_host/dummy.h" #include "opencl_host/host.h" #include "league/database/database.h" +#include "league/database/db_search.h" #include "league/ll_formulas.h" #include "poset/lattice.h" #include "poset/kernel/lattice_kernel.h" -#include "../kernel/metric_ADPS.h" +#include "../kernel/metric_area.h" #include "../db_input.h" #include "../opencl_bind.h" #include "league/unit_test/find_max.h" @@ -125,7 +126,7 @@ static result_t TestBuildPathCPU(ideal_lattice *lattice, item_t *db_filtered, si for (i = 0; i < pass_size; ++i) { linext_nth(lattice, le_buf, linext_offset + i, 0); - r = metric_ADPS(le_buf, db_filtered, cfg, &info, linext_offset + i); + r = metric_area(le_buf, db_filtered, cfg, &info, linext_offset + i); if (local_best.metric < r.metric) local_best = r; @@ -165,7 +166,7 @@ PrintExtension(ideal_lattice *il, item_t *db_filtered, itemid_t *idmap, result_t { itemid_t index = db_filtered[le_buf[i]].id; //assert(index == expected[i]); - assert(dbi_find(index) == idmap[le_buf[i]-1]); + assert(db_find(index) == idmap[le_buf[i]-1]); printf("%s, ", db_names[idmap[le_buf[i] - 1]]); } @@ -220,7 +221,7 @@ static void TestBuildpathALL() cfg.enemy_mr = 100; cfg.build_maxcost = 15000; cfg.build_maxinventory = 6; - + cfg.metric_type = METRIC_SUSTAIN; max = tests[i](&lattice, db_filtered, db_len, &cfg); printf("max=%f,i=%d\n", max.metric, max.index); PrintExtension(&lattice, db_filtered, global_idx, max); @@ -237,7 +238,7 @@ extern void unittest_db_input(); void unittest_buildpath() { - glbinit_lattice(); + //glbinit_lattice(); unittest_lattice(1); unittest_db_input(); unittest_opencl(); diff --git a/src/league/build_path/unit_test/unittest_kernel.cl b/src/league/build_path/unit_test/unittest_kernel.cl index aeba598..15a7c3c 100644 --- a/src/league/build_path/unit_test/unittest_kernel.cl +++ b/src/league/build_path/unit_test/unittest_kernel.cl @@ -1,4 +1,4 @@ -#include "league/build_path/kernel/metric_ADPS.c" +#include "league/build_path/kernel/metric_area.c" __kernel void kunittest_mem(DB, __global itemid_t id[], diff --git a/src/league/build_path/unit_test/unittest_opencl.c b/src/league/build_path/unit_test/unittest_opencl.c index 995dd9d..ec07b33 100644 --- a/src/league/build_path/unit_test/unittest_opencl.c +++ b/src/league/build_path/unit_test/unittest_opencl.c @@ -324,6 +324,7 @@ void unittest_opencl_llformulas() cfg.enemy_mr = 100; cfg.build_maxcost = 15000; cfg.build_maxinventory = 6; + cfg.metric_type = METRIC_ALL_IN; ka_mconst(func, "db_items", 0, test_db, sizeof test_db); ka_mconst(func, "build_path", 0, build_path, sizeof build_path); diff --git a/src/league/database/db_search.c b/src/league/database/db_search.c new file mode 100644 index 0000000..9a911a9 --- /dev/null +++ b/src/league/database/db_search.c @@ -0,0 +1,26 @@ +#include +#include +#include "league/database/db_search.h" +#include "league/database/database.h" + +static int item_cmp_id(const void * a, const void * b) +{ + return ((item_t*)a)->id - ((item_t*)b)->id; +} + +#define PTR2INDEX(E, B, SIZE) \ + (!E) ? -1 : ((uintptr_t)E - (uintptr_t)B) / SIZE; + + +// TODO: verify db is actually sorted somewhere. +size_t db_find(itemid_t item) // ASSUMES: database is sorted +{ + item_t key; + uintptr_t entry; + + key.id = item; + entry = (uintptr_t)bsearch(&key, db_items, DB_LEN, sizeof key, &item_cmp_id); + + assert(entry); + return PTR2INDEX(entry, db_items, sizeof key); +} diff --git a/src/league/database/db_search.h b/src/league/database/db_search.h new file mode 100644 index 0000000..7df07bc --- /dev/null +++ b/src/league/database/db_search.h @@ -0,0 +1,7 @@ +#ifndef _LEAGUE_DATABASE_SEARCH_H_ +#define _LEAGUE_DATABASE_SEARCH_H_ +#include "league/database/db_layout.h" + +size_t db_find(itemid_t item); + +#endif diff --git a/src/league/ll_formulas.h b/src/league/ll_formulas.h index 6d35026..5cca9a6 100644 --- a/src/league/ll_formulas.h +++ b/src/league/ll_formulas.h @@ -1,6 +1,8 @@ #ifndef _LL_FORMULAS_H_ #define _LL_FORMULAS_H_ +#include "league/database/db_layout.h" + #ifndef __OPENCL_VERSION__ #define min(a,b) (((a) < (b)) ? (a) : (b)) #endif @@ -17,7 +19,7 @@ typedef struct uint_t build_maxcost; uint_t build_maxinventory; - + enum llf_metric { METRIC_SUSTAIN, METRIC_ALL_IN } metric_type; } llf_criteria; @@ -111,7 +113,14 @@ static inline float llf_sustain(llf_criteria *cfg, VECTOR(*X)) (F_SPELLVAMP(*X) * cast_rate * cast_dmg); } - +static float llf_metric(llf_criteria *cfg, VECTOR(*X)) +{ + switch (cfg->metric_type){ + case METRIC_ALL_IN : return llf_dmgtotal(cfg, X); + case METRIC_SUSTAIN : return llf_sustain(cfg, X); + default : return 0; + } +} #endif diff --git a/src/league/main.c b/src/league/main.c new file mode 100644 index 0000000..d2df526 --- /dev/null +++ b/src/league/main.c @@ -0,0 +1,69 @@ +#include + +#include "opencl_host/function.h" +#include "league/database/db_search.h" +#include "league/database/database.h" +#include "league/build_path/opencl_bind.h" +#include "league/build_combo/item_filter.h" +#include "league/build_combo/build_combo.h" +#include "poset/lattice.h" +#include "combinatorics/factorial.h" + +extern void unittest_league(); +#define UTBC_DEF "-DUNIT_TEST -ID:/GitRoot/llio/src/" +#define UTBC_SRC "#include \"league/build_combo/kernel/k_build_combo.c\"" +#define INVENTORY_N 5 + +void +run(opencl_function *build_combo, opencl_function *build_path) +{ + item_t *items; + size_t items_n; + setmax_t combo[INVENTORY_N]; + llf_criteria cfg = { 0 }; + + cfg.time_frame = 3; + cfg.ad_ratio = 3.4f; + cfg.ap_ratio = 0; + cfg.level = 18; + cfg.enemy_armor = 100; + cfg.enemy_mr = 100; + cfg.build_maxcost = 15000; + cfg.build_maxinventory = 6; + cfg.metric_type = METRIC_ALL_IN; + + items = item_filter(FC_SUSTAIN|FC_PHYSICAL_DAMAGE, &items_n); + build_comboGPU(build_combo, &cfg, items, items_n, INVENTORY_N, combo); + for (size_t i = 0; i < INVENTORY_N; ++i) + combo[i] = db_find(items[combo[i]].id); + + printf("combo: \n\t%s,\n\t%s,\n\t%s,\n\t%s,\n\t%s\n", + db_names[combo[0]], + db_names[combo[1]], + db_names[combo[2]], + db_names[combo[3]], + db_names[combo[4]] + ); +} + + +int +main() +{ + opencl_context gpu; + opencl_function *build_combo, *build_path; + + glbinit_lattice(); + unittest_league(); + + opencl_init(&gpu, 1); + + build_combo = opencl_build(&gpu, "build_combo", UTBC_SRC, UTBC_DEF); + build_path = clbp_init(&gpu); + + run(build_combo, build_path); + + opencl_function_free(build_combo, 1); + opencl_function_free(build_path, CLBP_KERNEL_N); + opencl_free(&gpu); +} \ No newline at end of file diff --git a/src/league/unit_test/unittest_league.c b/src/league/unit_test/unittest_league.c index 4a55886..82aac4a 100644 --- a/src/league/unit_test/unittest_league.c +++ b/src/league/unit_test/unittest_league.c @@ -2,11 +2,10 @@ extern void unittest_buildpath(); extern void unittest_combination(); extern void unittest_buildcombo(); -int main() +void unittest_league() { unittest_buildpath(); unittest_combination(); unittest_buildcombo(); - return 0; } diff --git a/src/opencl_host/function.c b/src/opencl_host/function.c index 9f11766..0d0fb4c 100644 --- a/src/opencl_host/function.c +++ b/src/opencl_host/function.c @@ -221,3 +221,117 @@ opencl_buildv(opencl_context *ctx, const char *kernel_function[], size_t kernel_ return func; } + + + +#define ROUND_UP(N,M) ((N + M - 1) / M * M) + +static bool_t IsPow2(ulong_t x) +{ + return (x != 0) && ((x & (x - 1)) == 0); +} + +static size_t FloorPow2(size_t x) +{ + size_t y; + + do { + y = x; + x = x & (x - 1); + } while (x != 0); + + return y; +} + +#define GPU_SATURATION 256 + +opencl_allocinfo +sum_allocinfo(opencl_allocinfo *nfo, size_t nfo_n) +{ + opencl_allocinfo ai = { 0 }; + + for (size_t i = 0; i < nfo_n; ++i) + { +#define MERGE(F) \ + ai.F.constant += nfo[i].F.constant; \ + ai.F.global += nfo[i].F.global; \ + ai.F.local += nfo[i].F.local; + + MERGE(fixed) + MERGE(scale_pass) + MERGE(scale_workgroup) + MERGE(scale_reduce) +#undef MERGE + } + + return ai; +} + +opencl_workset +opencl_workcfg(opencl_context *ctx, ulong_t total_work, opencl_allocinfo nfo) +{ + size_t local_size, pass_size, saturation; + cl_ulong global_limit; + cl_ulong iterations; + opencl_workset work; + //NOFAIL(clGetKernelWorkGroupInfo(kernel, ctx->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof hint, &hint, 0)); + + { + const cl_ulong max_local = ctx->cfg_max_local_storage; + const cl_ulong max_global = ctx->cfg_max_global_storage; + const cl_ulong max_constant = ctx->cfg_max_const_storage; + + #define CONSTRAINT(F,F2,DEFAULT) \ + ((nfo.F2.F) \ + ? FloorPow2((max_##F - nfo.fixed.F) / nfo.F2.F) \ + : DEFAULT) + + #define MINIMUM(CRIT,F2,D) \ + min(CRIT(local, F2, D), min(CRIT(constant, F2, D), CRIT(global, F2, D))) + + local_size = MINIMUM(CONSTRAINT, scale_workgroup, ctx->cfg_max_workgroup_size); + global_limit = MINIMUM(CONSTRAINT, scale_pass, total_work); + } + + saturation = min(max(1, FloorPow2(global_limit / ctx->cfg_compute_units)), GPU_SATURATION); + pass_size = (ctx->cfg_compute_units * saturation * local_size); + + #define ALLOC_SIZE(F) ( \ + nfo.scale_workgroup.F*local_size + \ + nfo.scale_pass.F*pass_size + \ + nfo.fixed.F \ + ) + + printf("OpenCL: %s: %llu\n", "const_alloc", ALLOC_SIZE(constant)); + printf("OpenCL: %s: %llu\n", "global_alloc", ALLOC_SIZE(global)); + printf("OpenCL: %s: %llu\n", "local_alloc", ALLOC_SIZE(local)); + + assert(ALLOC_SIZE(constant) < ctx->cfg_max_const_storage); + assert(ALLOC_SIZE(global) < ctx->cfg_max_global_storage); + assert(ALLOC_SIZE(local) < ctx->cfg_max_local_storage); + + if ((ulong_t)pass_size > total_work) + pass_size = (size_t)ROUND_UP(total_work, 2); + + if (local_size > pass_size) + local_size = FloorPow2(pass_size); + + total_work = ROUND_UP(total_work, local_size); + pass_size = ROUND_UP(pass_size, local_size); + iterations = (cl_ulong)ceil((double)total_work / (double)pass_size); + + printf("OpenCL: n=%llu, pass_size=%u, local_size=%u, iterations=%llu\n", total_work, pass_size, local_size, iterations); + + assert(iterations * (cl_ulong)pass_size >= total_work); + assert(IsPow2(local_size)); // NOTE: the reduction algorithm in the kernel needs a power of 2 local size + + work.iterations = iterations; + work.local_size = local_size; + work.pass_size = pass_size; + work.total = total_work; + + return work; +} +#undef CONSTRAINT +#undef MINIMUM +#undef ALLOC_SIZE diff --git a/src/opencl_host/function.h b/src/opencl_host/function.h index 5a84b25..fc310a7 100644 --- a/src/opencl_host/function.h +++ b/src/opencl_host/function.h @@ -14,15 +14,15 @@ typedef struct } opencl_workset; -typedef struct +struct opencl_allocinfo_mem { - struct opencl_allocinfo_mem - { - cl_ulong local; - cl_ulong global; - cl_ulong constant; - }; + cl_ulong local; + cl_ulong global; + cl_ulong constant; +}; +typedef struct +{ struct opencl_allocinfo_mem fixed; struct opencl_allocinfo_mem scale_workgroup; struct opencl_allocinfo_mem scale_pass; @@ -35,6 +35,8 @@ typedef struct void opencl_memcheck(opencl_context *ctx, cl_kernel, opencl_kernel_params *args, opencl_workset *work); cl_ulong opencl_run(opencl_function *function, size_t func_n, bool_t multipass, ulong_t offset, opencl_workset *work); void opencl_function_free(opencl_function *func, size_t func_n); +opencl_workset opencl_workcfg(opencl_context *ctx, ulong_t total_work, opencl_allocinfo nfo); +opencl_allocinfo sum_allocinfo(opencl_allocinfo *nfo, size_t nfo_n); opencl_function * opencl_buildfile(opencl_context *ctx, const char *kernel_function, const char *source_file, const char *build_flags); opencl_function * opencl_buildfilev(opencl_context *ctx, const char *kernel_function[], size_t kernel_n, const char *source_file[], size_t source_n, const char *build_flags); diff --git a/src/opencl_host/host.c b/src/opencl_host/host.c index ee14f99..a709f4f 100644 --- a/src/opencl_host/host.c +++ b/src/opencl_host/host.c @@ -68,6 +68,9 @@ opencl_run(opencl_context *ctx, cl_kernel kernel, opencl_kernel_params *args, op } */ + + + void opencl_init(opencl_context *ctx, int profiling) { size_t i; diff --git a/src/poset/lattice.c b/src/poset/lattice.c index 8cadb03..16ca743 100644 --- a/src/poset/lattice.c +++ b/src/poset/lattice.c @@ -436,7 +436,8 @@ void MapToLinext(ideal_lattice *lattice) size_t i, slen = lattice->ctx.vertex_count * lattice->ctx.max_neighbors; for (i = 0; i < slen; ++i) - lattice->ideals[i] = lattice->ctx.linear_extension[lattice->ideals[i] - 1]; + if (lattice->neighbors[i] != INVALID_NEIGHBOR) + lattice->ideals[i] = lattice->ctx.linear_extension[lattice->ideals[i] - 1]; } int lattice_create(ideal_t p_relations[][2], size_t p_reln, size_t n, ideal_lattice *lattice)