Changeset: a00466f6b369 for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=a00466f6b369
Modified Files:
monetdb5/extras/bwd/91_bwd.sql
monetdb5/extras/bwd/bwd.c
monetdb5/extras/bwd/cl_program_utilities.c
monetdb5/extras/bwd/cl_program_utilities.h
monetdb5/extras/bwd/operations.c
monetdb5/extras/bwd/optimizer.c
monetdb5/extras/bwd/utilities.c
monetdb5/extras/bwd/utilities.h
Branch: bwd
Log Message:
* introduced (global) prefix compression for approximations
* made code a little bit more C89 compliant
* fixed minor bugs
Unterschiede (gekürzt von 619 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
@@ -1,11 +1,11 @@
-create function bwdecompose(col decimal(7,2), bits integer)
+create function bwdecomposedec(col decimal(9,2), bits integer)
returns varchar(4096) external name bwd.decompose;
create function bwdecomposedate(col date, bits integer)
returns varchar(4096) external name bwd.decompose;
-create function bwdecompose(col integer, bits integer)
+create function bwdecomposeint(col integer, bits integer)
returns varchar(4096) external name bwd.decompose;
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
@@ -102,13 +102,16 @@ str bwdecompose(bat * res, bat * subject
BAT* subject = BATdescriptor(*subjectBAT);
if (VIEWtparent(subject))
subject = BATdescriptor(abs(VIEWtparent(subject)));
- printf ("decomposing %s into %d and %d bits\n",
resolveBatToAttribute(subject->batCacheid, cntxt), approximationBits[0],
Tsize(subject)*8-*approximationBits);
- BATsetprop(subject, batRegistryIndex, TYPE_int,
(int[]){decomposeIntArray((int*)Tloc(subject, BUNfirst(subject)),
subject->batCount, *approximationBits)});
- BAT* result = BATnew(TYPE_void, TYPE_str, 0);
- char buffer[4096];
- snprintf(buffer, 4096, "successfully decomposed bat %d",
subject->batCacheid);
- BUNappend(result, buffer, 0);
- BBPkeepref(*res = result->batCacheid);
+ if(! batTailIsDecomposed(subject)){
+ printf ("decomposing %s into %d and %d bits\n",
resolveBatToAttribute(subject->batCacheid, cntxt), approximationBits[0],
Tsize(subject)*8-*approximationBits);
+ BATsetprop(subject, batRegistryIndex, TYPE_int,
(int[]){decomposeIntArray((int*)Tloc(subject, BUNfirst(subject)),
subject->batCount, *approximationBits)});
+ BAT* result = BATnew(TYPE_void, TYPE_str, 0);
+ char buffer[4096];
+ snprintf(buffer, 4096, "successfully decomposed bat %d",
subject->batCacheid);
+ BUNappend(result, buffer, 0);
+ BBPkeepref(*res = result->batCacheid);
+ }
+ BBPreleaseref(subject->batCacheid);
return MAL_SUCCEED;
}
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
@@ -57,7 +57,7 @@ cl_program getMultiplyProgram(unsigned i
return compileProgram(sourceCode, options);
}
-cl_program getProjectionLeftjoinProgram(unsigned int approximationBits){
+cl_program getProjectionLeftjoinProgram(unsigned int approximationBits,
unsigned int offsetBits){
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,"
@@ -72,7 +72,7 @@ cl_program getProjectionLeftjoinProgram(
/* " 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);
+ snprintf(options, 64, "-D approximationBytes=%d",
approximationBits/8-offsetBits/8);
return compileProgram(sourceCode, options);
}
@@ -91,21 +91,24 @@ cl_program getProjectionLeftjoinProgram(
/* return compileProgram(sourceCode, options); */
/* } */
-cl_program getUSelectProgram(int type, char* predicateOperation, char*
predicateOperation2, unsigned int approximationBits){
+cl_program getUSelectProgram(int type, char* predicateOperation, char*
predicateOperation2, unsigned int approximationBits, unsigned int offsetBits){
+
const char* sourceCodeTemplate = "__kernel void uselect (\n"
"__global struct{int count; int padding; int positions[];}*
outputHead,\n"
- "__global struct{int count; int padding; char values[];}*
outputTail,\n"
- "__global const struct{int count; int padding; char values[];}*
approximationTail\n"
+ "__global struct{int count; int base; char values[];}*
outputTail,\n"
+ "__global const struct{int count; int base; char values[];}*
approximationTail\n"
",\n"
"const %1$s operand,\n"
"const %1$s operand2\n"
") {\n"
" __global const char* approximation =
approximationTail->values;"
- " %1$s value = 0;\n"
+ " %1$s value = approximationTail->base;\n"
" const size_t inputOffset = get_global_id(0)*%4$d;\n"
" for(int i = 0; i < %4$d; i++)\n"
- " value += (approximation[inputOffset + i] <<
((i+sizeof(%1$s) - %4$d)*8));\n"
+ " 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" */
" if((value %2$s operand)"
" && (%5$d || value %3$s operand2)"
" )"
@@ -122,7 +125,7 @@ cl_program getUSelectProgram(int type, c
" }\n"
"}";
char* sourceCode = malloc(16384);
- snprintf(sourceCode, 16384, sourceCodeTemplate, (str[]){[TYPE_int] =
"int"}[type], approximateOperation(predicateOperation),
approximateOperation(predicateOperation2), approximationBits/8,
predicateOperation2 == NULL?1:0);
+ snprintf(sourceCode, 16384, sourceCodeTemplate, (str[]){[TYPE_int] =
"int"}[type], approximateOperation(predicateOperation),
approximateOperation(predicateOperation2), approximationBits/8-offsetBits/8,
predicateOperation2 == NULL?1:0, offsetBits/8);
cl_program program = compileProgram(sourceCode,"");
free(sourceCode);
return program;
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
@@ -8,7 +8,7 @@
#endif
-cl_program getUSelectProgram(int type, char* predicateOperation, char*
predicateOperation2, unsigned int approximationBits);
-cl_program getProjectionLeftjoinProgram(unsigned int approximationBits);
+cl_program getUSelectProgram(int type, char* predicateOperation, char*
predicateOperation2, unsigned int approximationBits, unsigned int offsetBits);
+cl_program getProjectionLeftjoinProgram(unsigned int approximationBits,
unsigned int offsetBits);
/* 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
@@ -95,7 +95,7 @@ str BWDLeftJoinApproximate(bat * res, ba
if(!leftColumn) leftColumn =
batHeadApproximation(BATmirror(left));
cl_mem rightColumn = batTailApproximation(right);
if(!rightColumn) rightColumn =
batHeadApproximation(BATmirror(right));
- cl_program program =
getProjectionLeftjoinProgram(batTailApproximationBits(right));
+ cl_program program =
getProjectionLeftjoinProgram(batTailApproximationBits(right),
batTailOffsetBits(right));
cl_int err;
cl_kernel projectKernel = clCreateKernel(program, "project",
&err);
@@ -105,24 +105,29 @@ str BWDLeftJoinApproximate(bat * res, ba
if((err = clEnqueueReadBuffer(getCommandQueue(), leftColumn,
CL_TRUE, 0, sizeof(int), &headCount , 0, NULL, NULL)))
printf("#%s, clEnqueueReadBuffer: %s;\n", __func__,
clError(err));
+ clTail 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));
+
BAT* result = BATnew(TYPE_void, ATOMtype(right->ttype), 0);
{
const unsigned int newIndex =
getNextFreeDecomposedBATSlotIndex();
DecomposedBATSlot* slot =
getDecomposedBATSlotForIndex(newIndex);
BATsetprop(result, batRegistryIndex, TYPE_int,
(int[]){newIndex});
slot->approximationBits =
batTailApproximationBits(right);
+ slot->tailOffsetBits = batTailOffsetBits(right);
slot->residuals = NULL;
slot->tailPositions = leftColumn;
slot->tailApproximation =
clCreateBuffer(getCLContext(), CL_MEM_READ_WRITE,
headCount*slot->approximationBits/8+sizeof(clTail), NULL, &err);
assert(slot->tailPositions);
if(err) printf("#%s, clCreateBuffer: %s;\n", __func__,
clError(err));
- clEnqueueWriteBuffer(getCommandQueue(),
slot->tailApproximation, CL_TRUE, 0, sizeof(int), &headCount, 0, NULL, NULL);
// I wonder what is faster transfering a single integer to the GPU or running a
kernel that initializes a value
+ clEnqueueWriteBuffer(getCommandQueue(),
slot->tailApproximation, CL_TRUE, 0, sizeof(clTail), &newTailDefinition, 0,
NULL, NULL);
if(err) printf("#%s, clCreateBuffer: %s;\n", __func__,
clError(err));
}
-
- for (int i = 0; i < 3; ++i) {
+ int i;
+ for (i = 0; i < 3; ++i) {
if((err = clSetKernelArg(projectKernel, i,
sizeof(cl_mem), &((cl_mem[]){
batTailApproximation(result), leftColumn, rightColumn}[i]
)))) printf("#%s,
clSetKernelArg(%d): %s;\n", __func__, 0, clError(err));
@@ -150,7 +155,7 @@ static inline int decompressIntValue(int
const int offset = (approximationbits/8)*i;
- const int compressedValue = *(int*)&(compressedTail->elements[offset]);
+ const int compressedValue =
compressedTail->base+*(int*)&(compressedTail->elements[offset]);
const int deCompressedValue = (compressedValue << residualBits) +
(*(int*)&residuals[residualI*residualBytes] & residualMask);
return deCompressedValue;
}
@@ -203,8 +208,8 @@ str BWDLeftJoinRefine(bat * res, bat * l
oid* positionRegion = (oid*) Tloc(left, BUNfirst(left));
size_t refinementCount = 0;
const unsigned char* residuals = batTailResiduals(right);
-
- for (int i = 0; i < supersetPositionsColumn->count; ++i) {
+ int i;
+ for (i = 0; i < supersetPositionsColumn->count; ++i) {
if(supersetPositionsColumn->positions[i] ==
*positionRegion){
refinementRegion[refinementCount++] =
decompressIntValue(i, approximationBits, supersetApproximateValuesColumn,
residuals, *positionRegion);
positionRegion++;
@@ -254,6 +259,7 @@ static inline str uselect(bat *res, bat
DecomposedBATSlot* slot =
getDecomposedBATSlotForIndex(newIndex);
BATsetprop(result, batRegistryIndex, TYPE_int,
(int[]){newIndex});
slot->approximationBits =
batTailApproximationBits(data);
+ slot->tailOffsetBits = batTailOffsetBits(data);
slot->residuals = NULL;
cl_int err;
slot->tailApproximation =
clCreateBuffer(getCLContext(), CL_MEM_READ_WRITE,
dataCount*slot->approximationBits/8+sizeof(clTail), NULL, &err);
@@ -268,7 +274,7 @@ static inline str uselect(bat *res, bat
}
- cl_program program =
getUSelectProgram(ATOMstorage(BATttype(data)), OP, OP2?OP2:"==",
batTailApproximationBits(data));
+ cl_program program =
getUSelectProgram(ATOMstorage(BATttype(data)), OP, OP2?OP2:"==",
batTailApproximationBits(data), batTailOffsetBits(data));
int err;
@@ -276,14 +282,15 @@ static inline str uselect(bat *res, bat
if(err) printf("#%s, clCreateKernel: %s;\n", __func__,
clError(err));
cl_mem buffers[3] = {batHeadApproximation(result),
batTailApproximation(result),
- batTailApproximation(data)};
- for (int i = 0; i < 3; ++i)
+
batTailApproximation(data)};
+ int i;
+ for ( i = 0; i < 3; ++i)
if((err = clSetKernelArg(selectKernel, i,
sizeof(cl_mem), &(buffers[i]))))
printf("#%s, clSetKernelArg(%d): %s;\n",
__func__, i, clError(err));
const int parameters[2] = {
((*(int*)val)>>batTailResidualBits(data))<<batTailResidualBits(data),
val2?(((*(int*)val2)>>batTailResidualBits(data))<<batTailResidualBits(data)):0};
- for (int i = 0; i < 2; ++i)
+ for (i = 0; i < 2; ++i)
if((err = clSetKernelArg(selectKernel, 3+i,
sizeof(int), &(parameters[i])))) // type specific
printf("#%s, clSetKernelArg(%d): %s;\n",
__func__, 3+i, clError(err));
@@ -302,7 +309,7 @@ static inline str uselect(bat *res, bat
if(err) printf("#%s, clEnqueueReadBuffer:
%s;\n", __func__, clError(err));
}
}
- 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)":"");
+ if(1) 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)":"");
BBPkeepref(*res = result->batCacheid);
BBPreleaseref(*bid);
return MAL_SUCCEED;
@@ -364,79 +371,88 @@ str uselectrefine(bat *res, bat *bid, pt
const unsigned int residualBytes =
batTailResidualBits(data)/8;
int i = 0, j = 0;
-
+ /* struct timespec before, after; */
+ /* clock_gettime(CLOCK_THREAD_CPUTIME_ID, &before); */
if(1){
+ const int tailApproximationBytes =
batTailApproximationBits(approximation)/8;
+ const int tailResidualBits =
batTailResidualBits(approximation);
+ const int value1 = *(int*)val;
+ const int value2 = val2?*(int*)val2:0;
#define refineLoopDoubleOperator(comparator, comparator2)
\
- while(i < candidateCount) {
\
- const int index =
compressedHead->positions[i];
\
- const int offset =
(batTailApproximationBits(approximation)/8)*i++; \
- const int compressedValue =
*(int*)&(compressedTail->elements[offset]); \
- const int deCompressedValue =
\
- (compressedValue <<
batTailResidualBits(approximation)) \
+ while(i < candidateCount) {
\
+ const int index =
compressedHead->positions[i]; \
+ const int offset =
tailApproximationBytes*i++; \
+ const int compressedValue =
compressedTail->base+*(int*)&(compressedTail->elements[offset]); \
+ const int deCompressedValue =
\
+ (compressedValue <<
tailResidualBits)
\
+
(*(int*)&residuals[index*residualBytes] & residualMask); \
- if(deCompressedValue comparator
*(int*)val &&
\
- deCompressedValue comparator2
*(int*)val2) { \
- positionRegion[j++] = index;
\
- }
\
- }
+ {
\
+ positionRegion[j] =
index;
\
+ j+= ((deCompressedValue
comparator value1) & \
+
(deCompressedValue comparator2 value2));
\
+ }
\
+ }
#define refineLoopSingleOperator(comparator)
\
- while(i < candidateCount) {
\
+ while(i < candidateCount) {
\
const int index = compressedHead->positions[i];
\
- const int offset =
(batTailApproximationBits(approximation)/8)*i++; \
- const int compressedValue =
*(int*)&(compressedTail->elements[offset]); \
+ const int offset =
(tailApproximationBytes)*i++;
\
+ const int compressedValue =
compressedTail->base+*(int*)&(compressedTail->elements[offset]); \
const int deCompressedValue =
\
- (compressedValue <<
batTailResidualBits(approximation)) \
+ (compressedValue << tailResidualBits)
\
+
(*(int*)&residuals[index*residualBytes] & residualMask); \
- if(deCompressedValue comparator *(int*)val) {
\
+ if(deCompressedValue comparator value1) {
\
positionRegion[j++] = index;
\
}
\
}
-#define refineLoop(comparator, comparator2)
\
- if(comparator2 == NULL) {
\
- refineLoopSingleOperator(comparator);
\
- } else {
\
- switch (comparator2[0]){
\
- case '<':
\
- switch((comparator2)[1]){
\
- case '\0':
\
-
refineLoopDoubleOperator(comparator, <); \
- case '=':
\
-
refineLoopDoubleOperator(comparator, <=); \
- }
\
- case '>':
\
- switch((comparator2)[1]){
\
- case '\0':
\
-
refineLoopDoubleOperator(comparator, >); \
- case '=':
\
-
refineLoopDoubleOperator(comparator, <=); \
- }
\
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
