Changeset: 40cd4922a0de for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=40cd4922a0de
Modified Files:
monetdb5/extras/bwd/cl_program_utilities.c
Branch: bwd
Log Message:
* still working on bitwise approximate processing
Unterschiede (59 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
@@ -14,6 +14,8 @@
#include <CL/cl.h>
#endif
+#define GPU_PRINTF(a) ""
+
#include "utilities.h"
char* approximateOperation(char* exactOperation){
@@ -158,7 +160,7 @@ cl_program getUSelectProgram(int type, c
" 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"
- " printf(\"===== %%d (global id: %%d)\\n\", value,
get_global_id(0));"
+ 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)"
@@ -178,7 +180,8 @@ cl_program getUSelectProgram(int type, c
" }\n"
"}",
- [2] = " __constant static const size_t targetTypeBits
= (sizeof(targetType)*8);"
+ [2] = " __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"
@@ -192,10 +195,14 @@ cl_program getUSelectProgram(int type, c
" size_t offset =
(get_global_id(0)*approximationBits)%%targetTypeBits;\n"
" {\n"
" __global const unsigned int* vals =
approximationTail->values;\n"
- " 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))<<residualBits);\n"
- " 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"
- " +
((offset+approximationBits)<=targetTypeBits)*(vals[slot]>>(targetTypeBits-offset-approximationBits))));\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"
+ /* " 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"
"\n"
" if((value %2$s operand)"
@@ -220,7 +227,7 @@ cl_program getUSelectProgram(int type, c
char* sourceCode = malloc(16384);
const str typeNames[] = {[TYPE_int] = "int"};
const size_t typeSizes[] = {[TYPE_int] = sizeof(int)};
- 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)*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);
{
char options[256];
cl_program program;
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
