Changeset: 9d0af5c9ca3f for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=9d0af5c9ca3f
Modified Files:
monetdb5/extras/bwd/91_bwd.sql
monetdb5/extras/bwd/bwd.c
monetdb5/extras/bwd/operations.c
monetdb5/extras/bwd/opt_bwd.mal
monetdb5/extras/bwd/optimizer.c
monetdb5/extras/bwd/utilities.c
monetdb5/extras/bwd/utilities.h
Branch: bwd
Log Message:
* introduced lot's of hacks to get tpch q14 running (q1, q6 and q14 run through
now)
Unterschiede (gekürzt von 587 auf 300 Zeilen):
diff --git a/monetdb5/extras/bwd/91_bwd.sql b/monetdb5/extras/bwd/91_bwd.sql
--- a/monetdb5/extras/bwd/91_bwd.sql
+++ b/monetdb5/extras/bwd/91_bwd.sql
@@ -20,7 +20,7 @@ returns varchar(4096) external name bwd.
create function bwdecomposeint(col integer, bits integer)
returns varchar(4096) external name bwd.decompose;
-create function bwdecomposefk(schema varchar(24), tabl varchar(24), columnn
varchar(24), bits integer)
+create function bwdecomposefk(scema varchar(24), tabl varchar(24), columnn
varchar(24), bits integer)
returns varchar(4096) external name bwd.decomposefk;
diff --git a/monetdb5/extras/bwd/bwd.c b/monetdb5/extras/bwd/bwd.c
--- a/monetdb5/extras/bwd/bwd.c
+++ b/monetdb5/extras/bwd/bwd.c
@@ -118,6 +118,9 @@ str bwdecompose(bat * res, bat * subject
} else if (Tsize(subject) == 4){
BATsetprop(subject, batRegistryIndex, TYPE_int,
(int[]){decomposeIntArray((int*)Tloc(subject, BUNfirst(subject)),
subject->batCount, *approximationBits)});
snprintf(buffer, 4096, "successfully decomposed integer bat
%d", subject->batCacheid);
+ } else if (BATttype(subject) == TYPE_oid && !BATtvoid(subject)){
+ BATsetprop(subject, batRegistryIndex, TYPE_lng,
(lng[]){transferOIDArray((oid*)Tloc(subject, BUNfirst(subject)),
subject->batCount, *approximationBits)});
+ snprintf(buffer, 4096, "successfully decomposed long integer
bat %d", subject->batCacheid);
} else if (Tsize(subject) == 8){
BATsetprop(subject, batRegistryIndex, TYPE_lng,
(lng[]){decomposeLngArray((lng*)Tloc(subject, BUNfirst(subject)),
subject->batCount, *approximationBits)});
snprintf(buffer, 4096, "successfully decomposed long integer
bat %d", subject->batCacheid);
diff --git a/monetdb5/extras/bwd/operations.c b/monetdb5/extras/bwd/operations.c
--- a/monetdb5/extras/bwd/operations.c
+++ b/monetdb5/extras/bwd/operations.c
@@ -68,6 +68,7 @@ clTail getTailHeader(cl_mem memoryObject
clHead* getPositionsColumn(cl_mem memoryObject, clHead* buffer, size_t*
bufferSize){
cl_int err = 0;
+ assert(memoryObject);
if(!buffer){
if(eagerBufferLoading){
clGetMemObjectInfo(memoryObject, CL_MEM_SIZE,
sizeof(size_t), bufferSize, NULL);
@@ -154,13 +155,19 @@ size_t calculatedBufferSize(size_t headC
#define NOT_DECOMPOSED GDKstrdup("input bat not decomposed")
str BWDLeftJoinApproximate(bat * res, bat * l, bat * r){
BAT *left = BATdescriptor(*l), *right = BATdescriptor(*r);
- if(!batTailIsDecomposed(right)){
+ str result = NOT_IMPLEMENTED;
+ if(!batTailIsDecomposed(right) && !(BAThvoid(right) &&
BATtvoid(right))){
printf("bwd.%s : bat is not decomposed: %d, no approximation
can be provided\n", __func__, *r);
BBPreleaseref(right->batCacheid);
return MAL_SUCCEED;
}
-
- if(BAThvoid(left) && BAThvoid(right) && right->tseqbase != oid_nil){
+ if(BAThvoid(left) && BAThvoid(right) && right->T->type == TYPE_void){
+ BBPincref(*res = *l, 1);
+ result = MAL_SUCCEED;
+ } else if(BAThvoid(left) && left->T->type == TYPE_void &&
BAThvoid(right)){
+ BBPincref(*res = *r, 1);
+ result = MAL_SUCCEED;
+ } else if(BAThvoid(left) && BAThvoid(right) && right->tseqbase !=
oid_nil){
cl_int err;
int headCount; // :-)
cl_kernel projectKernel;
@@ -170,6 +177,7 @@ str BWDLeftJoinApproximate(bat * res, ba
int i;
if(!leftColumn) leftColumn =
batHeadApproximation(BATmirror(left));
if(!rightColumn) rightColumn =
batHeadApproximation(BATmirror(right));
+ assert(leftColumn && rightColumn);
projectKernel =
clCreateKernel(getProjectionLeftjoinProgram(batTailApproximationBits(right),
batTailOffsetBits(right)), "project", &err);
if(err) printf("#%s, clCreateKernel: %s;\n", __func__,
clError(err));
@@ -211,12 +219,21 @@ str BWDLeftJoinApproximate(bat * res, ba
printf("#%s, clEnqueueNDRangeKernel: %s, %s:%u;\n",
__func__, clError(err), __FILE__, __LINE__);
if (synchronousGPU) clFinish(getCommandQueue());
BBPkeepref((*res = result->batCacheid));
+ return MAL_SUCCEED;
+ } else {
+ printf("bwd.BWDLeftJoinApproximate, %s (BAThvoid(left): %d,
BATtvoid(left): %d, BAThvoid(right): %d, right->tseqbase: %ld)", "this case
isn't implemented yet", BAThvoid(left), BATtvoid(left), BATtvoid(right),
left->tseqbase);
+ }
+ {
+ BAT* approximation = BATdescriptor(*res);
+ if(batTailApproximation(approximation))
assert(batTailPositions(approximation));
+
+ BBPreleaseref(*res);
+ }
BBPreleaseref(left->batCacheid);
BBPreleaseref(right->batCacheid);
- return MAL_SUCCEED;
- }
- printf("bwd.BWDLeftJoinApproximate, %s (BAThvoid(left): %d,
BATtvoid(left): %d, BAThvoid(right): %d, right->tseqbase: %ld)", "this case
isn't implemented yet", BAThvoid(left), BATtvoid(left), BATtvoid(right),
left->tseqbase);
- return NOT_IMPLEMENTED;
+
+
+ return result;
};
static inline int decompress_int_Value(const int approximationI, const int
approximationBits, const int offsetBits, const clTail* compressedTail, const
unsigned int* residuals, const unsigned int residualI){
@@ -256,6 +273,23 @@ static inline int decompress_int_Value(c
return deCompressedValue;
}
+static inline lng decompress_lng_Value_without_residual(const int
approximationI, const int approximationBits, const clTail* compressedTail){
+ const int approximationMask = (1<<(approximationBits))-1;
+ const unsigned int* vals = (unsigned int*) compressedTail->elements;
+ const unsigned int targetTypeBits = sizeof(int)*8;
+ const unsigned int slotI =
(approximationBits)*approximationI/(sizeof(int)*8);
+ const unsigned int offset =
((approximationBits)*approximationI)%(sizeof(int)*8);
+
+ const unsigned int compressedValue =
+ (((offset+approximationBits)<=targetTypeBits)*
+ (vals[slotI] >> (targetTypeBits-offset-approximationBits)) &
approximationMask)
+ + (((offset+approximationBits)>targetTypeBits)?
+
(((vals[slotI]<<(approximationBits-targetTypeBits+offset))
+ +
(vals[slotI+1]>>(targetTypeBits-(approximationBits-targetTypeBits+offset))))&approximationMask):0);
+
+ return compressedTail->base+ compressedValue;
+}
+
static inline int decompress_int_Value_without_residual(const int
approximationI, const int approximationBits, const clTail* compressedTail){
const int approximationMask = (1<<(approximationBits))-1;
const unsigned int* vals = (unsigned int*) compressedTail->elements;
@@ -336,6 +370,7 @@ str BWDLeftJoinRefine(bat * res, bat * l
BAT* right = BATdescriptor(*r);
BAT* refinement = NULL;
BAT* approximation = BATdescriptor(*approx);
+ str result = MAL_SUCCEED;
const size_t approximationBits =
batTailApproximationBits(right);
const size_t offsetBits = batTailOffsetBits(right);
@@ -343,22 +378,12 @@ str BWDLeftJoinRefine(bat * res, bat * l
/* BATnew(approximation->htype, approximation->ttype,
approximation->batCount); */
- if(0 && BAThvoid(left) && BATtvoid(left) && BAThvoid(right) &&
right->tseqbase != oid_nil){
- // not very likely optimization, hence disabled
- const unsigned char* residuals =
batTailResiduals(right);
- const unsigned int residualBytes =
(32-approximationBits)/8;
- uint i;
- register int* outputRegion = (int*) Tloc(refinement,
BUNfirst(refinement));
- const int* approximationRegion = (int*)
Tloc(approximation, BUNfirst(approximation));
- const size_t offset = left->tseqbase;
-
- if(!batTailIsDecomposed(right))
- throw (MAL, "bwd.BWDLeftJoinApproximate", "bat
is not decomposed: %d", *r);
- refinement = BATnew(ATOMtype(left->htype),
ATOMtype(right->ttype), left->batCount);
- printf ("%s returning bat of type [%s, %s]\n",
__func__, typeNames[BAThtype(refinement)], typeNames[BATttype(refinement)]);
- BATsetcount(refinement, approximation->batCount);
- for (i = 0; i < approximation->batCount; ++i)
- outputRegion[i] = approximationRegion[i] +
((*((unsigned int*) (residuals + (i+offset)*residualBytes))) >>
approximationBits);
+ if(BAThvoid(left) && BAThvoid(right) && BATtvoid(right)){
+ BBPincref(*res = *l, 1);
+ result = MAL_SUCCEED;
+ } else if(BAThvoid(left) && left->T->type == TYPE_void &&
BAThvoid(right)){
+ BBPincref(*res = *r, 1);
+ result = MAL_SUCCEED;
} else if(BAThvoid(right) && right->tseqbase != oid_nil){
// we essentially perform a merge-join of the precise
positions (left tail) and the approximate positions to get the approximate
values of the super result set
// this is okay because the approximate positions are a
superset of the precise positions and the position lists have the same
permutation
@@ -377,16 +402,19 @@ str BWDLeftJoinRefine(bat * res, bat * l
if(batTailResidualBits(approximation) == 0){
#define expansionLoop(type) {
\
- unsigned int i;
\
+ unsigned int i;
\
type* refinementRegion = (type*)
Tloc(refinement, \
BUNfirst(refinement)); \
for (i = 0; i< left->batCount; ++i)
\
refinementRegion[i] =
decompress_##type##_Value_without_residual( \
i, approximationBits,
supersetApproximateValuesColumn \
);
\
- }
+ }
switch(right->T->width){
+ case 8:
+ expansionLoop(lng);
+ break;
case 4:
expansionLoop(int);
break;
@@ -408,53 +436,57 @@ str BWDLeftJoinRefine(bat * res, bat * l
#define refinementLoop(type) {
\
type* refinementRegion = (type*)
Tloc(refinement, \
-
BUNfirst(refinement)); \
+
BUNfirst(refinement)); \
for (i = 0; i <
supersetPositionsColumn->count; ++i) { \
if(supersetPositionsColumn->positions[i] == *positionRegion){ \
refinementRegion[refinementCount++] =
\
-
decompress_##type##_Value(i, approximationBits, offsetBits, \
-
supersetApproximateValuesColumn, \
-
residuals,
*positionRegion); \
- positionRegion++;
\
+
decompress_##type##_Value(i, approximationBits, offsetBits, \
+
supersetApproximateValuesColumn, \
+
residuals, *positionRegion); \
+
positionRegion++;
\
}
\
}
\
}
switch(right->T->width){
- case 4:
- refinementLoop(int);
- break;
- case 1:
- refinementLoop(char);
- break;
- default:
- throw (MAL, "bwd problem", "this type
(%d-byte) for %s is not implemented, %s:%d", ATOMsize(right->ttype), __func__,
__FILE__, __LINE__);
- }
+ case 4:
+ refinementLoop(int);
+ break;
+ case 1:
+ refinementLoop(char);
+ break;
+ default:
+ throw (MAL, "bwd problem", "this type (%d-byte)
for %s is not implemented, %s:%d", ATOMsize(right->ttype), __func__, __FILE__,
__LINE__);
+ }
#undef refinementLoop
BATsetcount(refinement, refinementCount);
BATseqbase(refinement, 0);
}
- GDKfree(supersetPositionsColumn);
- GDKfree(supersetApproximateValuesColumn);
+ GDKfree(supersetPositionsColumn);
+ GDKfree(supersetApproximateValuesColumn);
- if(!BAThvoid(left)){ // propagate left head column
+ if(!BAThvoid(left)){ // propagate left head
column
BATseqbase(BATmirror(refinement), oid_nil);
refinement = VIEWcreate(BATmark(left,
oid_nil),refinement);
}
- if(BATttype(right) == TYPE_str){ // propagate right heap
+ if(BATttype(right) == TYPE_str){ // propagate
right heap
if (right->batRestricted == BAT_READ) {
- assert(right->T->vheap->parentid > 0);
- BBPshare(right->T->vheap->parentid);
- refinement->T->vheap = right->T->vheap;
- }
+ assert(right->T->vheap->parentid > 0);
+ BBPshare(right->T->vheap->parentid);
+ refinement->T->vheap = right->T->vheap;
}
-
+ }
+ if(refinement)
+ BBPkeepref(*res =
refinement->batCacheid);
+ } else {
+ printf("unimplemented case in %s, %s:%d", __func__,
__FILE__, __LINE__);
+ abort();
}
-
- BBPkeepref(*res = refinement->batCacheid);
BBPreleaseref(approximation->batCacheid);
- return MAL_SUCCEED;
+ BBPreleaseref(left->batCacheid);
+ BBPreleaseref(right->batCacheid);
+ return result;
}
};
@@ -480,7 +512,7 @@ static inline str uselect(bat *res, bat
assert(ATOMstorage(BATttype(data)) == TYPE_int); // type specific
if(!batTailIsDecomposed(data))
throw (MAL, "bwd problem", "while running %s, I noticed that
bat is not decomposed: %s", __func__, resolveBatToAttribute(data->batCacheid,
cntxt));
- result = BATnew(BAThtype(data), tailType, 0); //TODO: not sure if we
actually need to allocate space here (count could be 0)
+ result = BATnew(data->H->type==TYPE_void?TYPE_oid:data->H->type,
tailType, 0); //TODO: not sure if we actually need to allocate space here
(count could be 0)
{
const clTail dataHeader =
getTailHeader(batTailApproximation(data));
@@ -806,37 +838,37 @@ str BWDThetauselectRefine(bat *res, bat
}
str BWDtmark(int *res, int *bid, oid *base){
str returnValue = ALGtmark(res, bid, base);
- if(*res){ // propagate head approximation
+ BAT* input = BATdescriptor(*bid);
+ if(*res && input->H->type != TYPE_void){ // propagate head approximation
cl_int err;
- BAT* input;
const unsigned int newIndex =
getNextFreeDecomposedBATSlotIndex();
DecomposedBATSlot* slot =
getDecomposedBATSlotForIndex(newIndex);
BAT* result = BATdescriptor(*res);
BATsetprop(result, batRegistryIndex, TYPE_int,
(int[]){newIndex});
BBPreleaseref(*res);
- input = BATdescriptor(*bid);
slot->headApproximation = batHeadApproximation(input);
- if(0) printf ("retaining %p\n", slot->headApproximation);
+ if(!slot->headApproximation) slot->headApproximation =
batTailPositions(BATmirror(input));
+ assert(slot->headApproximation);
+
+ if(1) printf ("retaining %p\n", slot->headApproximation);
if((err=clRetainMemObject(slot->headApproximation)))
printf("#%s, clRetainMemObject: %s;\n", __func__,
clError(err));
-
- BBPreleaseref(*bid);
}
+ BBPreleaseref(*bid);
return returnValue;
}
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
