Changeset: ed790da2de07 for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=ed790da2de07
Modified Files:
monetdb5/extras/bwd/operations.c
Branch: bwd
Log Message:
* fixed a bug in the multicolumn grouping that switched histogram and group
assignment bats
* more verbose error logging (including source file and line)
* working on decompression of accurate char-values
Unterschiede (gekürzt von 310 auf 300 Zeilen):
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
@@ -51,7 +51,7 @@ int getCount(cl_mem memoryObject){
err = clEnqueueReadBuffer(getCommandQueue(), memoryObject, CL_TRUE, 0,
sizeof(int), &result, 0, NULL, NULL);
if(err)
- printf("#%s, clEnqueueReadBuffer: %s;\n", __func__,
clError(err));
+ printf("#%s, clEnqueueReadBuffer: %s, %s:%u;\n", __func__,
clError(err), __FILE__, __LINE__);
return result;
}
@@ -61,7 +61,7 @@ clTail getTailHeader(cl_mem memoryObject
err = clEnqueueReadBuffer(getCommandQueue(), memoryObject, CL_TRUE, 0,
sizeof(clHead), &result, 0, NULL, NULL);
if(err)
- printf("#%s, clEnqueueReadBuffer: %s;\n", __func__,
clError(err));
+ printf("#%s, clEnqueueReadBuffer: %s, %s:%u;\n", __func__,
clError(err), __FILE__, __LINE__);
return result;
}
@@ -83,7 +83,7 @@ clHead* getPositionsColumn(cl_mem memory
err = clEnqueueReadBuffer(getCommandQueue(), memoryObject,
CL_TRUE, 0, *bufferSize, buffer, 0, NULL, NULL); }
if(err){
- printf("#%s, clEnqueueReadBuffer: %s;\n", __func__,
clError(err));
+ printf("#%s, clEnqueueReadBuffer: %s, %s:%u;\n", __func__,
clError(err), __FILE__, __LINE__);
}
return buffer;
}
@@ -94,7 +94,7 @@ clTail* getApproximateValuesColumn(cl_me
clGetMemObjectInfo(memoryObject, CL_MEM_SIZE, sizeof(size_t),
bufferSize, NULL);
else
err = clEnqueueReadBuffer(getCommandQueue(), memoryObject,
CL_TRUE, 0, *bufferSize, buffer, 0, NULL, NULL);
- if(err) printf("#%s, clEnqueueReadBuffer: %s;\n", __func__,
clError(err));
+ if(err) printf("#%s, clEnqueueReadBuffer: %s, %s:%u;\n", __func__,
clError(err), __FILE__, __LINE__);
return buffer;
}
@@ -120,6 +120,27 @@ cl_int clEnqueueFillBuffer(cl_command_qu
};
#endif
+static inline cl_int bwdEnqueueNDRangeKernel ( cl_command_queue command_queue,
+
cl_kernel kernel,
+
cl_uint work_dim,
+
const size_t
*global_work_offset,
+
const size_t
*global_work_size,
+
const size_t
*local_work_size,
+
cl_uint
num_events_in_wait_list,
+
const cl_event
*event_wait_list,
+
cl_event *event){
+ return clEnqueueNDRangeKernel ( command_queue,
+
kernel,
+
work_dim,
+
global_work_offset,
+
global_work_size,
+
local_work_size,
+
num_events_in_wait_list,
+
event_wait_list,
+
event);
+
+ }
+
const unsigned int zeroIntPattern[1] = {};
size_t calculatedBufferSize(size_t headCount, size_t approximationBits){
@@ -169,11 +190,13 @@ str BWDLeftJoinApproximate(bat * res, ba
slot->tailPositions = leftColumn;
if((err=clRetainMemObject(leftColumn)))
printf("#%s, clRetainMemObject: %s;\n",
__func__, clError(err));
- const size_t tailSize =
calculatedBufferSize(headCount,slot->approximationBits)+sizeof(clTail);
- slot->tailApproximation =
bwdClCreateBuffer(getCLContext(), CL_MEM_READ_WRITE, tailSize, NULL, &err);
+ slot->tailApproximation =
bwdClCreateBuffer(getCLContext(), CL_MEM_READ_WRITE,
calculatedBufferSize(headCount,slot->approximationBits)+sizeof(clTail), NULL,
&err);
if(err) printf("#%s, bwdClCreateBuffer: %s;\n",
__func__, clError(err));
- err = clEnqueueFillBuffer(getCommandQueue(),
slot->tailApproximation, zeroIntPattern, sizeof(int), 0, tailSize, 0, NULL,
NULL);
+ err = clEnqueueWriteBuffer(getCommandQueue(),
slot->tailApproximation, CL_TRUE, 0, sizeof(clTail), &newTailDefinition, 0,
NULL, NULL);
+ if(err) printf("#%s, clEnqueueWriteBuffer: %s;\n",
__func__, clError(err));
+
+ err = clEnqueueFillBuffer(getCommandQueue(),
slot->tailApproximation, zeroIntPattern, sizeof(int), sizeof(clTail),
calculatedBufferSize(headCount,slot->approximationBits), 0, NULL, NULL);
if(err) printf("#%s, clEnqueueFillBuffer: %s;\n",
__func__, clError(err));
}
for (i = 0; i < 3; ++i) {
@@ -181,8 +204,8 @@ str BWDLeftJoinApproximate(bat * res, ba
batTailApproximation(result), leftColumn, rightColumn}[i]
)))) printf("#%s,
clSetKernelArg(%d): %s;\n", __func__, 0, clError(err));
}
- 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((err = bwdEnqueueNDRangeKernel(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, %s:%u;\n",
__func__, clError(err), __FILE__, __LINE__);
if (synchronousGPU) clFinish(getCommandQueue());
BBPkeepref((*res = result->batCacheid));
BBPreleaseref(left->batCacheid);
@@ -230,6 +253,39 @@ static inline int decompress_int_Value(c
return deCompressedValue;
}
+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;
+ 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 char decompress_char_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 char decompress_char_Value(const int approximationI, const int
approximationBits, const int offsetBits, const clTail* compressedTail, const
unsigned int* residuals, const unsigned int residualI){
const unsigned int residualBits = 8-approximationBits-offsetBits;
const unsigned int residualMask = (1 << residualBits)-1;
@@ -315,7 +371,29 @@ str BWDLeftJoinRefine(bat * res, bat * l
supersetApproximateValuesColumn =
getApproximateValuesColumn(batTailApproximation(approximation),
GDKmalloc(bufferSize), &bufferSize);
refinement = BATnew(TYPE_void, ATOMtype(right->ttype),
left->batCount);
- {
+ if(batTailResidualBits(right) == 0){
+#define expansionLoop(type) {
\
+ 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 4:
+ expansionLoop(int);
+ break;
+ case 1:
+ expansionLoop(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 expansionLoop
+ } else {
oid* positionRegion = (oid*) Tloc(left,
BUNfirst(left));
size_t refinementCount = 0;
const unsigned int* residuals = (unsigned
int*)batTailResiduals(right);
@@ -347,7 +425,7 @@ str BWDLeftJoinRefine(bat * res, bat * l
#undef refinementLoop
BATsetcount(refinement, refinementCount);
BATseqbase(refinement, 0);
- }
+ }
GDKfree(supersetPositionsColumn);
GDKfree(supersetApproximateValuesColumn);
@@ -447,9 +525,9 @@ static inline str uselect(bat *res, bat
if((err =
clSetKernelArg(selectKernel, bufferI+i, sizeof(int), &(parameters[i])))) //
type specific
printf("#%s,
clSetKernelArg(%d): %s;\n", __func__, bufferI+i, clError(err));
- err =
clEnqueueNDRangeKernel(getCommandQueue(), selectKernel, 1, (const size_t[]){0},
(const size_t[]){ceil(dataCount/((float)WORK_GROUP_SIZE))*(WORK_GROUP_SIZE)},
(const size_t[]){WORK_GROUP_SIZE}, 0, NULL, NULL);
+ err =
bwdEnqueueNDRangeKernel(getCommandQueue(), selectKernel, 1, (const
size_t[]){0}, (const
size_t[]){ceil(dataCount/((float)WORK_GROUP_SIZE))*(WORK_GROUP_SIZE)}, (const
size_t[]){WORK_GROUP_SIZE}, 0, NULL, NULL);
if(err)
- printf("#%s,
clEnqueueNDRangeKernel: %s;\n", __func__, clError(err));
+ printf("#%s,
clEnqueueNDRangeKernel: %s, %s:%u;\n", __func__, clError(err), __FILE__,
__LINE__);
if (synchronousGPU)
clFinish(getCommandQueue());
if(1){
clTail* compressedTail;
@@ -460,7 +538,7 @@ static inline str uselect(bat *res, bat
, NULL);
compressedTail =
GDKmalloc(resultSize);
err =
clEnqueueReadBuffer(getCommandQueue(), batTailApproximation(result), CL_TRUE,
0, resultSize, compressedTail , 0, NULL, NULL);
- if(err) printf("#%s,
clEnqueueReadBuffer: %s;\n", __func__, clError(err));
+ if(err) printf("#%s,
clEnqueueReadBuffer: %s, %s:%u;\n", __func__, clError(err), __FILE__, __LINE__);
(void) compressedTail;
}
}
@@ -636,7 +714,7 @@ str uselectrefine(bat *res, bat *bid, pt
clGetMemObjectInfo(headApproximation,
CL_MEM_SIZE, sizeof(size_t), &approximationSize, NULL);
compressedHead = GDKmalloc(approximationSize);
err = clEnqueueReadBuffer(getCommandQueue(),
batHeadApproximation(approximation), CL_TRUE, 0, approximationSize,
compressedHead , 0, NULL, NULL);
- if(err) printf("#%s, clEnqueueReadBuffer:
%s;\n", __func__, clError(err));
+ if(err) printf("#%s, clEnqueueReadBuffer: %s,
%s:%u;\n", __func__, clError(err), __FILE__, __LINE__);
}
{
const size_t candidateCount =
compressedHead->count;
@@ -649,6 +727,7 @@ str uselectrefine(bat *res, bat *bid, pt
unsigned int i;
for (i = 0; i < candidateCount;
++i)
positionRegion[i] =
headPositions[i];
+ if(1) printf ("%s not reduced
the result from %zd to %d because all bits are gpu resident\n", __func__,
candidateCount, i);
}
BATsetcount(result, candidateCount);
BATseqbase(BATmirror(result), oid_nil);
@@ -670,7 +749,7 @@ str uselectrefine(bat *res, bat *bid, pt
, NULL);
compressedTail =
GDKmalloc(approximationSize);
err =
clEnqueueReadBuffer(getCommandQueue(), batTailApproximation(approximation),
CL_TRUE, 0, approximationSize, compressedTail , 0, NULL, NULL);
- if(err) printf("#%s,
clEnqueueReadBuffer: %s;\n", __func__, clError(err));
+ if(err) printf("#%s,
clEnqueueReadBuffer: %s, %s:%u;\n", __func__, clError(err), __FILE__, __LINE__);
}
{
@@ -794,11 +873,11 @@ str BWDSemijoinApproximate(int *res, int
if(err) printf("#%s, clCreateKernel: %s;\n", __func__,
clError(err));
if((err = clEnqueueReadBuffer(getCommandQueue(),
positionColumn, CL_TRUE, 0, sizeof(int), &headCount , 0, NULL, NULL)))
- printf("#%s, clEnqueueReadBuffer: %s;\n",
__func__, clError(err));
+ printf("#%s, clEnqueueReadBuffer: %s,
%s:%u;\n", __func__, clError(err), __FILE__, __LINE__);
newTailDefinition.count = headCount; // :-)
if((err = clEnqueueReadBuffer(getCommandQueue(),
rightColumn, CL_TRUE, sizeof(int), sizeof(int), &newTailDefinition.base , 0,
NULL, NULL)))
- printf("#%s, clEnqueueReadBuffer: %s;\n",
__func__, clError(err));
+ printf("#%s, clEnqueueReadBuffer: %s,
%s:%u;\n", __func__, clError(err), __FILE__, __LINE__);
{
@@ -833,8 +912,8 @@ str BWDSemijoinApproximate(int *res, int
batTailApproximation(result), positionColumn, rightColumn}[i]
))))
printf("#%s, clSetKernelArg(%d): %s;\n", __func__, 0, clError(err));
if(headCount)
- 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((err =
bwdEnqueueNDRangeKernel(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, %s:%u;\n", __func__, clError(err), __FILE__,
__LINE__);
if(0) printf ("returning approximation with
tail positions %p\n", batTailPositions(result));
if (synchronousGPU) clFinish(getCommandQueue());
@@ -1248,8 +1327,8 @@ str BWDGroupApproximate(int *rethisto, i
batTailApproximation(b)}[i]
))))
printf("#%s, clSetKernelArg(%d): %s;\n", __func__, i, clError(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));
+ if((err =
bwdEnqueueNDRangeKernel(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, %s:%u;\n", __func__, clError(err), __FILE__, __LINE__);
if (synchronousGPU) clFinish(getCommandQueue());
}
@@ -1265,8 +1344,8 @@ str BWDGroupApproximate(int *rethisto, i
newHead, newTail}[i]
))))
printf("#%s, clSetKernelArg(%d): %s;\n", __func__, i, clError(err));
}
- if((err =
clEnqueueNDRangeKernel(getCommandQueue(), compactKernel, 1, (const
size_t[]){0}, (const size_t[]){histogramSize}, (const size_t[]){2}, 0, NULL,
NULL)))
- printf("#%s, clEnqueueNDRangeKernel:
%s;\n", __func__, clError(err));
+ if((err =
bwdEnqueueNDRangeKernel(getCommandQueue(), compactKernel, 1, (const
size_t[]){0}, (const size_t[]){histogramSize}, (const size_t[]){2}, 0, NULL,
NULL)))
+ printf("#%s, clEnqueueNDRangeKernel:
%s, %s:%u;\n", __func__, clError(err), __FILE__, __LINE__);
if (synchronousGPU) clFinish(getCommandQueue());
clReleaseMemObject(histogramSlot->tailApproximation);
histogramSlot->tailApproximation = newTail;
@@ -1284,8 +1363,8 @@ str BWDGroupApproximate(int *rethisto, i
}
str BWDMulticolumnGroupApproximate(Client cntxt, MalBlkPtr mb, MalStkPtr stk,
InstrPtr pci){
- int *retbid = (int*) getArgReference(stk,pci,0);
- int *rethisto = (int*) getArgReference(stk,pci,1);
+ int *rethisto = (int*) getArgReference(stk,pci,0);
+ int *retbid = (int*) getArgReference(stk,pci,1);
str result = MAL_SUCCEED;
int i;
BAT **inputs = (BAT**) GDKzalloc(sizeof(BAT*) * pci->argc);
@@ -1368,8 +1447,8 @@ str BWDMulticolumnGroupApproximate(Clien
}
- 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));
+ if((err =
bwdEnqueueNDRangeKernel(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, %s:%u;\n", __func__, clError(err), __FILE__, __LINE__);
if (synchronousGPU) clFinish(getCommandQueue());
}
@@ -1385,14 +1464,16 @@ str BWDMulticolumnGroupApproximate(Clien
newHead, newTail}[i]
))))
printf("#%s, clSetKernelArg(%d): %s;\n", __func__, i, clError(err));
}
- if((err =
clEnqueueNDRangeKernel(getCommandQueue(), compactKernel, 1, (const
size_t[]){0}, (const size_t[]){histogramSize}, (const size_t[]){2}, 0, NULL,
NULL)))
- printf("#%s, clEnqueueNDRangeKernel:
%s;\n", __func__, clError(err));
+ if((err =
bwdEnqueueNDRangeKernel(getCommandQueue(), compactKernel, 1, (const
size_t[]){0}, (const size_t[]){histogramSize}, (const size_t[]){2}, 0, NULL,
NULL)))
+ printf("#%s, clEnqueueNDRangeKernel:
%s, %s:%u;\n", __func__, clError(err), __FILE__, __LINE__);
if (synchronousGPU) clFinish(getCommandQueue());
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
