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 + Global
GPU Projection: t0.cat, t1.ax
Logic Parameter: SafetyLimit: 2147483647, KeyDistSalt: 9
Extra: outer-bulk-exec
-> Custom Scan (GpuJoin) on public.t0 (cost=25099.24..2826221.34 rows=93721454 width=12)
Output: t0.cat, t1.ax
GPU Projection: t0.cat::text, t1.ax::double precision
Depth 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, >s->pfm.tv_build_start, >s->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, >s->pfm.tv_build_start, >s->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
扫一扫
1614
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
