add device-quals pull-up on bulkloading
pull-up outer-quals on gpupreagg

This enhancement allows to pull-up qualifiers of outer relation on
GpuPreAgg if underlying GpuScan has no host-only qualifiers.
It reduces number of opencl interaction and improve data density
of PCI-E data transfer.

postgres=# explain select cat, avg(x) from t0 where x < 50.0 group by cat;
INFO:  Agg cost = 379300.23..379300.31 nrows (10048333 => 179)
                                    QUERY PLAN
-----------------------------------------------------------------------------------
 HashAggregate  (cost=379300.31..379300.23 rows=7 width=12)
   Group Key: cat
   ->  Custom (GpuPreAgg)  (cost=10335.27..379299.33 rows=179 width=112)
         Bulkload: On
         Device Filter: (x < 50::double precision)
         ->  Custom (GpuScan) on t0  (cost=0.00..266667.00 rows=20000024 width=12)
 Planning time: 2.130 ms
(7 rows)
cosmetic change: kills warnning because of unused variables
revert 6b7fdc7da17b974d5ab04719a7a559076a03b0ce.

These were variables in used only when --enable-cassert.
This fixup changes these variables are not used, but equivalent checks.
bugfix: MemoryContextAllocHuge() instead of palloc() for large chunk

multihash_preload_khashtable may allocate more than 1GB chunk if
zone-length and table datasize are large, however, palloc()
internally has hard limitation by 1GB.
So, we had to use MemoryContextAllocHuge() to allocate more than
the boundary.
bugfix: bitonic_sorting logic didn't pay attention for kernel resource consumption
code cleanup on bitonic_compute_workgroup_size
fixbug: gpuscan exited relation scan unintentionally.

NOTE: In case when it scans on a large hole (that is
continuous blocks contain invisible tuples only; may
be created by DELETE with relaxed condition),
pgstrom_data_store_insert_block() may return negative
value without valid tuples, even though we don't reach
either end of relation or chunk.
So, we need to check whether we actually touched on
the end-of-relation. If not, retry scanning.
bugfix: plan->qual contains bare expressions, not RestrictInfo, once PlannedStmt gets constructed.
bugfix: bulk-loading from sub-plan with projection didn't work.

bulk-loading mode runs projection of the sub-plan in kernel space,
so we have to make a projection if CpuReCheck error leads fallback
routine.
In this case, input tuple-slot shall be set on the ExprContext of
the ProjectionInfo of the "sub-plan".
bugfix: GpuHashJoin didn't compute correct hash-value if data-width of both side are not identical.

todo: more graceful hash-join decision is needed on integration
with upstreamed custom-scan interface.
add #ifdef CL_VERSION_x_x to avoid deprecated / non supported functions
Merge branch 'master' of ../pg_strom