PostgreSQL异构引擎pg-strom源码解读

以源码test目录下testquery.sql样例SQL来解读pg-strom源码
SELECT cat, count(*), avg(ax) FROM t0 NATURAL JOIN t1 NATURAL JOIN t2 NATURAL JOIN t3 NATURAL JOIN t4 NATURAL JOIN t5 NATURAL JOIN t6 NATURAL JOIN t7 NATURAL JOIN t8 GROUP BY cat;

EXPLAIN执行计划

HashAggregate  (cost=3529132.25..3529132.57 rows=26 width=12)Output: t0.cat, pgstrom.count((pgstrom.nrows())), pgstrom.avg((pgstrom.nrows((t1.ax IS NOT NULL))), (pgstrom.psum(t1.ax)))Group Key: t0.cat->  Custom Scan (GpuPreAgg)  (cost=29099.24..2859511.64 rows=234 width=48)Output: t0.cat, pgstrom.nrows(), pgstrom.nrows((t1.ax IS NOT NULL)), pgstrom.psum(t1.ax)Reduction: Local + GlobalGPU Projection: t0.cat, t1.axLogic Parameter: SafetyLimit: 2147483647, KeyDistSalt: 9Extra: outer-bulk-exec->  Custom Scan (GpuJoin) on public.t0  (cost=25099.24..2826221.34 rows=93721454 width=12)Output: t0.cat, t1.axGPU Projection: t0.cat::text, t1.ax::double precisionDepth 1: GpuHashJoin, HashKeys: (t0.fid)JoinQuals: (t0.fid = t6.fid)Nrows (in/out: 97.75%), KDS-Hash (size: 13.47MB, nbatches: 1)Depth 2: GpuHashJoin, HashKeys: (t0.cid)JoinQuals: (t0.cid = t3.cid)Nrows (in/out: 98.34%), KDS-Hash (size: 13.47MB, nbatches: 1)Depth 3: GpuHashJoin, HashKeys: (t0.bid)JoinQuals: (t0.bid = t2.bid)Nrows (in/out: 98.76%), KDS-Hash (size: 13.47MB, nbatches: 1)Depth 4: GpuHashJoin, HashKeys: (t0.eid)JoinQuals: (t0.eid = t5.eid)Nrows (in/out: 98.97%), KDS-Hash (size: 13.47MB, nbatches: 1)Depth 5: GpuHashJoin, HashKeys: (t0.aid)JoinQuals: (t0.aid = t1.aid)Nrows (in/out: 100.00%), KDS-Hash (size: 13.47MB, nbatches: 1)Depth 6: GpuHashJoin, HashKeys: (t0.did)JoinQuals: (t0.did = t4.did)Nrows (in/out: 100.00%), KDS-Hash (size: 13.47MB, nbatches: 1)Depth 7: GpuHashJoin, HashKeys: (t0.gid)JoinQuals: (t0.gid = t7.gid)Nrows (in/out: 100.00%), KDS-Hash (size: 13.47MB, nbatches: 1)Depth 8: GpuHashJoin, HashKeys: (t0.hid)JoinQuals: (t0.hid = t8.hid)Nrows (in/out: 99.74%), KDS-Hash (size: 13.47MB, nbatches: 1)Extra: bulk-exec-support, row-format->  Seq Scan on public.t6  (cost=0.00..1935.00 rows=100000 width=4)Output: t6.fid, t6.ftext, t6.fx->  Seq Scan on public.t3  (cost=0.00..1935.00 rows=100000 width=4)Output: t3.cid, t3.ctext, t3.cx->  Seq Scan on public.t2  (cost=0.00..1935.00 rows=100000 width=4)Output: t2.bid, t2.btext, t2.bx->  Seq Scan on public.t5  (cost=0.00..1935.00 rows=100000 width=4)Output: t5.eid, t5.etext, t5.ex->  Seq Scan on public.t1  (cost=0.00..1935.00 rows=100000 width=12)Output: t1.aid, t1.atext, t1.ax->  Seq Scan on public.t4  (cost=0.00..1935.00 rows=100000 width=4)Output: t4.did, t4.dtext, t4.dx->  Seq Scan on public.t7  (cost=0.00..1935.00 rows=100000 width=4)Output: t7.gid, t7.gtext, t7.gx->  Seq Scan on public.t8  (cost=0.00..1935.00 rows=100000 width=4)Output: t8.hid, t8.htext, t8.hx

数据库内核源码执行流程

src/backend/main/main.c/main(int argc, char *argv[]) ..PostmasterMain(argc, argv);->src/backend/postmaster/postmaster.c/PostmasterMain(int argc, char *argv[]) ..pqsignal(SIGUSR1, sigusr1_handler);->src/backend/postmaster/postmaster.c/static void sigusr1_handler(SIGNAL_ARGS) ..StartAutovacuumWorker();->src/backend/postmaster/postmaster.c/static void StartAutovacuumWorker(void) ..bn->pid = StartAutoVacWorker();->src/backend/postmaster/autovacuum.c/int StartAutoVacWorker(void) ..AutoVacWorkerMain(0, NULL);->src/backend/postmaster/autovacuum.c/NON_EXEC_STATIC void AutoVacWorkerMain(int argc, char *argv[]) ..do_autovacuum();->src/backend/postmaster/autovacuum.c/static void do_autovacuum(void) ..ScanKeyInit(&key, Anum_pg_class_relkind, BTEqualStrategyNumber, F_CHAREQ, CharGetDatum(RELKIND_TOASTVALUE));->src/backend/access/common/scankey.c/void ScanKeyInit(ScanKey entry, AttrNumber attributeNumber, StrategyNumber strategy, RegProcedure procedure, Datum argument) ..fmgr_info(procedure, &entry->sk_func);->src/backend/utils/fmgr/fmgr.c/void fmgr_info(Oid functionId, FmgrInfo *finfo) ..fmgr_info_cxt_security(functionId, finfo, CurrentMemoryContext, false);->src/backend/utils/fmgr/fmgr.c/static void fmgr_info_cxt_security(Oid functionId, FmgrInfo *finfo, MemoryContext mcxt, bool ignore_security)->src/backend/postmaster/postmaster.c/PostmasterMain(int argc, char *argv[]) ..InitializeGUCOptions();->src/backend/utils/misc/guc.c/void InitializeGUCOptions(void) ..build_guc_variables();->src/backend/utils/misc/guc.c/void build_guc_variables(void) ..for (i = 0; ConfigureNamesString[i].gen.name; i++);->src/backend/utils/misc/guc.c/static struct config_string ConfigureNamesString[]->src/backend/postmaster/postmaster.c/PostmasterMain(int argc, char *argv[]) ..SelectConfigFiles(userDoption, progname)->src/backend/utils/misc/guc.c/SetConfigOption("config_file", fname, PGC_POSTMASTER, PGC_S_OVERRIDE) ..SetConfigOption("config_file", fname, PGC_POSTMASTER, PGC_S_OVERRIDE);->src/backend/utils/misc/guc.c/(void) set_config_option(name, value, context, source, GUC_ACTION_SET, true, 0, false);->src/backend/utils/misc/guc.c/static struct config_generic *find_option(const char *name, bool create_placeholders, int elevel)->src/backend/utils/misc/guc.c/(void) set_config_option(name, value, context, source, GUC_ACTION_SET, true, 0, false) ..switch (record->context);->src/backend/utils/misc/guc.c/(void) set_config_option(name, value, context, source, GUC_ACTION_SET, true, 0, false) ..switch (record->vartype);->src/backend/postmaster/postmaster.c/PostmasterMain(int argc, char *argv[]) ..process_shared_preload_libraries(); // load pg-strom->src/backend/utils/init/miscinit.c/void process_shared_preload_libraries(void) ..load_libraries(shared_preload_libraries_string,  "shared_preload_libraries", false);->src/backend/utils/init/miscinit.c/static void load_libraries(const char *libraries, const char *gucname, bool restricted) ..load_file(filename, restricted);->src/backend/utils/fmgr/dfmgr.c/void load_file(const char *filename, bool restricted) ..(void) internal_load_library(fullname);->src/backend/utils/fmgr/dfmgr.c/static void *internal_load_library(const char *libname) ..(*PG_init) ();->pg_strom/src/main.c/void _PG_init(void)->src/backend/postmaster/postmaster.c/PostmasterMain(int argc, char *argv[]) ..ServerLoop();->src/backend/postmaster/postmaster.c/ServerLoop(void) ..initMasks(&readmask);->src/backend/postmaster/postmaster.c/ServerLoop(void) ..for (;;)->src/backend/postmaster/postmaster.c/ServerLoop(void) ..select(nSockets, &rmask, NULL, NULL, &timeout); #include <sys/select.h>->src/backend/postmaster/postmaster.c/ServerLoop(void) ..ConnCreate(ListenSocket[i]);->src/backend/postmaster/postmaster.c/ServerLoop(void) ..BackendStartup(port);->src/backend/postmaster/postmaster.c/ServerLoop(void) ..AutoVacPID = StartAutoVacLauncher();->src/backend/postmaster/autovacuum.c/int StartAutoVacLauncher(void) ..AutoVacLauncherMain(0, NULL);->src/backend/postmaster/autovacuum.c/NON_EXEC_STATIC void AutoVacLauncherMain(int argc, char *argv[]) ..->src/backend/postmaster/postmaster.c/BackendStartup(Port *port) ..PostmasterRandom();->src/backend/postmaster/postmaster.c/BackendStartup(Port *port) ..canAcceptConnections();->src/backend/postmaster/postmaster.c/BackendStartup(Port *port) ..fork_process(); #include "postmaster/fork_process.h"->src/backend/postmaster/postmaster.c/BackendStartup(Port *port) ..InitPostmasterChild();->src/backend/postmaster/postmaster.c/BackendStartup(Port *port) ..ClosePostmasterPorts(false);->src/backend/postmaster/postmaster.c/BackendStartup(Port *port) ..BackendInitialize(port);->src/backend/postmaster/postmaster.c/BackendStartup(Port *port) ..BackendRun(port);->src/backend/postmaster/postmaster.c/BackendRun(Port *port) ..PostgresMain(ac, av, port->database_name, port->user_name);->src/backend/tcop/postgres.c/PostgresMain(int argc, char *argv[],const char *dbname,const char *username) ..BaseInit();->src/backend/tcop/postgres.c/PostgresMain(int argc, char *argv[],const char *dbname,const char *username) ..InitProcess();->src/backend/tcop/postgres.c/PostgresMain(int argc, char *argv[],const char *dbname,const char *username) ..InitPostgres(dbname, InvalidOid, username, InvalidOid, NULL);->src/backend/tcop/postgres.c/PostgresMain(int argc, char *argv[],const char *dbname,const char *username) ..switch (firstchar)->src/backend/tcop/postgres.c/PostgresMain(int argc, char *argv[],const char *dbname,const char *username) ..case 'P':->src/backend/tcop/postgres.c/static void exec_parse_message(const char *query_string,const char *stmt_name,Oid *paramTypes,int numParams) ..parsetree_list = pg_parse_query(query_string);->src/backend/tcop/postgres.c/List *pg_parse_query(const char *query_string) ..raw_parsetree_list = raw_parser(query_string);->src/backend/parser/parser.c/List *raw_parser(const char *str) ..yyresult = base_yyparse(yyscanner);->src/backend/tcop/postgres.c/static void exec_parse_message(const char *query_string,const char *stmt_name,Oid *paramTypes,int numParams) ..querytree_list = pg_rewrite_query(query);->src/backend/tcop/postgres.c/static List *pg_rewrite_query(Query *query) ..querytree_list = QueryRewrite(query);->src/backend/rewrite/rewriteHandler.c/List *QueryRewrite(Query *parsetree) ..querylist = RewriteQuery(parsetree, NIL);->src/backend/rewrite/rewriteHandler.c/static List *RewriteQuery(Query *parsetree, List *rewrite_events)->src/backend/rewrite/rewriteHandler.c/List *QueryRewrite(Query *parsetree) ..query = fireRIRrules(query, NIL, false);->src/backend/rewrite/rewriteHandler.c/static Query *fireRIRrules(Query *parsetree, List *activeRIRs, bool forUpdatePushedDown)->src/backend/tcop/postgres.c/PostgresMain(int argc, char *argv[],const char *dbname,const char *username) ..case 'B':->src/backend/tcop/postgres.c/static void exec_bind_message(StringInfo input_message) ..portal = CreatePortal(portal_name, true, true);->src/backend/tcop/postgres.c/static void exec_bind_message(StringInfo input_message) ..PortalStart(portal, params, 0, InvalidSnapshot);->src/backend/tcop/pquery.c/void PortalStart(Portal portal, ParamListInfo params, int eflags, Snapshot snapshot)->src/backend/tcop/postgres.c/PostgresMain(int argc, char *argv[],const char *dbname,const char *username) ..case 'D':->src/backend/tcop/postgres.c/PostgresMain(int argc, char *argv[],const char *dbname,const char *username) ..case 'E':->src/backend/tcop/postgres.c/static void exec_execute_message(const char *portal_name, long max_rows) ..completed = PortalRun(portal,max_rows,true,receiver,receiver,completionTag);EXPLAIN [ ANALYZE ] [ VERBOSE ]->src/backend/tcop/pquery.c/bool PortalRun(Portal portal, long count, bool isTopLevel,DestReceiver *dest, DestReceiver *altdest,char *completionTag) ..if (portal->strategy != PORTAL_ONE_SELECT && !portal->holdStore) FillPortalStore(portal, isTopLevel);->src/backend/tcop/pquery.c/static void FillPortalStore(Portal portal, bool isTopLevel) ..case PORTAL_UTIL_SELECT:PortalRunUtility(portal, (Node *) linitial(portal->stmts), isTopLevel, treceiver, completionTag);->src/backend/tcop/pquery.c/static void PortalRunUtility(Portal portal, Node *utilityStmt, bool isTopLevel, DestReceiver *dest, char *completionTag) ..ProcessUtility(utilityStmt, portal->sourceText, isTopLevel ? PROCESS_UTILITY_TOPLEVEL : PROCESS_UTILITY_QUERY, portal->portalParams, dest, completionTag);->src/backend/tcop/utility.c/void ProcessUtility(Node *parsetree, const char *queryString, ProcessUtilityContext context, ParamListInfo params, DestReceiver *dest, char *completionTag) ..standard_ProcessUtility(parsetree, queryString, context, params, dest, completionTag);->src/backend/tcop/utility.c/void standard_ProcessUtility(Node *parsetree, const char *queryString, ProcessUtilityContext context, ParamListInfo params, DestReceiver *dest, char *completionTag) ..case T_ExplainStmt:ExplainQuery((ExplainStmt *) parsetree, queryString, params, dest);->src/backend/commands/explain.c/void ExplainQuery(ExplainStmt *stmt, const char *queryString, ParamListInfo params, DestReceiver *dest) ..ExplainOneQuery((Query *) lfirst(l), NULL, es, queryString, params);->src/backend/commands/explain.c/static void ExplainOneQuery(Query *query, IntoClause *into, ExplainState *es, const char *queryString, ParamListInfo params) ..plan = pg_plan_query(query, 0, params);->src/backend/tcop/postgres.c/PlannedStmt *pg_plan_query(Query *querytree, int cursorOptions, ParamListInfo boundParams) ..plan = planner(querytree, cursorOptions, boundParams);->src/backend/optimizer/plan/planner.c/PlannedStmt *planner(Query *parse, int cursorOptions, ParamListInfo boundParams) ..result = (*planner_hook) (parse, cursorOptions, boundParams);->pg_strom/src/main.c/static PlannedStmt *pgstrom_planner_entrypoint(Query *parse, int cursorOptions, ParamListInfo boundParams) ..result = standard_planner(parse, cursorOptions, boundParams);->src/backend/optimizer/plan/planner.c/PlannedStmt *standard_planner(Query *parse, int cursorOptions, ParamListInfo boundParams) ..top_plan = subquery_planner(glob, parse, NULL, false, tuple_fraction, &root); /* PlannerInfo begin */->src/backend/optimizer/plan/planner.c/Plan *subquery_planner(PlannerGlobal *glob, Query *parse, PlannerInfo *parent_root, bool hasRecursion, double tuple_fraction, PlannerInfo **subroot) ..plan = grouping_planner(root, tuple_fraction);->src/backend/optimizer/plan/planner.c/static Plan *grouping_planner(PlannerInfo *root, double tuple_fraction) ..final_rel = query_planner(root, sub_tlist, standard_qp_callback, &qp_extra);->src/backend/optimizer/plan/planmain.c/RelOptInfo *query_planner(PlannerInfo *root, List *tlist, query_pathkeys_callback qp_callback, void *qp_extra) ..final_rel = make_one_rel(root, joinlist);->src/backend/optimizer/path/allpaths.c/RelOptInfo *make_one_rel(PlannerInfo *root, List *joinlist) ..set_base_rel_sizes(root);->src/backend/optimizer/path/allpaths.c/static void set_base_rel_sizes(PlannerInfo *root) ..set_rel_size(root, rel, rti, root->simple_rte_array[rti]);->src/backend/optimizer/path/allpaths.c/static void set_rel_size(PlannerInfo *root, RelOptInfo *rel, Index rti, RangeTblEntry *rte) ..set_plain_rel_size(root, rel, rte); /* Plain relation */->src/backend/optimizer/path/allpaths.c/static void set_plain_rel_size(PlannerInfo *root, RelOptInfo *rel, RangeTblEntry *rte) ..check_partial_indexes(root, rel); // T_SeqScan No T_IndexScan->src/backend/optimizer/path/allpaths.c/static void set_plain_rel_size(PlannerInfo *root, RelOptInfo *rel, RangeTblEntry *rte) ..set_baserel_size_estimates(root, rel);->src/backend/optimizer/path/costsize.c/void set_baserel_size_estimates(PlannerInfo *root, RelOptInfo *rel) ..nrows = rel->tuples * clauselist_selectivity(root, rel->baserestrictinfo, 0, JOIN_INNER, NULL);->src/backend/optimizer/path/costsize.c/void set_baserel_size_estimates(PlannerInfo *root, RelOptInfo *rel) ..cost_qual_eval(&rel->baserestrictcost, rel->baserestrictinfo, root);->src/backend/optimizer/path/costsize.c/void cost_qual_eval(QualCost *cost, List *quals, PlannerInfo *root) ..cost_qual_eval_walker(qual, &context);->src/backend/optimizer/path/costsize.c/static bool cost_qual_eval_walker(Node *node, cost_qual_eval_context *context) ..context->total.per_tuple += get_func_cost(((FuncExpr *) node)->funcid) * cpu_operator_cost;->src/backend/utils/cache/lsyscache.c/float4 get_func_cost(Oid funcid) ..result = ((Form_pg_proc) GETSTRUCT(tp))->procost;->src/backend/optimizer/path/costsize.c/void set_baserel_size_estimates(PlannerInfo *root, RelOptInfo *rel) ..set_rel_width(root, rel);->src/backend/optimizer/path/costsize.c/static void set_rel_width(PlannerInfo *root, RelOptInfo *rel) ..item_width = get_attavgwidth(reloid, var->varattno);->src/backend/utils/cache/lsyscache.c/int32 get_attavgwidth(Oid relid, AttrNumber attnum) ..tp = SearchSysCache3(STATRELATTINH, ObjectIdGetDatum(relid), Int16GetDatum(attnum), BoolGetDatum(false));->src/backend/optimizer/path/allpaths.c/RelOptInfo *make_one_rel(PlannerInfo *root, List *joinlist) ..set_base_rel_pathlists(root);->src/backend/optimizer/path/allpaths.c/static void set_base_rel_pathlists(PlannerInfo *root) ..set_rel_pathlist(root, rel, rti, root->simple_rte_array[rti]);->src/backend/optimizer/path/allpaths.c/static void set_rel_pathlist(PlannerInfo *root, RelOptInfo *rel, Index rti, RangeTblEntry *rte) ..if (set_rel_pathlist_hook) (*set_rel_pathlist_hook) (root, rel, rti, rte);->pg_strom/src/gpuscan.c/void pgstrom_init_gpuscan(void) ..set_rel_pathlist_hook = gpuscan_add_scan_path;->src/backend/optimizer/path/allpaths.c/static void set_rel_pathlist(PlannerInfo *root, RelOptInfo *rel, Index rti, RangeTblEntry *rte) ..set_cheapest(rel);->src/backend/optimizer/util/pathnode.c/void set_cheapest(RelOptInfo *parent_rel)->src/backend/optimizer/path/allpaths.c/RelOptInfo *make_one_rel(PlannerInfo *root, List *joinlist) ..rel = make_rel_from_joinlist(root, joinlist);->src/backend/optimizer/path/allpaths.c/static RelOptInfo *make_rel_from_joinlist(PlannerInfo *root, List *joinlist) ..return standard_join_search(root, levels_needed, initial_rels);->src/backend/optimizer/path/allpaths.c/RelOptInfo *standard_join_search(PlannerInfo *root, int levels_needed, List *initial_rels) ..join_search_one_level(root, lev);->src/backend/optimizer/path/joinrels.c/void join_search_one_level(PlannerInfo *root, int level) ..(void) make_join_rel(root, old_rel, new_rel);->src/backend/optimizer/path/joinrels.c/RelOptInfo *make_join_rel(PlannerInfo *root, RelOptInfo *rel1, RelOptInfo *rel2) ..add_paths_to_joinrel(root, joinrel, rel1, rel2, JOIN_INNER, sjinfo, restrictlist);->src/backend/optimizer/path/joinpath.c/void add_paths_to_joinrel(PlannerInfo *root, RelOptInfo *joinrel, RelOptInfo *outerrel, RelOptInfo *innerrel, JoinType jointype, SpecialJoinInfo *sjinfo, List *restrictlist) ..if (set_join_pathlist_hook) set_join_pathlist_hook(root, joinrel, outerrel, innerrel, jointype, &extra);->pg_strom/src/gpujoin.c/pgstrom_init_gpujoin(void) ..set_join_pathlist_hook = gpujoin_add_join_path;->pg_strom/src/gpujoin.c/static void gpujoin_add_join_path(PlannerInfo *root, RelOptInfo *joinrel, RelOptInfo *outerrel, RelOptInfo *innerrel, JoinType jointype, JoinPathExtraData *extra) ..create_gpujoin_path( root, joinrel, outer_path, inner_path_list, final_tlist, param_info, required_outer);->pg_strom/src/gpujoin.c/static void create_gpujoin_path(PlannerInfo *root, RelOptInfo *joinrel, Path *outer_path, List *inner_path_list, List *final_tlist, ParamPathInfo *param_info, Relids required_outer)->src/backend/optimizer/path/allpaths.c/RelOptInfo *standard_join_search(PlannerInfo *root, int levels_needed, List *initial_rels) ..set_cheapest(rel);->src/backend/optimizer/util/pathnode.c/void set_cheapest(RelOptInfo *parent_rel)->src/backend/optimizer/plan/planner.c/static Plan *grouping_planner(PlannerInfo *root, double tuple_fraction) ..result_plan = create_plan(root, best_path);->src/backend/optimizer/plan/createplan.c/Plan *create_plan(PlannerInfo *root, Path *best_path) ..plan = create_plan_recurse(root, best_path);->src/backend/optimizer/plan/createplan.c/static Plan *create_plan_recurse(PlannerInfo *root, Path *best_path) ..plan = create_scan_plan(root, best_path);->src/backend/optimizer/plan/createplan.c/static Plan *create_scan_plan(PlannerInfo *root, Path *best_path) ..plan = (Plan *) create_customscan_plan(root, (CustomPath *) best_path, tlist, scan_clauses);->src/backend/optimizer/plan/createplan.c/static CustomScan *create_customscan_plan(PlannerInfo *root, CustomPath *best_path, List *tlist, List *scan_clauses) ..cplan = (CustomScan *) best_path->methods->PlanCustomPath(root, rel, best_path, tlist, scan_clauses, custom_plans);->pg_strom/src/gpujoin.c/void pgstrom_init_gpujoin(void) ..gpujoin_path_methods.PlanCustomPath = create_gpujoin_plan;->pg_strom/src/gpujoin.c/static Plan *create_gpujoin_plan(PlannerInfo *root, RelOptInfo *rel, CustomPath *best_path, List *tlist, List *clauses, List *custom_plans)->pg_strom/src/main.c/static PlannedStmt *pgstrom_planner_entrypoint(Query *parse, int cursorOptions, ParamListInfo boundParams) ..pgstrom_recursive_grafter(result, NULL, &result->planTree);->src/backend/commands/explain.c/static void ExplainOneQuery(Query *query, IntoClause *into, ExplainState *es, const char *queryString, ParamListInfo params) ..ExplainOnePlan(plan, into, es, queryString, params, &planduration);->src/backend/commands/explain.c/void ExplainOnePlan(PlannedStmt *plannedstmt, IntoClause *into, ExplainState *es, const char *queryString, ParamListInfo params, const instr_time *planduration) ..ExecutorStart(queryDesc, eflags);->src/backend/executor/execMain.c/void ExecutorStart(QueryDesc *queryDesc, int eflags) ..standard_ExecutorStart(queryDesc, eflags);->src/backend/executor/execMain.c/void standard_ExecutorStart(QueryDesc *queryDesc, int eflags) ..InitPlan(queryDesc, eflags);->src/backend/executor/execMain.c/static void InitPlan(QueryDesc *queryDesc, int eflags) ..planstate = ExecInitNode(plan, estate, eflags);->src/backend/executor/execProcnode.c/PlanState *ExecInitNode(Plan *node, EState *estate, int eflags) ..case T_Agg:result = (PlanState *) ExecInitAgg((Agg *) node, estate, eflags);->src/backend/executor/nodeAgg.c/AggState *ExecInitAgg(Agg *node, EState *estate, int eflags)->src/backend/executor/execProcnode.c/PlanState *ExecInitNode(Plan *node, EState *estate, int eflags) ..case T_CustomScan:result = (PlanState *) ExecInitCustomScan((CustomScan *) node, estate, eflags);->src/backend/executor/nodeCustom.c/CustomScanState *ExecInitCustomScan(CustomScan *cscan, EState *estate, int eflags) ..css = (CustomScanState *) cscan->methods->CreateCustomScanState(cscan);->pg_strom/src/gpujoin.c/void pgstrom_init_gpujoin(void) ..gpujoin_plan_methods.CreateCustomScanState = gpujoin_create_scan_state;->pg_strom/src/gpujoin.c/static Node *gpujoin_create_scan_state(CustomScan *node)->src/backend/executor/nodeCustom.c/CustomScanState *ExecInitCustomScan(CustomScan *cscan, EState *estate, int eflags) ..css->ss.ps.targetlist = (List *) ExecInitExpr((Expr *) cscan->scan.plan.targetlist, (PlanState *) css);->src/backend/executor/nodeCustom.c/CustomScanState *ExecInitCustomScan(CustomScan *cscan, EState *estate, int eflags) ..css->ss.ps.qual = (List *) ExecInitExpr((Expr *) cscan->scan.plan.qual, (PlanState *) css);->src/backend/executor/nodeCustom.c/CustomScanState *ExecInitCustomScan(CustomScan *cscan, EState *estate, int eflags) ..ExecInitScanTupleSlot(estate, &css->ss);->src/backend/executor/nodeCustom.c/CustomScanState *ExecInitCustomScan(CustomScan *cscan, EState *estate, int eflags) ..ExecInitResultTupleSlot(estate, &css->ss.ps);->src/backend/executor/nodeCustom.c/CustomScanState *ExecInitCustomScan(CustomScan *cscan, EState *estate, int eflags) ..css->methods->BeginCustomScan(css, estate, eflags);->pg_strom/src/gpujoin.c/gpujoin_exec_methods.BeginCustomScan = gpujoin_begin;->pg_strom/src/gpujoin.c/static void gpujoin_begin(CustomScanState *node, EState *estate, int eflags) ..gcontext = pgstrom_get_gpucontext();->pg_strom/src/cuda_control.c/GpuContext *pgstrom_get_gpucontext(void) ..gcontext = pgstrom_create_gpucontext(CurrentResourceOwner, &context_reused);->pg_strom/src/cuda_control.c/static GpuContext *pgstrom_create_gpucontext(ResourceOwner resowner, bool *context_reused) ..pgstrom_init_cuda();->pg_strom/src/cuda_control.c/static void pgstrom_init_cuda(void) ..rc = cuInit(0);->pg_strom/src/cuda_control.c/static void pgstrom_init_cuda(void) ..cuDeviceGet(&device, ordinal);->pg_strom/src/cuda_control.c/static GpuContext *pgstrom_create_gpucontext(ResourceOwner resowner, bool *context_reused) ..rc = cuCtxCreate(&cuda_context_temp[index], CU_CTX_SCHED_AUTO, cuda_devices[index]);->pg_strom/src/cuda_control.c/static GpuContext *pgstrom_create_gpucontext(ResourceOwner resowner, bool *context_reused) ..rc = cuCtxSetCacheConfig(CU_FUNC_CACHE_PREFER_SHARED);->pg_strom/src/gpujoin.c/static void gpujoin_begin(CustomScanState *node, EState *estate, int eflags) ..gjs->gts.cb_task_process = gpujoin_task_process;->pg_strom/src/gpujoin.c/static void gpujoin_begin(CustomScanState *node, EState *estate, int eflags) ..pgstrom_load_cuda_program(&gjs->gts, true);->pg_strom/src/cuda_program.c/bool pgstrom_load_cuda_program(GpuTaskState *gts, bool is_preload) ..    cuda_modules = __pgstrom_load_cuda_program( gts->gcontext, gts->extra_flags, gts->kern_source, gts->kern_define, is_preload, true, &gts->pfm.tv_build_start, &gts->pfm.tv_build_end);->pg_strom/src/cuda_program.c/static CUmodule *__pgstrom_load_cuda_program( GpuContext *gcontext, cl_uint extra_flags, const char *kern_source, const char *kern_define, bool is_preload, bool with_async_build, struct timeval *tv_build_start, struct timeval *tv_build_end) ..worker.bgw_main = pgstrom_build_cuda_program_bgw_main;->pg_strom/src/cuda_program.c/static void pgstrom_build_cuda_program_bgw_main(Datum cuda_program) ..pgstrom_build_cuda_program(entry);->pg_strom/src/cuda_program.c/static void pgstrom_build_cuda_program(program_cache_entry *entry) ..__build_cuda_program(entry);->pg_strom/src/cuda_program.c/static void __build_cuda_program(program_cache_entry *old_entry) ..rc = nvrtcCreateProgram(&program, source, "pg_strom", 0, NULL, NULL);->pg_strom/src/cuda_program.c/static void __build_cuda_program(program_cache_entry *old_entry) ..rc = nvrtcCompileProgram(program, opt_index, options);->pg_strom/src/cuda_program.c/static void __build_cuda_program(program_cache_entry *old_entry) ..rc = nvrtcGetPTXSize(program, &ptx_length);->pg_strom/src/cuda_program.c/static void __build_cuda_program(program_cache_entry *old_entry) ..rc = nvrtcGetPTX(program, ptx_image);->pg_strom/src/cuda_program.c/static void __build_cuda_program(program_cache_entry *old_entry) ..rc = nvrtcGetProgramLogSize(program, &length);->pg_strom/src/cuda_program.c/static void __build_cuda_program(program_cache_entry *old_entry) ..rc = nvrtcGetProgramLog(program, build_log);->pg_strom/src/gpujoin.c/static void gpujoin_begin(CustomScanState *node, EState *estate, int eflags) ..pgstrom_init_gputaskstate(gcontext, &gjs->gts, estate);->pg_strom/src/gpujoin.c/static void gpujoin_begin(CustomScanState *node, EState *estate, int eflags) ..gjs->gts.cb_task_process = gpujoin_task_process;->pg_strom/src/gpujoin.c/static void gpujoin_begin(CustomScanState *node, EState *estate, int eflags) ..gjs->gts.cb_next_chunk = gpujoin_next_chunk;->pg_strom/src/gpujoin.c/static void gpujoin_begin(CustomScanState *node, EState *estate, int eflags) ..gjs->gts.cb_next_tuple = gpujoin_next_tuple;->pg_strom/src/gpujoin.c/static void gpujoin_begin(CustomScanState *node, EState *estate, int eflags) ..gjs->gts.cb_bulk_exec = pgstrom_exec_chunk_gputask;->src/backend/executor/execProcnode.c/PlanState *ExecInitNode(Plan *node, EState *estate, int eflags) ..case T_SeqScan:result = (PlanState *) ExecInitSeqScan((SeqScan *) node, estate, eflags);->src/backend/executor/nodeSeqscan.c/SeqScanState *ExecInitSeqScan(SeqScan *node, EState *estate, int eflags)->src/backend/commands/explain.c/void ExplainOnePlan(PlannedStmt *plannedstmt, IntoClause *into, ExplainState *es, const char *queryString, ParamListInfo params, const instr_time *planduration)..ExplainPrintPlan(es, queryDesc);->src/backend/commands/explain.c/void ExplainPrintPlan(ExplainState *es, QueryDesc *queryDesc) ..ExplainNode(queryDesc->planstate, NIL, NULL, NULL, es);->src/backend/commands/explain.c/static void ExplainNode(PlanState *planstate, List *ancestors, const char *relationship, const char *plan_name, ExplainState *es) ..case T_CustomScan:if (((Scan *) plan)->scanrelid > 0) ExplainScanTarget((Scan *) plan, es);->src/backend/commands/explain.c/static void ExplainScanTarget(Scan *plan, ExplainState *es) ..ExplainTargetRel((Plan *) plan, plan->scanrelid, es);->src/backend/commands/explain.c/static void ExplainTargetRel(Plan *plan, Index rti, ExplainState *es)->src/backend/tcop/pquery.c/bool PortalRun(Portal portal, long count, bool isTopLevel,DestReceiver *dest, DestReceiver *altdest,char *completionTag) ..nprocessed = PortalRunSelect(portal, true, count, dest);->src/backend/tcop/pquery.c/static long PortalRunSelect(Portal portal,bool forward,long count,DestReceiver *dest) ..ExecutorRun(queryDesc, direction, count);->src/backend/executor/execMain.c/void ExecutorRun(QueryDesc *queryDesc, ScanDirection direction, long count) ..standard_ExecutorRun(queryDesc, direction, count);->src/backend/executor/execMain.c/void standard_ExecutorRun(QueryDesc *queryDesc, ScanDirection direction, long count) ..ExecutePlan(estate, queryDesc->planstate, operation, sendTuples, count, direction, dest);->src/backend/executor/execMain.c/static void ExecutePlan(EState *estate,PlanState *planstate,CmdType operation,bool sendTuples,long numberTuples,ScanDirection direction,DestReceiver *dest) ..slot = ExecProcNode(planstate);->src/backend/executor/execProcnode.c/TupleTableSlot *ExecProcNode(PlanState *node) ..case T_SeqScanState:result = ExecSeqScan((SeqScanState *) node);->src/backend/executor/nodeSeqscan.c/TupleTableSlot *ExecSeqScan(SeqScanState *node) ..return ExecScan((ScanState *) node, (ExecScanAccessMtd) SeqNext, (ExecScanRecheckMtd) SeqRecheck);->src/backend/executor/nodeSeqscan.c/static TupleTableSlot *SeqNext(SeqScanState *node) ..tuple = heap_getnext(scandesc, direction);->src/backend/access/heap/heapam.c/HeapTuple heap_getnext(HeapScanDesc scan, ScanDirection direction) ..heapgettup_pagemode(scan, direction, scan->rs_nkeys, scan->rs_key);->src/backend/access/heap/heapam.c/static void heapgettup_pagemode(HeapScanDesc scan, ScanDirection dir, int nkeys, ScanKey key) ..heapgetpage(scan, page);->src/backend/access/heap/heapam.c/void heapgetpage(HeapScanDesc scan, BlockNumber page) ..scan->rs_cbuf = ReadBufferExtended(scan->rs_rd, MAIN_FORKNUM, page, RBM_NORMAL, scan->rs_strategy);->src/backend/storage/buffer/bufmgr.c/Buffer ReadBufferExtended(Relation reln, ForkNumber forkNum, BlockNumber blockNum, ReadBufferMode mode, BufferAccessStrategy strategy) ..buf = ReadBuffer_common(reln->rd_smgr, reln->rd_rel->relpersistence, forkNum, blockNum, mode, strategy, &hit);->src/backend/storage/buffer/bufmgr.c/static Buffer ReadBuffer_common(SMgrRelation smgr, char relpersistence, ForkNumber forkNum, BlockNumber blockNum, ReadBufferMode mode, BufferAccessStrategy strategy, bool *hit) ..smgrread(smgr, forkNum, blockNum, (char *) bufBlock);->src/backend/storage/smgr/smgr.c/void smgrread(SMgrRelation reln, ForkNumber forknum, BlockNumber blocknum, char *buffer) ..(*(smgrsw[reln->smgr_which].smgr_read)) (reln, forknum, blocknum, buffer);->src/backend/storage/smgr/md.c/void mdread(SMgrRelation reln, ForkNumber forknum, BlockNumber blocknum, char *buffer) ..nbytes = FileRead(v->mdfd_vfd, buffer, BLCKSZ);->src/backend/executor/nodeSeqscan.c/static TupleTableSlot *SeqNext(SeqScanState *node) ..ExecStoreTuple(tuple, slot, scandesc->rs_cbuf, false);->src/backend/executor/execProcnode.c/TupleTableSlot *ExecProcNode(PlanState *node) ..case T_CustomScanState:result = ExecCustomScan((CustomScanState *) node);->src/backend/executor/nodeCustom.c/TupleTableSlot *ExecCustomScan(CustomScanState *node) ..return node->methods->ExecCustomScan(node);->pg_strom/src/gpujoin.c/static TupleTableSlot *gpujoin_exec(CustomScanState *node) ..return ExecScan(&node->ss, (ExecScanAccessMtd) pgstrom_exec_gputask, (ExecScanRecheckMtd) pgstrom_recheck_gputask);->pg_strom/src/cuda_control.c/TupleTableSlot *pgstrom_exec_gputask(GpuTaskState *gts) ..while (!gts->curr_task || !(slot = gts->cb_next_tuple(gts)))->pg_strom/src/gpujoin.c/static TupleTableSlot *gpujoin_next_tuple(GpuTaskState *gts) ..pgstrom_fetch_data_store(slot, pds_dst, index, &gjs->curr_tuple);->pg_strom/src/datastore.c/bool pgstrom_fetch_data_store(TupleTableSlot *slot, pgstrom_data_store *pds, size_t row_index, HeapTuple tuple) ..return kern_fetch_data_store(slot, pds->kds, row_index, tuple);->pg_strom/src/datastore.c/bool kern_fetch_data_store(TupleTableSlot *slot, kern_data_store *kds, size_t row_index, HeapTuple tuple) ..ExecStoreTuple(tuple, slot, InvalidBuffer, false);->pg_strom/src/cuda_control.c/TupleTableSlot *pgstrom_exec_gputask(GpuTaskState *gts) ..gtask = pgstrom_fetch_gputask(gts);->pg_strom/src/cuda_control.c/GpuTask *pgstrom_fetch_gputask(GpuTaskState *gts) ..do {gtask = gts->cb_next_chunk(gts);launch_pending_tasks(gts);} while (waitfor_ready_tasks(gts));->pg_strom/src/cuda_control.c/static void launch_pending_tasks(GpuTaskState *gts) ..if (!pgstrom_load_cuda_program(gts, false))->pg_strom/src/cuda_program.c/bool pgstrom_load_cuda_program(GpuTaskState *gts, bool is_preload) ..cuda_modules = __pgstrom_load_cuda_program(gts->gcontext, gts->extra_flags, gts->kern_source, gts->kern_define, is_preload, true, &gts->pfm.tv_build_start, &gts->pfm.tv_build_end);->pg_strom/src/cuda_program.c/static CUmodule *__pgstrom_load_cuda_program( GpuContext *gcontext, cl_uint extra_flags, const char *kern_source, const char *kern_define, bool is_preload, bool with_async_build, struct timeval *tv_build_start, struct timeval *tv_build_end) ..worker.bgw_main = pgstrom_build_cuda_program_bgw_main;->pg_strom/src/cuda_program.c/static void pgstrom_build_cuda_program_bgw_main(Datum cuda_program) ..pgstrom_build_cuda_program(entry);->pg_strom/src/cuda_program.c/static void pgstrom_build_cuda_program(program_cache_entry *entry) ..__build_cuda_program(entry);->pg_strom/src/cuda_program.c/static void __build_cuda_program(program_cache_entry *old_entry) ..link_cuda_libraries(ptx_image, ptx_length, old_entry->extra_flags, &bin_image, &bin_length);->pg_strom/src/cuda_program.c/static void __build_cuda_program(program_cache_entry *old_entry) ..pgstrom_wakeup_backends(old_entry->waiting_backends);->pg_strom/src/cuda_program.c/static void __build_cuda_program(program_cache_entry *old_entry) ..pgstrom_put_cuda_program(old_entry);->pg_strom/src/gpujoin.c/static GpuTask *gpujoin_next_chunk(GpuTaskState *gts) ..pmrels_new = gpujoin_inner_getnext(gjs);->pg_strom/src/gpujoin.c/static pgstrom_multirels *gpujoin_inner_getnext(GpuJoinState *gjs) ..if (!gpujoin_inner_preload(gjs))->pg_strom/src/gpujoin.c/static bool gpujoin_inner_preload(GpuJoinState *gjs) ..if (!(istate->hash_inner_keys != NIL ? gpujoin_inner_hash_preload(gjs, istate, &total_usage) : gpujoin_inner_heap_preload(gjs, istate, &total_usage)))->pg_strom/src/gpujoin.c/static bool gpujoin_inner_hash_preload(GpuJoinState *gjs, innerState *istate, Size *p_total_usage) ..scan_slot = ExecProcNode(istate->state);->pg_strom/src/gpujoin.c/static bool gpujoin_inner_hash_preload(GpuJoinState *gjs, innerState *istate, Size *p_total_usage) ..foreach (lc, istate->pds_list) {pgstrom_data_store *pds = lfirst(lc); add_extra_randomness(pds); PDS_build_hashtable(pds);}->pg_strom/src/datastore.c/void PDS_build_hashtable(pgstrom_data_store *pds)->pg_strom/src/gpujoin.c/static pgstrom_multirels *gpujoin_inner_getnext(GpuJoinState *gjs) ..return gpujoin_create_multirels(gjs);->pg_strom/src/gpujoin.c/static pgstrom_multirels *gpujoin_create_multirels(GpuJoinState *gjs)->pg_strom/src/gpujoin.c/static GpuTask *gpujoin_next_chunk(GpuTaskState *gts) ..pds = pgstrom_exec_scan_chunk(gts, pgstrom_chunk_size());->pg_strom/src/gpuscan.c/pgstrom_data_store *pgstrom_exec_scan_chunk(GpuTaskState *gts, Size chunk_length) ..if (PDS_insert_block(pds, base_rel, scan->rs_cblock, scan->rs_snapshot, scan->rs_strategy) < 0)->pg_strom/src/datastore.c/int PDS_insert_block(pgstrom_data_store *pds, Relation rel, BlockNumber blknum, Snapshot snapshot, BufferAccessStrategy strategy) ..buffer = ReadBufferExtended(rel, MAIN_FORKNUM, blknum, RBM_NORMAL, strategy);->src/backend/storage/buffer/bufmgr.c/Buffer ReadBufferExtended(Relation reln, ForkNumber forkNum, BlockNumber blockNum, ReadBufferMode mode, BufferAccessStrategy strategy) ..buf = ReadBuffer_common(reln->rd_smgr, reln->rd_rel->relpersistence, forkNum, blockNum, mode, strategy, &hit);->src/backend/storage/buffer/bufmgr.c/static Buffer ReadBuffer_common(SMgrRelation smgr, char relpersistence, ForkNumber forkNum, BlockNumber blockNum, ReadBufferMode mode, BufferAccessStrategy strategy, bool *hit) ..smgrread(smgr, forkNum, blockNum, (char *) bufBlock);->src/backend/storage/smgr/smgr.c/void smgrread(SMgrRelation reln, ForkNumber forknum, BlockNumber blocknum, char *buffer) ..(*(smgrsw[reln->smgr_which].smgr_read)) (reln, forknum, blocknum, buffer);->src/backend/storage/smgr/md.c/void mdread(SMgrRelation reln, ForkNumber forknum, BlockNumber blocknum, char *buffer) ..nbytes = FileRead(v->mdfd_vfd, buffer, BLCKSZ);->pg_strom/src/gpujoin.c/static GpuTask *gpujoin_next_chunk(GpuTaskState *gts) ..return gpujoin_create_task(gjs, gjs->curr_pmrels, pds, NULL);->pg_strom/src/gpujoin.c/static GpuTask *gpujoin_create_task(GpuJoinState *gjs, pgstrom_multirels *pmrels, pgstrom_data_store *pds_src, kern_join_scale *jscale_old) ..pgstrom_init_gputask(&gjs->gts, &pgjoin->task);->pg_strom/src/gpujoin.c/static GpuTask *gpujoin_create_task(GpuJoinState *gjs, pgstrom_multirels *pmrels, pgstrom_data_store *pds_src, kern_join_scale *jscale_old) ..multirels_get_buffer(pmrels, pgjoin);->pg_strom/src/cuda_control.c/static void launch_pending_tasks(GpuTaskState *gts) ..launch = gts->cb_task_process(gtask);->pg_strom/src/gpujoin.c/static bool gpujoin_task_process(GpuTask *gtask) ..rc = cuCtxPushCurrent(gtask->cuda_context);->pg_strom/src/gpujoin.c/static bool gpujoin_task_process(GpuTask *gtask) ..multirels_get_buffer(pgjoin->pmrels, pgjoin)->pg_strom/src/gpujoin.c/static bool multirels_get_buffer(pgstrom_multirels *pmrels, pgstrom_gpujoin *pgjoin) ..m_kmrels = gpuMemAlloc(&pgjoin->task, pmrels->usage_length);->pg_strom/src/cuda_control.c/CUdeviceptr gpuMemAlloc(GpuTask *gtask, size_t bytesize) ..return __gpuMemAlloc(gtask->gts->gcontext, gtask->cuda_index, bytesize);->pg_strom/src/cuda_control.c/CUdeviceptr __gpuMemAlloc(GpuContext *gcontext, int cuda_index, size_t bytesize) ..rc = cuMemAlloc(&block_addr, required);->pg_strom/src/gpujoin.c/static bool multirels_get_buffer(pgstrom_multirels *pmrels, pgstrom_gpujoin *pgjoin) ..m_ojmaps = gpuMemAlloc(&pgjoin->task, length);->pg_strom/src/cuda_control.c/CUdeviceptr gpuMemAlloc(GpuTask *gtask, size_t bytesize) ..return __gpuMemAlloc(gtask->gts->gcontext, gtask->cuda_index, bytesize);->pg_strom/src/cuda_control.c/CUdeviceptr __gpuMemAlloc(GpuContext *gcontext, int cuda_index, size_t bytesize) ..rc = cuMemAlloc(&block_addr, required);->pg_strom/src/gpujoin.c/static bool gpujoin_task_process(GpuTask *gtask) ..status = __gpujoin_task_process(pgjoin);->pg_strom/src/gpujoin.c/static bool __gpujoin_task_process(pgstrom_gpujoin *pgjoin) ..rc = cuModuleGetFunction(&pgjoin->kern_main, pgjoin->task.cuda_module, "gpujoin_main");->pg_strom/src/cuda_gpujoin.h/KERNEL_FUNCTION(void) gpujoin_main(kern_gpujoin *kgjoin, kern_multirels *kmrels, cl_bool *outer_join_map, kern_data_store *kds_src, kern_data_store *kds_dst, cl_int cuda_index)->pg_strom/src/gpujoin.c/static bool __gpujoin_task_process(pgstrom_gpujoin *pgjoin) ..pgjoin->m_kgjoin = gpuMemAlloc(&pgjoin->task, total_length);->pg_strom/src/cuda_control.c/CUdeviceptr gpuMemAlloc(GpuTask *gtask, size_t bytesize) ..return __gpuMemAlloc(gtask->gts->gcontext, gtask->cuda_index, bytesize);->pg_strom/src/cuda_control.c/CUdeviceptr __gpuMemAlloc(GpuContext *gcontext, int cuda_index, size_t bytesize) ..rc = cuMemAlloc(&block_addr, required);->pg_strom/src/gpujoin.c/static bool __gpujoin_task_process(pgstrom_gpujoin *pgjoin) ..CUDA_EVENT_CREATE(pgjoin, ev_dma_send_start);->pg_strom/src/pg_strom.h/CUresult __rc = cuEventCreate(&(node)->ev_field, CU_EVENT_DEFAULT);->pg_strom/src/gpujoin.c/static bool __gpujoin_task_process(pgstrom_gpujoin *pgjoin) ..CUDA_EVENT_CREATE(pgjoin, ev_dma_send_stop);->pg_strom/src/pg_strom.h/CUresult __rc = cuEventCreate(&(node)->ev_field, CU_EVENT_DEFAULT);->pg_strom/src/gpujoin.c/static bool __gpujoin_task_process(pgstrom_gpujoin *pgjoin) ..CUDA_EVENT_CREATE(pgjoin, ev_dma_recv_start);->pg_strom/src/pg_strom.h/CUresult __rc = cuEventCreate(&(node)->ev_field, CU_EVENT_DEFAULT);->pg_strom/src/gpujoin.c/static bool __gpujoin_task_process(pgstrom_gpujoin *pgjoin) ..CUDA_EVENT_CREATE(pgjoin, ev_dma_recv_stop);->pg_strom/src/pg_strom.h/CUresult __rc = cuEventCreate(&(node)->ev_field, CU_EVENT_DEFAULT);->pg_strom/src/gpujoin.c/static bool __gpujoin_task_process(pgstrom_gpujoin *pgjoin) ..CUDA_EVENT_RECORD(pgjoin, ev_dma_send_start);->pg_strom/src/pg_strom.h/CUresult __rc = cuEventRecord((node)->ev_field, ((GpuTask *)(node))->cuda_stream);->pg_strom/src/gpujoin.c/static bool __gpujoin_task_process(pgstrom_gpujoin *pgjoin) ..multirels_send_buffer(pmrels, &pgjoin->task);->pg_strom/src/gpujoin.c/static void multirels_send_buffer(pgstrom_multirels *pmrels, GpuTask *gtask) ..rc = cuEventCreate(&ev_loaded, CU_EVENT_DEFAULT);->pg_strom/src/gpujoin.c/static void multirels_send_buffer(pgstrom_multirels *pmrels, GpuTask *gtask) ..rc = cuMemcpyHtoDAsync(m_kmrels, &pmrels->kern, length, cuda_stream);->pg_strom/src/gpujoin.c/static void multirels_send_buffer(pgstrom_multirels *pmrels, GpuTask *gtask) ..rc = cuMemcpyHtoDAsync(m_kmrels + offset, kds, kds->length, cuda_stream);->pg_strom/src/gpujoin.c/static void multirels_send_buffer(pgstrom_multirels *pmrels, GpuTask *gtask) ..rc = cuEventRecord(ev_loaded, cuda_stream);->pg_strom/src/gpujoin.c/static bool __gpujoin_task_process(pgstrom_gpujoin *pgjoin) ..rc = cuMemcpyHtoDAsync(pgjoin->m_kgjoin, &pgjoin->kern, length, pgjoin->task.cuda_stream);->pg_strom/src/gpujoin.c/static bool __gpujoin_task_process(pgstrom_gpujoin *pgjoin) ..rc = cuMemcpyHtoDAsync(pgjoin->m_kds_src, pds_src->kds, length, pgjoin->task.cuda_stream);->pg_strom/src/gpujoin.c/static bool __gpujoin_task_process(pgstrom_gpujoin *pgjoin) ..rc = cuMemcpyHtoDAsync(pgjoin->m_kds_dst, pds_dst->kds, length, pgjoin->task.cuda_stream);->pg_strom/src/gpujoin.c/static bool __gpujoin_task_process(pgstrom_gpujoin *pgjoin) ..CUDA_EVENT_RECORD(pgjoin, ev_dma_send_stop);->pg_strom/src/pg_strom.h/CUresult __rc = cuEventRecord((node)->ev_field, ((GpuTask *)(node))->cuda_stream);->pg_strom/src/gpujoin.c/static bool __gpujoin_task_process(pgstrom_gpujoin *pgjoin) ..rc = cuLaunchKernel(pgjoin->kern_main, 1, 1, 1, 1, 1, 1, sizeof(kern_errorbuf), pgjoin->task.cuda_stream, kern_args, NULL);->pg_strom/src/gpujoin.c/static bool __gpujoin_task_process(pgstrom_gpujoin *pgjoin) ..CUDA_EVENT_RECORD(pgjoin, ev_dma_recv_start);->pg_strom/src/pg_strom.h/CUresult __rc = cuEventRecord((node)->ev_field, ((GpuTask *)(node))->cuda_stream);->pg_strom/src/gpujoin.c/static bool __gpujoin_task_process(pgstrom_gpujoin *pgjoin) ..rc = cuMemcpyDtoHAsync(&pgjoin->kern, pgjoin->m_kgjoin, length, pgjoin->task.cuda_stream);->pg_strom/src/gpujoin.c/static bool __gpujoin_task_process(pgstrom_gpujoin *pgjoin) ..rc = cuMemcpyDtoHAsync(pds_dst->kds, pgjoin->m_kds_dst, length, pgjoin->task.cuda_stream);->pg_strom/src/gpujoin.c/static bool __gpujoin_task_process(pgstrom_gpujoin *pgjoin) ..CUDA_EVENT_RECORD(pgjoin, ev_dma_recv_stop);->pg_strom/src/pg_strom.h/CUresult __rc = cuEventRecord((node)->ev_field, ((GpuTask *)(node))->cuda_stream);->pg_strom/src/gpujoin.c/static bool __gpujoin_task_process(pgstrom_gpujoin *pgjoin) ..rc = cuStreamAddCallback(pgjoin->task.cuda_stream, gpujoin_task_respond, pgjoin, 0);->pg_strom/src/gpujoin.c/static bool gpujoin_task_process(GpuTask *gtask) ..rc = cuCtxPopCurrent(NULL);->src/backend/executor/execProcnode.c/TupleTableSlot *ExecProcNode(PlanState *node) ..case T_AggState:result = ExecAgg((AggState *) node);->src/backend/executor/nodeAgg.c/TupleTableSlot *ExecAgg(AggState *node) ..->src/backend/tcop/postgres.c/PostgresMain(int argc, char *argv[],const char *dbname,const char *username) ..case 'S':

                                                 --致敬异构数据库领域先驱KaiGai Kohei

PostgreSQL异构引擎pg-strom源码解读相关推荐

PostgreSQL 源码解读（160）- 查询#80(如何实现表达式解析)
本节介绍了PostgreSQL如何解析查询语句中的表达式列并计算得出该列的值.表达式列是指除关系定义中的系统列/定义列之外的其他投影列.比如: testdb=# create table t_expr ...
PostgreSQL 源码解读（147）- Storage Manager#3（fsm_search函数）
本节简单介绍了PostgreSQL在执行插入过程中与存储相关的函数RecordAndGetPageWithFreeSpace->fsm_search,该函数搜索FSM,找到有足够空闲空间(min ...
PostgreSQL 源码解读（156）- 后台进程#8（walsender#4）
上节介绍了PostgreSQL的后台进程walsender中的函数WalSndLoop->WaitLatchOrSocket->WaitEventSetWait->WaitEvent ...
PostgreSQL 源码解读（154）- 后台进程#6（walsender#2）
本节继续介绍PostgreSQL的后台进程walsender,重点介绍的是调用栈中的exec_replication_command和StartReplication函数. 调用栈如下: (gdb) ...
PostgreSQL 源码解读（155）- 后台进程#7（walsender#3）
本节继续介绍PostgreSQL的后台进程walsender,重点介绍的是调用栈中的函数WalSndLoop->WaitLatchOrSocket->WaitEventSetWait-&g ...
PostgreSQL 源码解读（153）- 后台进程#5（walsender#1）
本节简单介绍了PostgreSQL的后台进程walsender,该进程实质上是streaming replication环境中master节点上普通的backend进程,在standby节点启动时,s ...
PostgreSQL 源码解读（31）- 查询语句#16（查询优化-表达式预处理#1）
本节简单介绍了PG查询优化对表达式预处理中连接Var(RTE中的Var,其中RTE_KIND=RTE_JOIN)溯源的过程.处理逻辑在主函数subquery_planner中通过调用flatten_j ...
PostgreSQL 源码解读（35）- 查询语句#20（查询优化-简化Having和Grou...
本节简单介绍了PG查询优化中对Having和Group By子句的简化处理. 一.基本概念简化Having语句把Having中的约束条件,如满足可以提升到Where条件中的,则移动到Where子句 ...
PostgreSQL 源码解读（212）- 后台进程#11（checkpointer-SyncOneBuffer）
本节介绍了checkpoint中用于刷一个脏page的函数:SyncOneBuffer,该函数在syncing期间处理一个buffer. 一.数据结构宏定义 checkpoints request ...

PostgreSQL异构引擎pg-strom源码解读

PostgreSQL异构引擎pg-strom源码解读

PostgreSQL异构引擎pg-strom源码解读相关推荐

最新文章

热门文章