Changeset: 4af8c47a665c for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=4af8c47a665c
Modified Files:
monetdb5/extras/bwd/cl_program_utilities.c
monetdb5/extras/bwd/cl_program_utilities.h
monetdb5/extras/bwd/operations.c
monetdb5/extras/bwd/utilities.c
Branch: bwd
Log Message:
* started implementation of approximate grouping
Unterschiede (196 Zeilen):
diff --git a/monetdb5/extras/bwd/cl_program_utilities.c
b/monetdb5/extras/bwd/cl_program_utilities.c
--- a/monetdb5/extras/bwd/cl_program_utilities.c
+++ b/monetdb5/extras/bwd/cl_program_utilities.c
@@ -235,3 +235,7 @@ cl_program getUSelectProgram(int type, c
return program;
}
}
+
+cl_program getGroupProgram(unsigned int approximationBits){
+ return NULL;
+}
diff --git a/monetdb5/extras/bwd/cl_program_utilities.h
b/monetdb5/extras/bwd/cl_program_utilities.h
--- a/monetdb5/extras/bwd/cl_program_utilities.h
+++ b/monetdb5/extras/bwd/cl_program_utilities.h
@@ -10,6 +10,7 @@
cl_program getUSelectProgram(int type, char* predicateOperation, char*
predicateOperation2, unsigned int approximationBits, unsigned int offsetBits,
char inputIsVoidHeaded);
cl_program getProjectionLeftjoinProgram(unsigned int approximationBits,
unsigned int offsetBits);
+cl_program getGroupProgram(unsigned int approximationBits);
cl_program compileProgram(const char* sourceCode, char* options);
/* cl_program getProjectionSemijoinProgram(unsigned int approximationBits); */
#endif /* _CL_PROGRAM_UTILITIES_H_ */
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
@@ -124,60 +124,28 @@ size_t calculatedBufferSize(size_t headC
return (((int)ceil(headCount*approximationBits/8.0))/8)*8+16;
}
+static char* NOT_IMPLEMENTED = "case not implemented";
+static char* NOT_DECOMPOSED = "input bat not decomposed";
str BWDLeftJoinApproximate(bat * res, bat * l, bat * r){
- BAT* left;
- BAT* right = BATdescriptor(*r);
+ BAT *left = BATdescriptor(*l), *right = BATdescriptor(*r);
if(!batTailIsDecomposed(right)){
printf("bwd.%s : bat is not decomposed: %d, no approximation
can be provided\n", __func__, *r);
BBPreleaseref(right->batCacheid);
- return MAL_SUCCEED;
+ return NOT_DECOMPOSED;
}
- left = BATdescriptor(*l);
- if(0) printf ("%s left bat (%d) is of type [%s%s, %s%s]\n", __func__,
left->batCacheid, typeNames[BAThtype(left)], BAThvoid(left)?" (void)":"",
typeNames[BATttype(left)], BATtvoid(left)?" (void)":"");
-
- if(0 && BAThvoid(left) && BATtvoid(left) && BAThvoid(right) &&
right->tseqbase != oid_nil){
- // the more I think about this, the more I feel this case can
never happen
- uint i;
- int* outputRegion;
- const size_t offset = left->tseqbase;
- const size_t approximationBytes =
batTailApproximationBits(right)/8;
- const unsigned int residualBits =
32-batTailApproximationBits(right);
- unsigned char* approximation =
GDKmalloc(left->batCount*approximationBytes);
- BAT* result = BATnew(ATOMtype(left->htype),
ATOMtype(right->ttype), left->batCount);
-
- BATsetcount(result, left->batCount);
- bzero(Tloc(result, BUNfirst(result)),
left->batCount*sizeof(int));
- outputRegion = (int*) Tloc(result, BUNfirst(result));
-
- {
- clEnqueueReadBuffer(getCommandQueue(),
batTailApproximation(right), CL_TRUE, offset*approximationBytes,
left->batCount*approximationBytes, approximation, 0, NULL, NULL);
- }
-
-
- for (i = 0; i < left->batCount; ++i)
- outputRegion[i] = (*((unsigned int*)(approximation +
(i+offset)*approximationBytes))) << residualBits;
-
- GDKfree(approximation);
- BBPkeepref((*res = result->batCacheid));
- BBPreleaseref(left->batCacheid);
- BBPreleaseref(right->batCacheid);
- return MAL_SUCCEED;
- } else if(BAThvoid(left) && BAThvoid(right) && right->tseqbase !=
oid_nil){
+ if(BAThvoid(left) && BAThvoid(right) && right->tseqbase != oid_nil){
cl_int err;
int headCount; // :-)
cl_kernel projectKernel;
- cl_program program;
clTail newTailDefinition;
- cl_mem leftColumn = batTailApproximation(left);
- cl_mem rightColumn = batTailApproximation(right);
+ cl_mem leftColumn = batTailApproximation(left), rightColumn =
batTailApproximation(right);
BAT* result;
int i;
if(!leftColumn) leftColumn =
batHeadApproximation(BATmirror(left));
if(!rightColumn) rightColumn =
batHeadApproximation(BATmirror(right));
- program =
getProjectionLeftjoinProgram(batTailApproximationBits(right),
batTailOffsetBits(right));
- projectKernel = clCreateKernel(program, "project", &err);
+ projectKernel =
clCreateKernel(getProjectionLeftjoinProgram(batTailApproximationBits(right),
batTailOffsetBits(right)), "project", &err);
if(err) printf("#%s, clCreateKernel: %s;\n", __func__,
clError(err));
if((err = clEnqueueReadBuffer(getCommandQueue(), leftColumn,
CL_TRUE, 0, sizeof(int), &headCount , 0, NULL, NULL)))
@@ -198,11 +166,9 @@ str BWDLeftJoinApproximate(bat * res, ba
slot->tailOffsetValue = batTailOffsetValue(right);
slot->residuals = NULL;
slot->tailPositions = leftColumn;
- if(CL_REFCOUNT_DEBUG) printf ("retaining %p\n",
leftColumn);
if((err=clRetainMemObject(leftColumn)))
printf("#%s, clRetainMemObject: %s;\n",
__func__, clError(err));
slot->tailApproximation =
bwdClCreateBuffer(getCLContext(), CL_MEM_READ_WRITE,
calculatedBufferSize(headCount,slot->approximationBits)+sizeof(clTail), NULL,
&err);
- if(CL_REFCOUNT_DEBUG) printf ("%s, result (%d) tail
approximation: %p\n", __func__, result->batCacheid ,
batTailApproximation(result));
assert(slot->tailPositions);
if(err) printf("#%s, bwdClCreateBuffer: %s;\n",
__func__, clError(err));
@@ -220,21 +186,13 @@ str BWDLeftJoinApproximate(bat * res, ba
if((err = clEnqueueNDRangeKernel(getCommandQueue(),
projectKernel, 1, (const size_t[]){0}, (const
size_t[]){ceil(headCount/((float)WORK_GROUP_SIZE))*WORK_GROUP_SIZE}, (const
size_t[]){WORK_GROUP_SIZE}, 0, NULL, NULL)))
printf("#%s, clEnqueueNDRangeKernel: %s;\n", __func__,
clError(err));
if (synchronousGPU) clFinish(getCommandQueue());
-
BBPkeepref((*res = result->batCacheid));
- if(0) printf ("%s, returning bat %d\n", __func__, *res);
- if(0) printf ("%s, batTailApproximation of result: %p\n",
__func__, batTailApproximation(result));
- if(0) printf ("%s, batHeadApproximation of result: %p\n",
__func__, batHeadApproximation(result));
- if(0) printf ("%s, bbpreleasing left: %d\n", __func__,
left->batCacheid);
BBPreleaseref(left->batCacheid);
- if(0) printf ("%s, bbpreleasing right: %d\n", __func__,
right->batCacheid);
BBPreleaseref(right->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);
- 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;
};
static inline int decompressIntValue(const int approximationI, const int
approximationBits, const int offsetBits, const clTail* compressedTail, const
unsigned int* residuals, const unsigned int residualI){
@@ -1108,14 +1066,50 @@ str BWDMulticolumnGroupApproximate(Clien
str BWDGroupApproximate(int *rethisto, int *retbid, int *bid){
// this function should be used as a template for new BWD-operators
- GRPgroup(rethisto, retbid, bid);
- // bn = BATnew(b->htype, TYPE_oid, BATcount(b));
+ str result = MAL_SUCCEED;
+ BAT *b = BATdescriptor(*bid);
+ if (b == NULL)
+ throw(MAL, __func__, RUNTIME_OBJECT_MISSING);
+
+ if(!batTailIsDecomposed(b)){
+ printf("bwd.%s : bat is not decomposed: %d, no approximation
can be provided\n", __func__, *bid);
+ result = NOT_DECOMPOSED;
+ } else{ // get input objects and parameters
+ cl_mem valueColumn = batTailApproximation(b);
+ clTail newTailDefinition;
+ cl_int err;
- throw(MAL, __func__, "not implemented yet: %s", __func__);
+ if((err = clEnqueueReadBuffer(getCommandQueue(), valueColumn,
CL_TRUE, 0, sizeof(int), &newTailDefinition.count , 0, NULL, NULL)))
+ printf("#%s, clEnqueueReadBuffer (b): %s;\n", __func__,
clError(err));
+
+ { // create output objects
+ BAT* bn = BATnew(b->htype, TYPE_oid, BATcount(b));
+ BAT* histo = BATnew(TYPE_oid,TYPE_wrd, 0);
+ const unsigned int newIndex =
getNextFreeDecomposedBATSlotIndex();
+ DecomposedBATSlot* slot =
getDecomposedBATSlotForIndex(newIndex);
+ BATsetprop(bn, batRegistryIndex, TYPE_int,
(int[]){newIndex});
+
+ slot->tailApproximation =
bwdClCreateBuffer(getCLContext(), CL_MEM_READ_WRITE,
calculatedBufferSize(newTailDefinition.count,slot->approximationBits)+sizeof(clTail),
NULL, &err);
+
+ { // run kernel
+ cl_kernel groupKernel =
clCreateKernel(getGroupProgram(batTailApproximationBits(b)), "group", &err);
+ if((err =
clEnqueueNDRangeKernel(getCommandQueue(), groupKernel, 1, (const size_t[]){0},
(const
size_t[]){ceil(newTailDefinition.count/((float)WORK_GROUP_SIZE))*WORK_GROUP_SIZE},
(const size_t[]){WORK_GROUP_SIZE}, 0, NULL, NULL)))
+ printf("#%s, clEnqueueNDRangeKernel:
%s;\n", __func__, clError(err));
+
+ }
+ BBPkeepref(*rethisto = histo->batCacheid);
+ BBPkeepref(*retbid = bn->batCacheid);
+ }
+ }
+ BBPreleaseref(b->batCacheid);
+ return result;
}
str BWDGroupRefine(int *rethisto, int *retbid, int *bid, int *apprrethisto,
int *apprretbid){
- throw(MAL, __func__, "not implemented yet: %s", __func__);
+ *rethisto = *apprrethisto;
+ *retbid = *apprretbid;
+ return MAL_SUCCEED;
+ /* throw(MAL, __func__, "not implemented yet: %s", __func__); */
}
str BWDMulticolumnGroupRefine(Client cntxt, MalBlkPtr mb, MalStkPtr stk,
InstrPtr pci){
diff --git a/monetdb5/extras/bwd/utilities.c b/monetdb5/extras/bwd/utilities.c
--- a/monetdb5/extras/bwd/utilities.c
+++ b/monetdb5/extras/bwd/utilities.c
@@ -270,6 +270,7 @@ int batTailOffsetValue(const BAT* subjec
}
size_t batTailResidualBits(const BAT* subject){
+ assert(ATOMsize(subject->ttype) > 0);
return ATOMsize(subject->ttype)*8 - batTailApproximationBits(subject) -
batTailOffsetBits(subject);
}
const unsigned char* batTailResiduals(const BAT* subject){
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
