Changeset: efb650c56939 for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=efb650c56939
Modified Files:
monetdb5/extras/bwd/cl_program_utilities.c
Branch: bwd
Log Message:
* true bitwise approximation of selections seems to work, let's get to the
refinement
Unterschiede (174 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
@@ -66,7 +66,6 @@ cl_program getMultiplyProgram(unsigned i
" outputTail->values[get_global_id(0)*approximationBytes+i]
= approximation[offset + i];\n"
" value += approximation[offset + i] << 8*i;"
" }\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);
@@ -79,7 +78,6 @@ cl_program getProjectionLeftjoinProgram(
if(programCache[deviceForThisThread][approximationBits/8-offsetBits/8]
== NULL){
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"
@@ -92,24 +90,24 @@ cl_program getProjectionLeftjoinProgram(
" size_t slot =
(inputIndex*approximationBits)/targetTypeBits;\n"
" size_t offset =
(inputIndex*approximationBits)%targetTypeBits;\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"
- " +
(((offset+approximationBits)<=targetTypeBits)*(vals[slot]>>(targetTypeBits-offset-approximationBits)))"
- " )&approximationMask);\n"
+
+ " const unsigned int delta = (("
+ "
(((offset+approximationBits)>targetTypeBits)?(((vals[slot]<<(approximationBits-targetTypeBits+offset))
+
(vals[slot+1]>>(targetTypeBits-(approximationBits-targetTypeBits+offset))))&approximationMask):0)\n"
+ " +
(((offset+approximationBits)<=targetTypeBits)*(vals[slot]>>(targetTypeBits-offset-approximationBits)))"
+ " )&approximationMask);\n"
" size_t outslot =
(index*approximationBits)/targetTypeBits;\n"
" size_t outoffset =
(index*approximationBits)%targetTypeBits;\n"
- " if(outoffset+approximationBits >
8*sizeof(accessType))"
- " atomic_add(&(outputTail->values[outslot]),
(delta << (approximationBits-(8*sizeof(accessType)-outoffset))));"
- " else"
- " atomic_add(&(outputTail->values[outslot]),
(delta << (8*sizeof(accessType)-outoffset-approximationBits)));"
- " int value =
approximationTail->base+(delta<<residualBits);"
- " if(index==2114)"
- " 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));"
+ " if(outoffset+approximationBits > targetTypeBits){"
+ " atomic_add(&(outputTail->values[outslot]),
(delta >> (outoffset+approximationBits-targetTypeBits)));"
+ " atomic_add(&(outputTail->values[outslot+1]),
(delta << (targetTypeBits-(outoffset+approximationBits-targetTypeBits))));"
+ " }else{"
+ " atomic_add(&(outputTail->values[outslot]),
(delta << (targetTypeBits-outoffset-approximationBits)));"
+ " }"
" }\n"
"}";
char options[256];
int type = TYPE_int;
- snprintf(options, 256, "-D approximationBytes=%d -D
residualBits=%lu", approximationBits/8-offsetBits/8,
typeSizes[type]*8-approximationBits);
+ snprintf(options, 256, "-D approximationBits=%d -D
residualBits=%lu", approximationBits,
typeSizes[type]*8-approximationBits-offsetBits);
programCache[deviceForThisThread][approximationBits/8-offsetBits/8] =
compileProgram(sourceCode, options);
}
return
programCache[deviceForThisThread][approximationBits/8-offsetBits/8];
@@ -134,17 +132,15 @@ cl_program getUSelectProgram(int type, c
" __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" */
" size_t slot =
(get_global_id(0)*approximationBits)/targetTypeBits;\n"
" size_t offset =
(get_global_id(0)*approximationBits)%%targetTypeBits;\n"
+ " const unsigned int delta = (("
+ "
(((offset+approximationBits)>targetTypeBits)?(((vals[slot]<<(approximationBits-targetTypeBits+offset))
+
(vals[slot+1]>>(targetTypeBits-(approximationBits-targetTypeBits+offset))))&approximationMask):0)\n"
+ " +
(((offset+approximationBits)<=targetTypeBits)*(vals[slot]>>(targetTypeBits-offset-approximationBits)))"
+ " )&approximationMask);\n"
+ " value += (delta<<residualBits);\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)"
@@ -161,48 +157,7 @@ cl_program getUSelectProgram(int type, c
" }\n"
"}",
- [1] = "__kernel void uselect (\n" // void-headed case
- "__global struct{int count; int padding; int positions[];}*
outputHead,\n"
- "#if %4$d == 2\n"
- "__global struct{int count; int base; unsigned short
values[];}* outputTail,\n"
- "__global const struct{int count; int base; unsigned short
values[];}* approximationTail,\n"
- "#else\n"
- "__global struct{int count; int base; unsigned char values[];}*
outputTail,\n"
- "__global const struct{int count; int base; unsigned char
values[];}* approximationTail,\n"
- "#endif\n"
- "const %1$s operand,\n"
- "const %1$s operand2\n"
- ") {\n"
- " if(get_global_id(0) < approximationTail->count){\n"
- " %1$s value = approximationTail->base;\n"
- "#if %4$d == 2\n"
- " value += (approximationTail->values[get_global_id(0)] <<
((sizeof(%1$s)-%6$d - %4$d)*8));\n"
- "#else\n"
- " const size_t inputOffset = get_global_id(0)*%4$d;\n"
- " for(int i = 0; i < %4$d; i++)\n"
- " value += (approximationTail->values[inputOffset + i] <<
((i+sizeof(%1$s)-%6$d - %4$d)*8));\n"
- "#endif\n"
- GPU_PRINTF(" printf(\"===== %%d (global id: %%d)\\n\", value,
get_global_id(0));")
- "\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
- " 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"
- " }\n"
- " }\n"
- "}",
-
- [2] = "#define accessType unsigned int\n"
+ [1] = "#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
@@ -217,12 +172,11 @@ cl_program getUSelectProgram(int type, c
" size_t slot =
(get_global_id(0)*approximationBits)/targetTypeBits;\n"
" size_t offset =
(get_global_id(0)*approximationBits)%%targetTypeBits;\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"
+ " const unsigned int delta = (("
+ "
(((offset+approximationBits)>targetTypeBits)?(((vals[slot]<<(approximationBits-targetTypeBits+offset))
+
(vals[slot+1]>>(targetTypeBits-(approximationBits-targetTypeBits+offset))))&approximationMask):0)\n"
" +
(((offset+approximationBits)<=targetTypeBits)*(vals[slot]>>(targetTypeBits-offset-approximationBits)))"
" )&approximationMask);\n"
- " value += (delta<<residualBits);"
- "\n"
+ " value += (delta<<residualBits);\n"
" if((value %2$s operand)"
" && (%5$d || value %3$s operand2)"
" ){\n"
@@ -231,27 +185,17 @@ cl_program getUSelectProgram(int type, c
" 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(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);
- 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);
+ snprintf(sourceCode, 16384, sourceCodeTemplates[!!inputIsVoidHeaded],
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;
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
