Changeset: 3557a6c06478 for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=3557a6c06478
Modified Files:
monetdb5/extras/bwd/cl_program_utilities.c
Branch: bwd
Log Message:
* working on true bitwise (in contrast to bytewise) decomposition
Unterschiede (113 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
@@ -106,7 +106,7 @@ 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"
+ [0] = "__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"
@@ -121,32 +121,23 @@ cl_program getUSelectProgram(int type, c
" 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"
- /* " if(get_global_id(0) < 10)\n" */
- /* " printf(\"%%d, i=%%d: value: %%d, chunk: %%d\\n\",
get_global_id(0), i, value, approximation[inputOffset + i]);\n" */
-
"}\n"
- /* " if(get_global_id(0) < 10)\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)"
" )"
"{\n"
" const int index = atomic_inc(&(outputHead->count));\n"
" atomic_inc(&(outputTail->count));\n" // TODO: this could
probably be done more efficiently
- /* " printf(\"selected value %%d from slot %%d into slot
%%d, head value_: %%d\\n\", value, get_global_id(0), index,
approximationHead->positions[get_global_id(0)]);\n" */
" const int offset = index * %4$d;\n"
" outputHead->positions[index] =
approximationHead->positions[get_global_id(0)];\n"
" for(int i = 0; i < %4$d; i++){\n"
" outputTail->values[offset+i] = approximation[inputOffset
+ i];\n"
- /* " printf(\"set byte %%d to %%d\\n\",
offset+i,outputTail->values[offset+i]);" */
" }\n"
" } \n"
- /* "else {\n" */
- /* " printf(\"not selected value %%d (base: %%d, outbase:
%%d) from slot %%d, head value: %%d\\n\", value, approximationTail->base,
outputTail->base, get_global_id(0),
approximationHead->positions[get_global_id(0)]);\n" */
- /* " }" */
" }\n"
"}",
- [1] = "__kernel void uselect (\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"
@@ -159,7 +150,6 @@ cl_program getUSelectProgram(int type, c
"const %1$s operand2\n"
") {\n"
" if(get_global_id(0) < approximationTail->count){\n"
- /* " __global const unsigned char* approximation =
approximationTail->values;" */
" %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"
@@ -185,14 +175,55 @@ cl_program getUSelectProgram(int type, c
"#endif\n"
" }\n"
" }\n"
+ "}",
+
+ [2] = "__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 const struct{int count; int base; unsigned char
values[];}* approximationTail,\n"
+ "const %1$s operand,\n"
+ "const %1$s operand2\n"
+ ") {\n"
+ " if(get_global_id(0) < approximationTail->count){\n"
+ " static const size_t targetTypeBits = (sizeof(targetType)*8);"
+ " targetType value = approximationTail->base;\n"
+ " size_t slot =
(get_global_id(0)*approximationBits)/targetTypeBits;"
+ " size_t offset =
(get_global_id(0)*approximationBits)%%targetTypeBits;"
+ " {"
+ " __global const unsigned char* vals =
approximationTail->values;"
+ " value +=
((offset+approximationBits>targetTypeBits)*((vals[slot]<<(targetTypeBits-offset))
+ vals[slot+1]>>(approximationBits-targetTypeBits+offset))\n"
+ " +
(offset+approximationBits>targetTypeBits)*(vals[slot]>>(targetTypeBits-offset-approximationBits)))&((1<<approximationBits)-1);"
+ " }"
+ "\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"
"}"
+
};
char* sourceCode = malloc(16384);
static const str typeNames[] = {[TYPE_int] = "int"};
static const size_t typeSizes[] = {[TYPE_int] = sizeof(int)};
- snprintf(sourceCode, 16384, sourceCodeTemplates[inputIsVoidHeaded!=0],
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);
{
- cl_program program = compileProgram(sourceCode,"");
+ char options[64];
+ cl_program program;
+ snprintf(options, 64, "-D targetType=int -D
approximationBits=%d", approximationBits-offsetBits);
+ printf ("%s\n", sourceCode);
+ program = compileProgram(sourceCode, options);
free(sourceCode);
return program;
}
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
