Changeset: d7365bc58da0 for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=d7365bc58da0
Modified Files:
monetdb5/extras/bwd/91_opt_bwd.mal
monetdb5/extras/bwd/bwd.c
monetdb5/extras/bwd/bwd.h
monetdb5/extras/bwd/cl_program_utilities.c
monetdb5/extras/bwd/operations.c
monetdb5/extras/bwd/opt_bwd.mal
monetdb5/extras/bwd/utilities.c
monetdb5/extras/bwd/utilities.h
Branch: bwd
Log Message:
* implemented gpu garbage collection (seems to work reasonably well)
Unterschiede (287 Zeilen):
diff --git a/monetdb5/extras/bwd/91_opt_bwd.mal
b/monetdb5/extras/bwd/91_opt_bwd.mal
--- a/monetdb5/extras/bwd/91_opt_bwd.mal
+++ b/monetdb5/extras/bwd/91_opt_bwd.mal
@@ -1,3 +1,5 @@
library bwd;
-include opt_bwd;
\ No newline at end of file
+include opt_bwd;
+
+bwd.init();
\ No newline at end of file
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
@@ -161,3 +161,35 @@ bwd_result_column_wrap(Client cntxt, Mal
{
return MAL_SUCCEED;
}
+
+void releaseMemObject(cl_mem object, char* description, bat subjectID){
+ cl_int err;
+ if(0){
+ cl_uint refcount;
+ size_t size;
+ if((err = clGetMemObjectInfo(object, CL_MEM_REFERENCE_COUNT,
sizeof(cl_uint), &refcount, &size)))
+ printf ("error when getting memobject %p info: %s\n",
object, clError(err));
+ printf ("decreasing refcount of %s(%d) %p (%d) by 1\n",
description?description:"", subjectID, object, refcount);
+ }
+
+ if((err = clReleaseMemObject(object)) != CL_SUCCESS)
+ printf ("error when relasing memobject %p: %s\n", object,
clError(err));
+}
+
+void releaseCLMemObject(bat* subjectID){
+ if(0) printf ("releasing bat %d\n", *subjectID);
+ BAT* subject = BBPquickdesc(*subjectID,0);
+ if(!isPersistentBAT(subject)){
+ if(batHeadApproximation(subject))
+ releaseMemObject(batHeadApproximation(subject),
"batHeadApproximation", *subjectID);
+ if(batTailApproximation(subject))
+ releaseMemObject(batTailApproximation(subject),
"batTailApproximation", *subjectID);
+ if(batHeadApproximation(subject) != batTailPositions(subject)
&& batTailPositions(subject))
+ releaseMemObject(batTailPositions(subject),
"batTailPositions", *subjectID);
+ }
+}
+
+str bwdinit(){
+ registerBBPDestroyCallback(releaseCLMemObject);
+ return MAL_SUCCEED;
+}
diff --git a/monetdb5/extras/bwd/bwd.h b/monetdb5/extras/bwd/bwd.h
--- a/monetdb5/extras/bwd/bwd.h
+++ b/monetdb5/extras/bwd/bwd.h
@@ -9,5 +9,7 @@ str deviceInfo(bat * resPlatform, bat *
str bwd_result_column_wrap(Client cntxt, MalBlkPtr mb, MalStkPtr stk, InstrPtr
pci);
char* resolveBatToAttribute(bat bid, Client cntxt);
+str bwdinit();
+
#endif /* _BWD_H_ */
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
@@ -108,7 +108,7 @@ cl_program getUSelectProgram(int type, c
"{\n"
" const int index = atomic_inc(&(outputHead->count));\n"
" atomic_inc(&(outputTail->count));\n" // TODO: this could
probably be done more efficiently
- /* " printf(\"selected value %%d from slot %%d into slot
%%d, head value: %d\\n\", value, get_global_id(0), index,
approximationHead->positions[get_global_id(0)]);\n" */
+ " printf(\"selected value %%d from slot %%d into slot %%d,
head value_: %%d\\n\", value, get_global_id(0), index,
approximationHead->positions[get_global_id(0)]);\n"
" const int offset = index * %4$d;\n"
" outputHead->positions[index] =
approximationHead->positions[get_global_id(0)];\n"
" for(int i = 0; i < %4$d; i++){\n"
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
@@ -27,6 +27,7 @@ static const int WORK_GROUP_SIZE=16;
#endif
#define MAX_INTERMEDIATE_RESULT_SIZE 16777216
+static const int CL_REFCOUNT_DEBUG = 0;
#pragma mark Actual MAL Operations Implementation
#ifdef __APPLE__
@@ -155,7 +156,11 @@ str BWDLeftJoinApproximate(bat * res, ba
slot->tailOffsetBits = batTailOffsetBits(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 =
clCreateBuffer(getCLContext(), CL_MEM_READ_WRITE,
headCount*slot->approximationBits/8+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, clCreateBuffer: %s;\n", __func__,
clError(err));
@@ -175,7 +180,12 @@ str BWDLeftJoinApproximate(bat * res, ba
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;
@@ -294,7 +304,7 @@ static inline str uselect(bat *res, bat
const clTail dataHeader = getTailHeader(batTailApproximation(data));
const int dataCount = dataHeader.count;
-
+ printf ("selecting out of %d values\n", dataCount);
if(dataCount){
{
const unsigned int newIndex =
getNextFreeDecomposedBATSlotIndex();
@@ -531,9 +541,9 @@ str uselectrefine(bat *res, bat *bid, pt
free(compressedHead);
free(compressedTail);
}
- BBPreleaseref(approximation->batCacheid);
}
+ BBPreleaseref(approximation->batCacheid);
}
if(0) printf ("%s result bat (%d) is of type [%s%s, %s%s]\n", __func__,
result->batCacheid, typeNames[BAThtype(result)], BAThvoid(result)?" (void)":"",
typeNames[BATttype(result)], BATtvoid(result)?" (void)":"");
@@ -557,6 +567,11 @@ str BWDtmark(int *res, int *bid, oid *ba
BBPreleaseref(*res);
BAT* input = BATdescriptor(*bid);
slot->headApproximation = batHeadApproximation(input);
+ if(0) printf ("retaining %p\n", slot->headApproximation);
+ cl_int err;
+ if((err=clRetainMemObject(slot->headApproximation)))
+ printf("#%s, clRetainMemObject: %s;\n", __func__,
clError(err));
+
BBPreleaseref(*bid);
}
return returnValue;
@@ -612,6 +627,10 @@ str BWDSemijoinApproximate(int *res, int
slot->residuals = NULL;
slot->tailPositions = positionColumn;
slot->headApproximation = positionColumn;
+ if(CL_REFCOUNT_DEBUG) printf ("retaining %p\n",
positionColumn);
+ if((err=clRetainMemObject(positionColumn)))
+ printf("#%s, clRetainMemObject: %s;\n",
__func__, clError(err));
+
slot->tailApproximation =
clCreateBuffer(getCLContext(), CL_MEM_READ_WRITE,
headCount*slot->approximationBits/8+sizeof(clTail), NULL, &err);
assert(slot->tailPositions);
@@ -641,13 +660,27 @@ str BWDSemijoinApproximate(int *res, int
if(0) printf ("first position: %d\n",
supersetPositionsColumn->positions[0]);
}
BBPkeepref((*res = result->batCacheid));
- BBPreleaseref(left->batCacheid);
- BBPreleaseref(right->batCacheid);
- return MAL_SUCCEED;
+ }
+ if(CL_REFCOUNT_DEBUG) {
+ int refCount = BBPrefs(left->batCacheid);
+ printf ("left(%d) hast %d references\n",left->batCacheid,
refCount);
+ }
+ BBPreleaseref(left->batCacheid);
+ if(CL_REFCOUNT_DEBUG) {
+ int refCount = BBPrefs(left->batCacheid);
+ printf ("left hast %d references\n", refCount);
+ }
+ if(CL_REFCOUNT_DEBUG) {
+ int refCount = BBPrefs(right->batCacheid);
+ printf ("right(%d) has %d references\n", right->batCacheid,
refCount);
+ }
-
- }
+ BBPreleaseref(right->batCacheid);
+ if(CL_REFCOUNT_DEBUG) {
+ int refCount = BBPrefs(right->batCacheid);
+ printf ("right has %d references after\n", refCount);
+ }
return MAL_SUCCEED;
}
diff --git a/monetdb5/extras/bwd/opt_bwd.mal b/monetdb5/extras/bwd/opt_bwd.mal
--- a/monetdb5/extras/bwd/opt_bwd.mal
+++ b/monetdb5/extras/bwd/opt_bwd.mal
@@ -96,6 +96,9 @@ address bwdecomposeScalar;
command decompose(column:date, approximationBits: int):void
address bwdecomposeScalar;
+command init():void
+address bwdinit;
+
module batbwd;
command decompose(col:bat[:oid,:int],bits:int):bat[:oid,:str]
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
@@ -139,6 +139,7 @@ static DecomposedBATSlot bwdRegistry[MAX
const char* batRegistryIndex = "bwd.batRegistryIndex";
const char batTailIsDecomposed(const BAT* subject){
+ if(0) printf ("bat %d is %s decomposed\n", subject->batCacheid,
BATgetprop(subject, batRegistryIndex) == NULL?"not":"");
return BATgetprop(subject, batRegistryIndex) != NULL;
}
@@ -188,6 +189,15 @@ const size_t batTailOffsetBits(const BAT
return
bwdRegistry[rightTailApproximationProperty->v.val.ival].tailOffsetBits;
}
+const int isPersistentBAT(const BAT* subject){
+ PROPrec* rightTailApproximationProperty;
+ if(!(rightTailApproximationProperty = BATgetprop(subject,
batRegistryIndex))){
+ if(0) printf("%s: bat hasn't been decomposed;\n",
__func__);
+ return 0;
+ }
+ return
bwdRegistry[rightTailApproximationProperty->v.val.ival].isPersistentBAT;
+}
+
const size_t batTailResidualBits(const BAT* subject){
return ATOMsize(subject->ttype)*8 - batTailApproximationBits(subject);
}
@@ -203,8 +213,9 @@ const unsigned char* batTailResiduals(co
const unsigned int getNextFreeDecomposedBATSlotIndex(){
size_t i = 0;
- while (i < MAX_DECOMPOSED_BATS && bwdRegistry[i].approximationBits > 0)
+ while (i < MAX_DECOMPOSED_BATS && bwdRegistry[i].used > 0)
i++;
+ bwdRegistry[i].used = 1;
return i;
};
@@ -263,6 +274,7 @@ const unsigned int decomposeIntArray(con
const int neededBits = ceil(log2(subjectBounds.max -
subjectBounds.min));
slot->tailOffsetBits = 32-ceil(neededBits/8.0)*8;
slot->approximationBits = approximationBits;
+ slot->isPersistentBAT = 1;
printf ("bounds: %d, %d, offset %zd bytes\n",subjectBounds.min,
subjectBounds.max, slot->tailOffsetBits/8);
@@ -281,6 +293,7 @@ const unsigned int decomposeIntArray(con
for (i = 0; i < size; ++i){
*((unsigned int*)(slot->residuals + i*residualBytes)) |=
((subject[i]-approximation->base)&residualMask);
*((unsigned
int*)(&approximation->elements[i*approximationBytes])) |=
((subject[i]-approximation->base) >> residualBits);
+ /* printf ("value: %d shifted approximation value: %d\n",
(subject[i]-approximation->base), ((subject[i]-approximation->base) >>
residualBits)); */
}
{
cl_int err;
@@ -297,10 +310,10 @@ const unsigned int decomposeIntArray(con
if(0){
int i, j;
- for (i = 0; i < size; ++i) {
+ for (i = 0; i < MIN(size,10); ++i) {
int approximationValue = 0, residualValue = 0;
for (j = 0; j < approximationBytes ; ++j){
- printf ("%d\n",
approximation->elements[i*approximationBytes+j]);
+ printf ("aproximationbyte %d: %d, ", j,
approximation->elements[i*approximationBytes+j]);
approximationValue =
approximation->elements[i*approximationBytes+j]<<(j*8);
}
diff --git a/monetdb5/extras/bwd/utilities.h b/monetdb5/extras/bwd/utilities.h
--- a/monetdb5/extras/bwd/utilities.h
+++ b/monetdb5/extras/bwd/utilities.h
@@ -7,6 +7,8 @@ typedef struct {
cl_mem headApproximation;
cl_mem tailApproximation;
cl_mem tailPositions;
+ int isPersistentBAT;
+ char used;
size_t approximationBits;
size_t tailOffsetBits;
unsigned char* residuals;
@@ -40,6 +42,8 @@ const size_t batTailApproximationBits(co
const size_t batTailOffsetBits(const BAT* subject) __attribute__((pure));
const size_t batTailResidualBits(const BAT* subject) __attribute__((pure));
const unsigned char* batTailResiduals(const BAT*) __attribute__((pure));
+const int isPersistentBAT(const BAT* subject) __attribute__((pure));
+
const unsigned int decomposeIntArray(const int* subject, const size_t size,
const size_t approximationBits);
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
