Changeset: 77160b3e8cfb for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=77160b3e8cfb
Modified Files:
monetdb5/extras/bwd/cl_program_utilities.c
monetdb5/extras/bwd/operations.c
Branch: bwd
Log Message:
* (hopefully) fixed projective semijoins
Unterschiede (166 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
@@ -61,17 +61,17 @@ cl_program getProjectionLeftjoinProgram(
const char* sourceCode = "__kernel void project (\n"
"__global struct{int count; int padding; char values[];}*
outputTail,\n"
"__global struct{int count; int padding; int positions[];}*
inputTail,"
- "__global const struct{int count; int padding; char values[];}*
approximationTail\n) {\n"
+ "__global const struct{int count; int base; char values[];}*
approximationTail\n) {\n"
" if(get_global_id(0) < inputTail->count){\n"
" __global const char* approximation =
approximationTail->values;"
" const int offset =
inputTail->positions[get_global_id(0)]*approximationBytes;\n"
- " int value = 0;"
+ " int value = approximationTail->base;"
" for(int i = 0; i < approximationBytes; i++){\n"
" outputTail->values[get_global_id(0)*approximationBytes+i]
= approximation[offset + i];\n"
" value += approximation[offset + i] << 8*i;"
" }\n"
+ /* " printf(\"projected value (%d): %d (base: %d + %d +
%d<<8), approximationBytes: %d\\n\", inputTail->positions[get_global_id(0)],
value, approximationTail->base, approximation[offset], approximation[offset +
1], approximationBytes);" */
" }\n"
- /* " printf(\"projected value (%d): %d\\n\",
inputTail->positions[get_global_id(0)], value<<(8*(4-approximationBytes)));" */
"}";
char options[64];
snprintf(options, 64, "-D approximationBytes=%d",
approximationBits/8-offsetBits/8);
@@ -111,7 +111,7 @@ cl_program getUSelectProgram(int type, c
" value += (approximation[inputOffset + i] <<
((i+sizeof(%1$s)-%6$d - %4$d)*8));\n"
"\n"
/* " if(get_global_id(0) < 10)\n" */
- /* " printf(\"%%d, base: %%d\\n\", value,
approximationTail->base);\n" */
+ /* " printf(\"%%d: %%d, base: %%d\\n\", get_global_id(0),
value, approximationTail->base);\n" */
" if((value %2$s operand)"
" && (%5$d || value %3$s operand2)"
" )"
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
@@ -19,6 +19,7 @@
static const int activateWorkInProgress = 1;
static const int synchronousGPU = 1;
+#define WORK_GROUP_SIZE 1
#pragma mark Actual MAL Operations Implementation
@@ -136,7 +137,7 @@ str BWDLeftJoinApproximate(bat * res, ba
)))) 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/16.0)*16}, (const size_t[]){16}, 0, NULL, NULL)))
+ 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());
@@ -151,16 +152,18 @@ str BWDLeftJoinApproximate(bat * res, ba
return MAL_SUCCEED;
};
-static inline int decompressIntValue(const int i, const int approximationbits,
const clTail* compressedTail, const unsigned char* residuals, const int
residualI){
+static inline int decompressIntValue(const int i, const int approximationbits,
const int offsetBits, const clTail* compressedTail, const unsigned char*
residuals, const int residualI){
/* const int index = compressedHead->positions[i];
*/
const int residualBits = 32-approximationbits;
const unsigned int residualMask = (1 << residualBits)-1;
const unsigned int residualBytes = residualBits/8;
+ const int approximationBytes = approximationbits/8-offsetBits/8;
+ const int approximationMask = (1<<(approximationBytes*8))-1;
- const int offset = (approximationbits/8)*i;
- const int compressedValue =
compressedTail->base+*(int*)&(compressedTail->elements[offset]);
- const int deCompressedValue = (compressedValue << residualBits) +
(*(int*)&residuals[residualI*residualBytes] & residualMask);
+ const int offset = approximationBytes*i;
+ const int compressedValue = *(int*)&(compressedTail->elements[offset])
& approximationMask;
+ const int deCompressedValue = compressedTail->base + (compressedValue
<< residualBits) + (*(int*)&residuals[residualI*residualBytes] & residualMask);
return deCompressedValue;
}
@@ -177,6 +180,7 @@ str BWDLeftJoinRefine(bat * res, bat * l
BAT* refinement;
BAT* approximation = BATdescriptor(*approx);
const register size_t approximationBits =
batTailApproximationBits(right);
+ const register size_t offsetBits = batTailOffsetBits(right);
/* 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
@@ -215,7 +219,7 @@ str BWDLeftJoinRefine(bat * res, bat * l
int i;
for (i = 0; i < supersetPositionsColumn->count; ++i) {
if(supersetPositionsColumn->positions[i] ==
*positionRegion){
- refinementRegion[refinementCount++] =
decompressIntValue(i, approximationBits, supersetApproximateValuesColumn,
residuals, *positionRegion);
+ refinementRegion[refinementCount++] =
decompressIntValue(i, approximationBits, offsetBits,
supersetApproximateValuesColumn, residuals, *positionRegion);
positionRegion++;
}
}
@@ -298,7 +302,7 @@ static inline str uselect(bat *res, bat
if((err = clSetKernelArg(selectKernel, 3+i,
sizeof(int), &(parameters[i])))) // type specific
printf("#%s, clSetKernelArg(%d): %s;\n",
__func__, 3+i, clError(err));
- err = clEnqueueNDRangeKernel(getCommandQueue(), selectKernel,
1, (const size_t[]){0}, (const size_t[]){ceil(dataCount/16.0)*16}, (const
size_t[]){16}, 0, NULL, NULL);
+ 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);
if(err) printf("#%s, clEnqueueNDRangeKernel: %s;\n", __func__,
clError(err));
if (synchronousGPU) clFinish(getCommandQueue());
if(0){
@@ -360,6 +364,8 @@ str uselectrefine(bat *res, bat *bid, pt
for (i = 0; i < candidateCount; ++i)
positionRegion[i] = headPositions[i];
BATsetcount(result, candidateCount);
+ BATseqbase(BATmirror(result), oid_nil);
+
} else {
@@ -527,7 +533,7 @@ str BWDSemijoinApproximate(int *res, int
cl_mem positionColumn = batTailPositions(right);
if(!positionColumn) positionColumn =
batHeadApproximation(BATmirror(right));
- cl_program program =
getProjectionLeftjoinProgram(batTailApproximationBits(right),
batTailOffsetBits(right)); // turns out this case is actually a simple
projection and a reuse of the right head column
+ cl_program program =
getProjectionLeftjoinProgram(batTailApproximationBits(left),
batTailOffsetBits(left)); // turns out this case is actually a simple
projection and a reuse of the right head column
cl_int err;
@@ -568,7 +574,7 @@ 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/16.0)*16}, (const size_t[]){16}, 0, NULL, NULL)))
+ 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(0) printf ("returning approximation with tail
positions %p\n", batTailPositions(result));
if (synchronousGPU) clFinish(getCommandQueue());
@@ -602,6 +608,7 @@ str BWDSemijoinRefine(int *res, int *lid
BAT* right = BATdescriptor(*rid);
BAT* approximation = BATdescriptor(*approx);
const size_t approximationBits = batTailApproximationBits(left);
+ const size_t offsetBits = batTailOffsetBits(left);
BAT* refinement;
if(BAThvoid(left) && ATOMstorage(BATttype(left)) == TYPE_int &&
BAThtype(right) == TYPE_oid){ // TODO: type specific
@@ -613,14 +620,15 @@ str BWDSemijoinRefine(int *res, int *lid
getPositionsColumn(batTailPositions(approximation), NULL,
&bufferSize);
clHead* supersetPositionsColumn =
getPositionsColumn(batTailPositions(approximation), malloc(bufferSize),
&bufferSize);
unsigned char* filteredResiduals = NULL;
+ getApproximateValuesColumn(batTailApproximation(approximation),
NULL, &bufferSize);
+ clTail* supersetApproximateValuesColumn =
getApproximateValuesColumn(batTailApproximation(approximation),
malloc(bufferSize), &bufferSize);
- if(supersetPositionsColumn->count == BATcount(right)){ // the
approximation was correct, no refinement necessary
- clEnqueueReadBuffer(getCommandQueue(),
batTailApproximation(approximation), CL_TRUE, sizeof(clTail),
BATcount(right)*sizeof(int), Tloc(refinement, BUNfirst(refinement)), 0, NULL,
NULL);
+ if(0 && supersetPositionsColumn->count == BATcount(right)){ //
the approximation was correct, no refinement necessary
+ /* clEnqueueReadBuffer(getCommandQueue(),
batTailApproximation(approximation), CL_TRUE, sizeof(clTail),
BATcount(right)*sizeof(int), Tloc(refinement, BUNfirst(refinement)), 0, NULL,
NULL); */
+
BATsetcount(refinement, supersetPositionsColumn->count);
} else {
-
getApproximateValuesColumn(batTailApproximation(approximation), NULL,
&bufferSize);
- clTail* supersetApproximateValuesColumn =
getApproximateValuesColumn(batTailApproximation(approximation),
malloc(bufferSize), &bufferSize);
size_t refinementCount = 0;
if(supersetPositionsColumn){
@@ -636,7 +644,7 @@ str BWDSemijoinRefine(int *res, int *lid
for ( i = j = 0; j < BATcount(right);
++j) {
while(supersetPositionsColumn->positions[i] != ((oid*)right->H->heap.base)[j])
i++;
-
refinementRegion[refinementCount++] = decompressIntValue(i, approximationBits,
supersetApproximateValuesColumn, residuals,
supersetPositionsColumn->positions[i]);
+
refinementRegion[refinementCount++] = decompressIntValue(i, approximationBits,
offsetBits, supersetApproximateValuesColumn, residuals,
supersetPositionsColumn->positions[i]);
}
}
}else {
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
