Changeset: b4228bb1e1e2 for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=b4228bb1e1e2
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/utilities.c
monetdb5/extras/bwd/utilities.h
Branch: bwd
Log Message:
* TPCH Q1 doesn't have any (obvious) gpu memory leaks any more
Unterschiede (gekürzt von 496 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
@@ -27,3 +27,6 @@ returns varchar(4096) external name bwd.
create function bwdevices()
returns table(platform int, device int, deviceDescription string,
platformDescription string) external name bwd.deviceInfo;
+
+create function bwdmemobjects()
+returns table(registryID int, batID int, approximationBits int, offsetBits
int, isPersistent int, tailPositionsRefcount int,
tailPositionsMemoryConsumption bigint, tailValuesRefcount int,
tailValuesMemoryConsumption bigint) external name bwd.memObjectInfo;
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
@@ -116,16 +116,16 @@ str bwdecompose(bat * res, bat * subject
if(batTailIsDecomposed(subject)) {
snprintf(buffer, 4096, "bat %d already decomposed",
subject->batCacheid);
} else if (Tsize(subject) == 4){
- BATsetprop(subject, batRegistryIndex, TYPE_int,
(int[]){decomposeIntArray((int*)Tloc(subject, BUNfirst(subject)),
subject->batCount, *approximationBits)});
+ BATsetprop(subject, batRegistryIndex, TYPE_int,
(int[]){decomposeIntArray((int*)Tloc(subject, BUNfirst(subject)),
subject->batCount, *approximationBits, subjectBAT)});
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)});
+ BATsetprop(subject, batRegistryIndex, TYPE_lng,
(lng[]){transferOIDArray((oid*)Tloc(subject, BUNfirst(subject)),
subject->batCount, *approximationBits, subjectBAT)});
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)});
+ BATsetprop(subject, batRegistryIndex, TYPE_lng,
(lng[]){decomposeLngArray((lng*)Tloc(subject, BUNfirst(subject)),
subject->batCount, *approximationBits, subjectBAT)});
snprintf(buffer, 4096, "successfully decomposed long integer
bat %d", subject->batCacheid);
} else if(Tsize(subject) == 1){
- BATsetprop(subject, batRegistryIndex, TYPE_int,
(int[]){decomposeVarchar1Array(Tloc(subject, BUNfirst(subject)),
subject->batCount, Tsize(subject), Tbase(subject), *approximationBits)});
+ BATsetprop(subject, batRegistryIndex, TYPE_int,
(int[]){decomposeVarchar1Array(Tloc(subject, BUNfirst(subject)),
subject->batCount, Tsize(subject), Tbase(subject), *approximationBits,
subjectBAT)});
snprintf(buffer, 4096, "successfully decomposed char
bat %d", subject->batCacheid);
} else {
snprintf(buffer, 4096, "bat %d cannot be decomposed because it
is not a 32 or 8-bit type", subject->batCacheid);
@@ -252,16 +252,29 @@ void releaseMemObject(cl_mem object, cha
}
void releaseCLMemObject(bat* subjectID){
- BAT* subject;
- if(0) printf ("releasing bat %d\n", *subjectID);
- 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);
+ BAT* subject = BBPquickdesc(*subjectID,0);
+ DecomposedBATSlot* slot = getDecomposedBATSlot(subject);
+ if(slot && !slot->isPersistentBAT){
+ if(0) printf ("releasing bat %d\n", *subjectID);
+ if(slot->tailApproximation){
+ releaseMemObject(slot->tailApproximation,
"batTailApproximation", *subjectID);
+ /* slot->tailApproximation = 0; */
+ }
+ if(slot->headApproximation != slot->tailPositions &&
slot->tailPositions){
+ releaseMemObject(slot->tailPositions,
"batTailPositions", *subjectID);
+ /* slot->tailPositions = 0; */
+ }
+ if(slot->headApproximation){
+ releaseMemObject(slot->headApproximation,
"batHeadApproximation", *subjectID);
+ /* slot->headApproximation = 0; */
+ }
+ if(slot->used){
+ /* getDecomposedBATSlot(subject)->used = 0; */
+ bzero(slot, sizeof(DecomposedBATSlot));
+ }
+
+
+
}
}
@@ -269,3 +282,96 @@ str bwdinit(void){
registerBBPDestroyCallback(releaseCLMemObject);
return MAL_SUCCEED;
}
+
+str memObjectInfo(bat* registryIDBID,
+ bat*
batIDBID,
+ bat*
approximationBitsBID,
+ bat*
offsetBitsBID,
+ bat*
isPersistentBID,
+ bat*
tailPositionsRefcountBID,
+ bat*
tailPositionsMemoryConsumptionBID,
+ bat*
tailValuesRefcountBID,
+ bat*
tailValuesMemoryConsumptionBID){
+
+ BAT* registryID = BATnew(TYPE_void, TYPE_int, MAX_DECOMPOSED_BATS);
+ BAT* batID = BATnew(TYPE_void, TYPE_int, MAX_DECOMPOSED_BATS);
+ BAT* approximationBits = BATnew(TYPE_void, TYPE_int,
MAX_DECOMPOSED_BATS);
+ BAT* offsetBits = BATnew(TYPE_void, TYPE_int, MAX_DECOMPOSED_BATS);
+ BAT* isPersistent = BATnew(TYPE_void, TYPE_int, MAX_DECOMPOSED_BATS);
+ BAT* tailPositionsRefcount = BATnew(TYPE_void, TYPE_int,
MAX_DECOMPOSED_BATS);
+ BAT* tailValuesRefcount = BATnew(TYPE_void, TYPE_int,
MAX_DECOMPOSED_BATS);
+
+ BAT* tailPositionsMemoryConsumption = BATnew(TYPE_void, TYPE_lng,
MAX_DECOMPOSED_BATS);
+ BAT* tailValuesMemoryConsumption = BATnew(TYPE_void, TYPE_lng,
MAX_DECOMPOSED_BATS);
+
+ BATseqbase(registryID, 0);
+ BATseqbase(batID, 0);
+ BATseqbase(approximationBits, 0);
+ BATseqbase(offsetBits, 0);
+ BATseqbase(isPersistent, 0);
+ BATseqbase(tailPositionsRefcount, 0);
+ BATseqbase(tailValuesRefcount, 0);
+ BATseqbase(tailPositionsMemoryConsumption, 0);
+ BATseqbase(tailValuesMemoryConsumption, 0);
+ {
+ int i;
+ for (i = 0; i<MAX_DECOMPOSED_BATS; ++i){
+ DecomposedBATSlot* slot =
getDecomposedBATSlotForIndex(i);
+ if(slot->used){
+ BUNappend(registryID, (int[]){i}, 0);
+ BUNappend(batID, (int[]){slot->used}, 0);
+ BUNappend(approximationBits,
(int[]){slot->approximationBits}, 0);
+ BUNappend(offsetBits,
(int[]){slot->tailOffsetBits}, 0);
+ BUNappend(isPersistent,
(int[]){slot->isPersistentBAT}, 0);
+ if(slot->tailPositions){
+ cl_uint refcount;
+ size_t size;
+ size_t memoryConsumption;
+ int err;
+ if((err =
clGetMemObjectInfo(slot->tailPositions, CL_MEM_REFERENCE_COUNT,
sizeof(cl_uint), &refcount, &size)))
+ printf ("error when getting
memobject %p info: %s\n", slot->tailPositions, clError(err));
+ BUNappend(tailPositionsRefcount,
(int[]){refcount}, 0);
+
+ if((err =
clGetMemObjectInfo(slot->tailApproximation?slot->tailApproximation:slot->tailApproximations[0],
CL_MEM_SIZE, sizeof(size_t), &memoryConsumption, &size)))
+ printf ("error when getting
memobject %p info: %s\n",
slot->tailApproximation?slot->tailApproximation:slot->tailApproximations[0],
clError(err));
+
BUNappend(tailPositionsMemoryConsumption, (long[]){memoryConsumption}, 0);
+
+ } else {
+ BUNappend(tailPositionsRefcount,
(int[]){-1}, 0);
+
BUNappend(tailPositionsMemoryConsumption, (long[]){0}, 0);
+ }
+ {
+ cl_uint refcount;
+ size_t size;
+ size_t memoryConsumption;
+ int err;
+
if(slot->tailApproximation?slot->tailApproximation:slot->tailApproximations[0]){
+ if((err =
clGetMemObjectInfo(slot->tailApproximation?slot->tailApproximation:slot->tailApproximations[0],
CL_MEM_REFERENCE_COUNT, sizeof(cl_uint), &refcount, &size)))
+ printf ("error when
getting memobject %p info: %s\n",
slot->tailApproximation?slot->tailApproximation:slot->tailApproximations[0],
clError(err));
+ BUNappend(tailValuesRefcount,
(int[]){refcount}, 0);
+
+ if((err =
clGetMemObjectInfo(slot->tailApproximation?slot->tailApproximation:slot->tailApproximations[0],
CL_MEM_SIZE, sizeof(size_t), &memoryConsumption, &size)))
+ printf ("error when
getting memobject %p info: %s\n",
slot->tailApproximation?slot->tailApproximation:slot->tailApproximations[0],
clError(err));
+
BUNappend(tailValuesMemoryConsumption, (long[]){memoryConsumption}, 0);
+
+ } else {
+ BUNappend(tailValuesRefcount,
(int[]){-1}, 0);
+
BUNappend(tailValuesMemoryConsumption, (long[]){0}, 0);
+ }
+
+ }
+ }
+ }
+
+ }
+ BBPkeepref((*registryIDBID = registryID->batCacheid));
+ BBPkeepref((*batIDBID = batID->batCacheid));
+ BBPkeepref((*approximationBitsBID = approximationBits->batCacheid));
+ BBPkeepref((*offsetBitsBID = offsetBits->batCacheid));
+ BBPkeepref((*isPersistentBID = isPersistent->batCacheid));
+ BBPkeepref((*tailPositionsRefcountBID =
tailPositionsRefcount->batCacheid));
+ BBPkeepref((*tailValuesRefcountBID = tailValuesRefcount->batCacheid));
+ BBPkeepref((*tailPositionsMemoryConsumptionBID =
tailPositionsMemoryConsumption->batCacheid));
+ BBPkeepref((*tailValuesMemoryConsumptionBID =
tailValuesMemoryConsumption->batCacheid));
+ return MAL_SUCCEED;
+}
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
@@ -172,7 +172,7 @@ 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);
- str result = NOT_IMPLEMENTED;
+ str resultCode = 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);
@@ -180,10 +180,10 @@ str BWDLeftJoinApproximate(bat * res, ba
}
if(BAThvoid(left) && BAThvoid(right) && right->T->type == TYPE_void){
BBPincref(*res = *l, 1);
- result = MAL_SUCCEED;
+ resultCode = MAL_SUCCEED;
} else if(BAThvoid(left) && left->T->type == TYPE_void &&
BAThvoid(right)){
BBPincref(*res = *r, 1);
- result = MAL_SUCCEED;
+ resultCode = MAL_SUCCEED;
} else if(BAThvoid(left) && BAThvoid(right) && right->tseqbase !=
oid_nil){
cl_int err;
int headCount; // :-)
@@ -208,7 +208,7 @@ str BWDLeftJoinApproximate(bat * res, ba
result = BATnew(TYPE_void, ATOMtype(right->ttype), 0);
{
- const unsigned int newIndex =
getNextFreeDecomposedBATSlotIndex();
+ const unsigned int newIndex =
getNextFreeDecomposedBATSlotIndex(result->batCacheid);
DecomposedBATSlot* slot =
getDecomposedBATSlotForIndex(newIndex);
BATsetprop(result, batRegistryIndex, TYPE_int,
(int[]){newIndex});
slot->approximationBits =
batTailApproximationBits(right);
@@ -237,11 +237,11 @@ 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;
+ resultCode = 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);
}
- {
+ if(0){
BAT* approximation = BATdescriptor(*res);
if(batTailApproximation(approximation))
assert(batTailPositions(approximation));
@@ -251,7 +251,7 @@ str BWDLeftJoinApproximate(bat * res, ba
BBPreleaseref(right->batCacheid);
- return result;
+ return resultCode;
};
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){
@@ -551,7 +551,7 @@ static inline str uselect(bat *res, bat
int err;
cl_kernel selectKernel;
{
- const unsigned int newIndex =
getNextFreeDecomposedBATSlotIndex();
+ const unsigned int newIndex =
getNextFreeDecomposedBATSlotIndex(result->batCacheid);
DecomposedBATSlot* slot =
getDecomposedBATSlotForIndex(newIndex);
cl_int err;
const int size =
MIN(dataCount,MAX_INTERMEDIATE_RESULT_SIZE)*sizeof(int)+sizeof(clHead);
@@ -884,7 +884,7 @@ str BWDtmark(int *res, int *bid, oid *ba
BAT* input = BATdescriptor(*bid);
if(*res && input->H->type != TYPE_void){ // propagate head approximation
cl_int err;
- const unsigned int newIndex =
getNextFreeDecomposedBATSlotIndex();
+ const unsigned int newIndex =
getNextFreeDecomposedBATSlotIndex(*res);
DecomposedBATSlot* slot =
getDecomposedBATSlotForIndex(newIndex);
BAT* result = BATdescriptor(*res);
BATsetprop(result, batRegistryIndex, TYPE_int,
(int[]){newIndex});
@@ -907,7 +907,7 @@ str BWDmirror(int *res, int *bid){
BAT* input = BATdescriptor(*bid);
if(*res && input->H->type != TYPE_void){ // propagate head approximation
cl_int err;
- const unsigned int newIndex =
getNextFreeDecomposedBATSlotIndex();
+ const unsigned int newIndex =
getNextFreeDecomposedBATSlotIndex(*res);
DecomposedBATSlot* slot =
getDecomposedBATSlotForIndex(newIndex);
BAT* result = BATdescriptor(*res);
BATsetprop(result, batRegistryIndex, TYPE_int,
(int[]){newIndex});
@@ -985,7 +985,7 @@ str BWDSemijoinApproximate(int *res, int
{
int* zeroIntPattern = GDKzalloc(4);
- const unsigned int newIndex =
getNextFreeDecomposedBATSlotIndex();
+ const unsigned int newIndex =
getNextFreeDecomposedBATSlotIndex(result->batCacheid);
DecomposedBATSlot* slot =
getDecomposedBATSlotForIndex(newIndex);
BATsetprop(result, batRegistryIndex, TYPE_int,
(int[]){newIndex});
slot->approximationBits =
batTailApproximationBits(left);
@@ -1210,7 +1210,7 @@ str BWDSemijoinRefine(int *res, int *lid
if(0) printf ("%s setting residuals of %d to %p\n",
__func__, refinement->batCacheid, getDecomposedBATSlot(refinement)->residuals);
- BATsetprop(refinement, batRegistryIndex, TYPE_int,
(int[]){getNextFreeDecomposedBATSlotIndex()});
+ BATsetprop(refinement, batRegistryIndex, TYPE_int,
(int[]){getNextFreeDecomposedBATSlotIndex(refinement->batCacheid)});
getDecomposedBATSlot(refinement)->residuals = (unsigned
char*)filteredResiduals;
getDecomposedBATSlot(refinement)->approximationBits =
batTailApproximationBits(left);
getDecomposedBATSlot(refinement)->tailOffsetBits =
batTailOffsetBits(left);
@@ -1331,34 +1331,34 @@ str BWDSelectNotNilRefine(int *ret, int
str BWDGroupedSumApproximate(int *retid,int *valuesID, int *groupsID, int
*histogramID){
str result = MAL_SUCCEED;
- BAT *groups = BATdescriptor(*groupsID);
BAT *values = BATdescriptor(*valuesID);
- BAT *histogram = BATdescriptor(*histogramID);
if(values == NULL){
- printf("bwd.%s : approximate value bat is not available: %d, no
approximation can be provided\n", __func__, groups->batCacheid);
- return MAL_SUCCEED;
+ printf("bwd.%s : approximate value bat is not available: %d, no
approximation can be provided\n", __func__, *valuesID);
+ result = MAL_SUCCEED;
+ } else {
+ BAT *groups = BATdescriptor(*groupsID);
+ BAT *histogram = BATdescriptor(*histogramID);
+ if (groups == NULL || histogram == NULL)
+ throw(MAL, __func__, RUNTIME_OBJECT_MISSING);
+ if(!batTailIsDecomposed(groups)){
+ printf("bwd.%s : groups bat is not decomposed: %d, no
approximation can be provided\n", __func__, groups->batCacheid);
+ result = NOT_DECOMPOSED;
+ } else if(!batTailIsDecomposed(values)){
+ printf("bwd.%s : values bat is not decomposed: %d, no
approximation can be provided\n", __func__, values->batCacheid);
+ result = NOT_DECOMPOSED;
+ } else if(!batTailIsDecomposed(histogram)){
+ printf("bwd.%s : histogram bat is not decomposed: %d,
no approximation can be provided\n", __func__, histogram->batCacheid);
+ result = NOT_DECOMPOSED;
+ } else {
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
