Changeset: 7f88e30cce54 for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=7f88e30cce54
Modified Files:
monetdb5/extras/bwd/cl_program_utilities.c
monetdb5/extras/bwd/operations.c
monetdb5/extras/bwd/utilities.c
Branch: bwd
Log Message:
* still working on true bitwise decomposition (largely fixing result
correctness bugs)
Unterschiede (gekürzt von 413 auf 300 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
@@ -25,6 +25,10 @@ char* approximateOperation(char* exactOp
return exactOperation;
}
+/* static const str typeNames[] = {[TYPE_int] = "int"}; */
+static const size_t typeSizes[] = {[TYPE_int] = sizeof(int)};
+
+
static inline void superverboseprintf(const char * format, ... ){};
cl_program compileProgram(const char* sourceCode, char* options){
@@ -49,6 +53,8 @@ cl_program compileProgram(const char* so
return program;
}
+
+
cl_program getMultiplyProgram(unsigned int approximationBits){
const char* sourceCode = "__kernel void project (\n"
"__global struct{int count; int base; unsigned char values[];}*
outputTail,\n"
@@ -71,35 +77,56 @@ cl_program getProjectionLeftjoinProgram(
static cl_program programCache[4][8] = {}; //TODO: limited to four
devices
const int deviceForThisThread = getGPUDeviceForThisThread();
if(programCache[deviceForThisThread][approximationBits/8-offsetBits/8]
== NULL){
- const char* sourceCode = "__kernel void project (\n"
- "#if approximationBytes == 2\n"
- "__global struct{int count; int base; unsigned short
values[];}* outputTail,\n"
- "#else\n"
- "__global struct{int count; int base; unsigned char
values[];}* outputTail,\n"
- "#endif\n"
- "__global struct{int count; int padding; int
positions[];}* inputTail,\n"
- "#if approximationBytes == 2\n"
- "__global const struct{int count; int base; unsigned
short values[];}* approximationTail\n) {\n"
- "#else\n"
- "__global const struct{int count; int base; unsigned
char values[];}* approximationTail\n) {\n"
- "#endif\n"
+ const char* sourceCode = ""
+ "#define accessType unsigned int\n"
+ " __constant static const size_t approximationBits =
(approximationBytes*8);\n"
+ " __constant static const size_t targetTypeBits =
(sizeof(accessType)*8);\n"
+ " __constant static const unsigned int
approximationMask = ((1<<approximationBits)-1);\n"
+ "__kernel void project (\n"
+ "__global struct{int count; int base; unsigned int
values[];}* outputTail,\n"
+ "__global struct{int count; int padding; unsigned int
positions[];}* inputTail,\n"
+ "__global const struct{int count; int base; unsigned
int values[];}* approximationTail\n) {\n"
" if(get_global_id(0) < inputTail->count){\n"
/* " int value = approximationTail->base;" */
- "#if approximationBytes == 2\n"
- " outputTail->values[get_global_id(0)] =
approximationTail->values[inputTail->positions[get_global_id(0)]];\n"
- "#else\n"
- " __global const unsigned char* approximation =
approximationTail->values;\n"
- " const int offset =
inputTail->positions[get_global_id(0)]*approximationBytes;\n"
- " for(int i = 0; i < approximationBytes; i++){\n"
- "
outputTail->values[get_global_id(0)*approximationBytes+i] =
approximation[offset + i];\n"
- " }\n"
- "#endif\n"
+ /* "#if approximationBytes == 2\n" */
+ /* " outputTail->values[get_global_id(0)] =
approximationTail->values[inputTail->positions[get_global_id(0)]];\n" */
+ /* "#else\n" */
+ /* " __global const unsigned char* approximation =
approximationTail->values;\n" */
+ /* " const int offset =
inputTail->positions[get_global_id(0)]*approximationBytes;\n" */
+ /* " for(int i = 0; i < approximationBytes; i++){\n" */
+ /* "
outputTail->values[get_global_id(0)*approximationBytes+i] =
approximation[offset + i];\n" */
+ /* " }\n" */
+ /* "#endif\n" */
+ " const size_t inputIndex =
inputTail->positions[get_global_id(0)];"
+ " size_t slot =
(inputIndex*approximationBits)/targetTypeBits;\n"
+ " size_t offset =
(inputIndex*approximationBits)%targetTypeBits;\n"
+ /* " printf(\"getting position %u\\n\",
get_global_id(0));\n" */
+ " __global const unsigned int* vals =
approximationTail->values;\n"
+ " unsigned int delta = (("
+ "
(((offset+approximationBits)>targetTypeBits)?((vals[slot]<<(targetTypeBits-offset))
+ vals[slot+1]>>(approximationBits-targetTypeBits+offset)):0)\n"
+ /* "0" */
+ " +
(((offset+approximationBits)<=targetTypeBits)*(vals[slot]>>(targetTypeBits-offset-approximationBits)))"
+ " )&approximationMask);\n"
+ /* " printf(\"projected delta #%u: %u\\n\",
get_global_id(0), delta);\n" */
+
+ " const size_t index = get_global_id(0);\n"
+ " size_t outslot =
(index*approximationBits)/targetTypeBits;\n"
+ " size_t outoffset =
(index*approximationBits)%targetTypeBits;\n"
+ " if(outoffset+approximationBits >
8*sizeof(accessType))"
+ " outputTail->values[outslot] |= (delta <<
(approximationBits-(8*sizeof(accessType)-outoffset)));"
+ " else"
+ " outputTail->values[outslot] |= (delta <<
(8*sizeof(accessType)-outoffset-approximationBits));"
/* " value += approximation[offset + i] << 8*(i+1);"
*/
+ " int value =
approximationTail->base+(delta<<residualBits);"
+ " if(index<10)"
+ " printf(\"%d->%d: offset+(%d>>x)=%d,
val[slot=%u]=%u, %d\\n\", inputIndex, get_global_id(0), delta, value, outslot,
outputTail->values[outslot],
(8*sizeof(accessType)-outoffset-approximationBits));"
+
/* " printf(\"projected value (%d): %d (base: %d + %d
+ %d<<8, outbase: %d), approximationBytes: %d\\n\",
inputTail->positions[get_global_id(0)], value, approximationTail->base,
approximation[offset], approximation[offset + 1], outputTail->base,
approximationBytes);" */
" }\n"
"}";
- char options[64];
- snprintf(options, 64, "-D approximationBytes=%d",
approximationBits/8-offsetBits/8);
+ char options[256];
+ int type = TYPE_int;
+ snprintf(options, 256, "-D approximationBytes=%d -D
residualBits=%lu", approximationBits/8-offsetBits/8,
typeSizes[type]*8-approximationBits);
programCache[deviceForThisThread][approximationBits/8-offsetBits/8] =
compileProgram(sourceCode, options);
}
return
programCache[deviceForThisThread][approximationBits/8-offsetBits/8];
@@ -108,7 +135,9 @@ cl_program getProjectionLeftjoinProgram(
cl_program getUSelectProgram(int type, char* predicateOperation, char*
predicateOperation2, unsigned int approximationBits, unsigned int offsetBits,
char inputIsVoidHeaded){
const char* sourceCodeTemplates[] = {
- [0] = "__kernel void uselect (\n" // non-void-headed case
+ [0] = " __constant static const size_t targetTypeBits =
(sizeof(targetType)*8);\n"
+ " __constant static const unsigned int approximationMask =
((1<<approximationBits)-1);\n"
+ "__kernel void uselect (\n" // non-void-headed case
"__global struct{int count; int padding; int positions[];}*
outputHead,\n"
"__global struct{int count; int base; unsigned char values[];}*
outputTail,\n"
"__global const struct{int count; int base; unsigned char
values[];}* approximationTail,\n"
@@ -118,12 +147,22 @@ cl_program getUSelectProgram(int type, c
"const %1$s operand2\n"
") {\n"
" if(get_global_id(0) < approximationTail->count){\n"
+ " __global const unsigned int* vals =
approximationTail->values;\n"
" __global const unsigned char* approximation =
approximationTail->values;"
" %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)-%6$d - %4$d)*8));\n"
- "}\n"
+ /* " for(int i = 0; i < %4$d; i++){\n" */
+ /* " value += (approximation[inputOffset + i] <<
((i+sizeof(%1$s)-%6$d - %4$d)*8));\n" */
+ /* "}\n" */
+ " size_t slot =
(get_global_id(0)*approximationBits)/targetTypeBits;\n"
+ " size_t offset =
(get_global_id(0)*approximationBits)%%targetTypeBits;\n"
+
+
+ " value += ((("
+ "
(((offset+approximationBits)>targetTypeBits)?((vals[slot]<<(targetTypeBits-offset))
+ vals[slot+1]>>(approximationBits-targetTypeBits+offset)):0)\n"
+ " +
(((offset+approximationBits)<=targetTypeBits)*(vals[slot]>>(targetTypeBits-offset-approximationBits)))"
+ " )&approximationMask)<<residualBits);\n"
+
" if((value %2$s operand)"
" && (%5$d || value %3$s operand2)"
" )"
@@ -180,11 +219,12 @@ cl_program getUSelectProgram(int type, c
" }\n"
"}",
- [2] = " __constant static const size_t targetTypeBits =
(sizeof(targetType)*8);\n"
+ [2] = "#define accessType unsigned int\n"
+ " __constant static const size_t targetTypeBits =
(sizeof(targetType)*8);\n"
" __constant static const unsigned int approximationMask =
((1<<approximationBits)-1);\n"
"__kernel void uselect (\n" // void-headed case
"__global struct{int count; int padding; int positions[];}*
outputHead,\n"
- "__global struct{int count; int base; unsigned char values[];}*
outputTail,\n"
+ "__global struct{int count; int base; accessType values[];}*
outputTail,\n"
"__global const struct{int count; int base; unsigned char
values[];}* approximationTail,\n"
"const %1$s operand,\n"
"const %1$s operand2\n"
@@ -193,41 +233,42 @@ cl_program getUSelectProgram(int type, c
" targetType value = approximationTail->base;\n"
" size_t slot =
(get_global_id(0)*approximationBits)/targetTypeBits;\n"
" size_t offset =
(get_global_id(0)*approximationBits)%%targetTypeBits;\n"
- " {\n"
" __global const unsigned int* vals =
approximationTail->values;\n"
- " value += ((("
+ " unsigned int delta = (("
"
(((offset+approximationBits)>targetTypeBits)?((vals[slot]<<(targetTypeBits-offset))
+ vals[slot+1]>>(approximationBits-targetTypeBits+offset)):0)\n"
" +
(((offset+approximationBits)<=targetTypeBits)*(vals[slot]>>(targetTypeBits-offset-approximationBits)))"
- " )&approximationMask)<<residualBits);\n"
- /* " value += (((\n" */
- /* " +
((vals[slot]>>(targetTypeBits-offset-approximationBits))))&((1<<approximationBits)-1))<<residualBits);\n"
*/
- GPU_PRINTF(" printf(\"===== %%d in slot %%d, offset %%d,
targetTypeBits: %%d, approximationBits: %%d (global id: %%d) = %%d\\n\", value,
slot, offset, targetTypeBits, approximationBits, get_global_id(0),
(((offset+approximationBits)>targetTypeBits)*((vals[slot]<<(targetTypeBits-offset))
+ vals[slot+1]>>(approximationBits-targetTypeBits+offset))\n")
- GPU_PRINTF(" +
((offset+approximationBits)<=targetTypeBits)*(vals[slot]>>(targetTypeBits-offset-approximationBits))));\n")
- " }\n"
+ " )&approximationMask);\n"
+ " value += (delta<<residualBits);"
"\n"
" if((value %2$s operand)"
" && (%5$d || value %3$s operand2)"
" ){\n"
" const int index = atomic_inc(&(outputHead->count));\n"
" atomic_inc(&(outputTail->count));\n" // TODO: this could
probably be done more efficiently
+ " size_t outslot =
(index*approximationBits)/targetTypeBits;\n"
+ " size_t outoffset =
(index*approximationBits)%%targetTypeBits;\n"
" outputHead->positions[index] = get_global_id(0);\n"
- "#if %4$d == 2\n"
- " outputTail->values[index] =
approximationTail->values[get_global_id(0)];\n"
- "#else\n"
- " const int offset = index * %4$d;\n"
- " for(int i = 0; i < %4$d; i++){\n"
- " outputTail->values[offset+i] =
approximationTail->values[inputOffset + i];\n"
- " }\n"
- "#endif\n"
+ /* "#if %4$d == 2\n" */
+ /* " outputTail->values[index] =
approximationTail->values[get_global_id(0)];\n" */
+ /* "#else\n" */
+ /* " const int offset = index * %4$d;\n" */
+ /* " for(int i = 0; i < %4$d; i++){\n" */
+ /* " outputTail->values[offset+i] =
approximationTail->values[inputOffset + i];\n" */
+ /* " }\n" */
+ /* "#endif\n" */
+ " if(outoffset+approximationBits > 8*sizeof(accessType))"
+ " outputTail->values[outslot] |= (delta <<
(approximationBits-(8*sizeof(accessType)-outoffset)));"
+ " else"
+ " outputTail->values[outslot] |= (delta <<
(8*sizeof(accessType)-outoffset-approximationBits));"
+ " if(index<10)"
+ " printf(\"%%d: offset+(%%d>>x)=%%d, val[slot=%%u]=%%u,
%%d\\n\", get_global_id(0), delta, value, outslot, outputTail->values[outslot],
(8*sizeof(accessType)-outoffset-approximationBits));"
" }\n"
" }\n"
"}"
};
char* sourceCode = malloc(16384);
- const str typeNames[] = {[TYPE_int] = "int"};
- const size_t typeSizes[] = {[TYPE_int] = sizeof(int)};
- snprintf(sourceCode, 16384,
sourceCodeTemplates[(!!inputIsVoidHeaded)*1], typeNames[type],
(approximationBits ==
8*typeSizes[type])?predicateOperation:(approximateOperation(predicateOperation)),
predicateOperation2?((approximationBits ==
8*typeSizes[type])?predicateOperation2:approximateOperation(predicateOperation2)):"==",
approximationBits/8-offsetBits/8, predicateOperation2 == NULL?1:0,
offsetBits/8);
+ snprintf(sourceCode, 16384,
sourceCodeTemplates[(!!inputIsVoidHeaded)*2], typeNames[type],
(approximationBits ==
8*typeSizes[type])?predicateOperation:(approximateOperation(predicateOperation)),
predicateOperation2?((approximationBits ==
8*typeSizes[type])?predicateOperation2:approximateOperation(predicateOperation2)):"==",
approximationBits/8-offsetBits/8, predicateOperation2 == NULL?1:0,
offsetBits/8);
{
char options[256];
cl_program program;
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
@@ -205,18 +205,30 @@ str BWDLeftJoinApproximate(bat * res, ba
return MAL_SUCCEED;
};
-static inline int decompressIntValue(const int i, const int approximationbits,
const int offsetBits, const clTail* compressedTail, const unsigned char*
residuals, const int residualI){
+static inline int decompressIntValue(const int approximationI, const int
approximationbits, const int offsetBits, const clTail* compressedTail, const
unsigned int* 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 unsigned int residualBytes = residualBits/8; */
const int approximationBytes = approximationbits/8-offsetBits/8;
const int approximationMask = (1<<(approximationBytes*8))-1;
- 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);
+ const int slotI =
(approximationBytes*8)*approximationI/(sizeof(int)*8);
+ const int offset =
((approximationBytes*8)*approximationI)%(sizeof(int)*8);
+ const unsigned int compressedValue = (((unsigned
int*)compressedTail->elements)[slotI] >>
(8*sizeof(int)-offset-(approximationBytes*8))) & approximationMask;
+ const size_t residualSlotI = (residualI*residualBits)/32;
+ const unsigned int residualOffset = (residualI*residualBits)%32;
+ const unsigned int residual = (((unsigned int*)residuals)[residualSlotI]
+
>>
(32-residualOffset-residualBits))&residualMask;
+ const int deCompressedValue = compressedTail->base+
+ (compressedValue << residualBits)
+ + residual;
+
+
+ /* 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;
}
@@ -275,7 +287,7 @@ str BWDLeftJoinRefine(bat * res, bat * l
oid* positionRegion = (oid*) Tloc(left,
BUNfirst(left));
size_t refinementCount = 0;
- const unsigned char* residuals =
batTailResiduals(right);
+ const unsigned int* residuals = (unsigned
int*)batTailResiduals(right);
uint i;
for (i = 0; i < supersetPositionsColumn->count;
++i) {
if(supersetPositionsColumn->positions[i] == *positionRegion){
@@ -389,7 +401,7 @@ static inline str uselect(bat *res, bat
if(err)
printf("#%s,
clEnqueueNDRangeKernel: %s;\n", __func__, clError(err));
if (synchronousGPU)
clFinish(getCommandQueue());
- if(0){
+ if(1){
clTail* compressedTail;
size_t resultSize;
printf ("%s result tail result:
%p\n", __func__, batTailApproximation(result));
@@ -399,6 +411,7 @@ static inline str uselect(bat *res, bat
compressedTail =
malloc(resultSize);
err =
clEnqueueReadBuffer(getCommandQueue(), batTailApproximation(result), CL_TRUE,
0, resultSize, compressedTail , 0, NULL, NULL);
if(err) printf("#%s,
clEnqueueReadBuffer: %s;\n", __func__, clError(err));
+ (void) compressedTail;
}
}
}
@@ -429,11 +442,18 @@ static inline unsigned int refinementLoo
#define refineLoopDoubleOperator(comparator, comparator2,
positionIsTruePositiveCondition, residualPosition) \
while(i < candidateCount) {
\
const unsigned int index = compressedHead->positions[i];
\
- const int offset = tailApproximationBytes*i;
\
- const int compressedValue =
*(int*)&(compressedTail->elements[offset]) & tailApproximationMask; \
+ const int slotI = tailApproximationBytes*8*i/(sizeof(int)*8);
\
+ const int offset =
(tailApproximationBytes*8*i)%(sizeof(int)*8); \
+ const unsigned int compressedValue = (((unsigned
int*)compressedTail->elements)[slotI] >>
(8*sizeof(int)-offset-tailApproximationBytes*8)) & tailApproximationMask; \
+ const size_t residualSlotI = (index*tailResidualBits)/32;
\
+ const unsigned int residualOffset =
(index*tailResidualBits)%32; \
+ const unsigned int residual = (((unsigned
int*)residuals)[residualSlotI] \
+
>>
(32-residualOffset-tailResidualBits))&residualMask; \
const int deCompressedValue = compressedTail->base+
\
(compressedValue << tailResidualBits)
\
- + (*(unsigned
int*)&residuals[residualPosition*residualBytes] & residualMask); \
+ + residual;
\
+ if(i<10) printf ("%u: %d (residual: %d)\n", index,
\
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
