Changeset: 30121cba2b0f for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=30121cba2b0f
Modified Files:
monetdb5/extras/bwd/cl_program_utilities.c
Branch: bwd
Log Message:
* switched from generating selection kernel c-code myself to using compiler
definitions
Unterschiede (124 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
@@ -114,49 +114,44 @@ cl_program getProjectionLeftjoinProgram(
}
cl_program getUSelectProgram(int type, char* predicateOperation, char*
predicateOperation2, unsigned int approximationBits, unsigned int offsetBits,
char inputIsVoidHeaded){
-
const char* sourceCodeTemplates[] = {
- [0] = " __constant static const size_t targetTypeBits =
(sizeof(targetType)*8);\n"
+ [0] = "#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" // 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"
- "__global const struct{int count; int padding; int
positions[];}* approximationHead\n"
- ",\n"
- "const %1$s operand,\n"
- "const %1$s operand2\n"
+ "__global const struct{int count; int padding; int
positions[];}* approximationHead,\n"
+ "const targetType operand,\n"
+ "const targetType 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"
+ " 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"
-
+ " 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"
-
-
- " if((value %2$s operand)"
- " && (%5$d || value %3$s operand2)"
- " )"
- "{\n"
+ " if((value firstOperator operand)"
+ " && (secondOperatorUnset || value secondOperator operand2)"
+ " ){\n"
" const int index = atomic_inc(&(outputHead->count));\n"
" atomic_inc(&(outputTail->count));\n" // TODO: this could
probably be done more efficiently
- " 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"
+ " 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));"
" }\n"
" } \n"
- " }\n"
- "}",
-
+ " }\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"
@@ -164,26 +159,26 @@ cl_program getUSelectProgram(int type, c
"__global struct{int count; int padding; int positions[];}*
outputHead,\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"
+ "const targetType operand,\n"
+ "const targetType operand2\n"
") {\n"
" if(get_global_id(0) < approximationTail->count){\n"
" 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"
+ " size_t offset =
(get_global_id(0)*approximationBits)%targetTypeBits;\n"
" __global const unsigned int* vals =
approximationTail->values;\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"
- " if((value %2$s operand)"
- " && (%5$d || value %3$s operand2)"
+ " if((value firstOperator operand)"
+ " && (secondOperatorUnset || value secondOperator 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"
+ " size_t outoffset =
(index*approximationBits)%targetTypeBits;\n"
" outputHead->positions[index] = get_global_id(0);\n"
" if(outoffset+approximationBits > 8*sizeof(accessType))"
" outputTail->values[outslot] |= (delta <<
(approximationBits-(8*sizeof(accessType)-outoffset)));"
@@ -194,15 +189,15 @@ cl_program getUSelectProgram(int type, c
"}"
};
- char* sourceCode = malloc(16384);
- 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* sourceCode = malloc(16384); */
+ /* 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;
- snprintf(options, 256, "-D targetType=int -D
approximationBits=%d -D residualBits=%lu", approximationBits,
typeSizes[type]*8-offsetBits-approximationBits);
+ snprintf(options, 256, "-DtargetType=%s -DapproximationBits=%d
-DresidualBits=%lu -DfirstOperator=%s -DsecondOperator=%s
-DsecondOperatorUnset=%d", typeNames[type], approximationBits,
typeSizes[type]*8-offsetBits-approximationBits, (approximationBits ==
8*typeSizes[type])?predicateOperation:(approximateOperation(predicateOperation)),
predicateOperation2?((approximationBits ==
8*typeSizes[type])?predicateOperation2:approximateOperation(predicateOperation2)):"==",
predicateOperation2 == NULL?1:0);
/* printf ("%s\n", sourceCode); */
- program = compileProgram(sourceCode, options);
- free(sourceCode);
+ program =
compileProgram(sourceCodeTemplates[!!inputIsVoidHeaded], options);
+ /* free(sourceCode); */
return program;
}
}
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
