Changeset: 0fd62512249d for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=0fd62512249d
Modified Files:
monetdb5/extras/bwd/bwd.c
monetdb5/extras/bwd/cl_program_utilities.c
monetdb5/extras/bwd/operations.c
Branch: bwd
Log Message:
* smaller fixes for correctness
Unterschiede (180 Zeilen):
diff --git a/monetdb5/extras/bwd/bwd.c b/monetdb5/extras/bwd/bwd.c
--- a/monetdb5/extras/bwd/bwd.c
+++ b/monetdb5/extras/bwd/bwd.c
@@ -102,15 +102,17 @@ str bwdecompose(bat * res, bat * subject
BAT* subject = BATdescriptor(*subjectBAT);
if (VIEWtparent(subject))
subject = BATdescriptor(abs(VIEWtparent(subject)));
+ BAT* result = BATnew(TYPE_void, TYPE_str, 0);
+ char buffer[4096];
if(! batTailIsDecomposed(subject)){
printf ("decomposing %s into %d and %d bits\n",
resolveBatToAttribute(subject->batCacheid, cntxt), approximationBits[0],
Tsize(subject)*8-*approximationBits);
BATsetprop(subject, batRegistryIndex, TYPE_int,
(int[]){decomposeIntArray((int*)Tloc(subject, BUNfirst(subject)),
subject->batCount, *approximationBits)});
- BAT* result = BATnew(TYPE_void, TYPE_str, 0);
- char buffer[4096];
snprintf(buffer, 4096, "successfully decomposed bat %d",
subject->batCacheid);
- BUNappend(result, buffer, 0);
- BBPkeepref(*res = result->batCacheid);
+ } else {
+ snprintf(buffer, 4096, "bat %d already decomposed",
subject->batCacheid);
}
+ BUNappend(result, buffer, 0);
+ BBPkeepref(*res = result->batCacheid);
BBPreleaseref(subject->batCacheid);
return MAL_SUCCEED;
}
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
@@ -24,6 +24,8 @@ char* approximateOperation(char* exactOp
static inline void superverboseprintf(const char * format, ... ){};
cl_program compileProgram(const char* sourceCode, char* options){
+ struct timespec before, after;
+ clock_gettime(CLOCK_THREAD_CPUTIME_ID, &before);
cl_int err;
cl_program program = clCreateProgramWithSource(getCLContext(), 1,
(const char*[]){sourceCode}, (size_t[]){strlen(sourceCode)}, &err);
if(err) THRprintf(GDKout, "#%s, clCreateProgramWithSource: %s;\n",
__func__, clError(err));
@@ -36,6 +38,9 @@ cl_program compileProgram(const char* so
clGetProgramBuildInfo(program, getDeviceID(),
CL_PROGRAM_BUILD_LOG, bufferSize, (buffer = alloca(bufferSize)), NULL);
printf("#%s, clBuildProgram log: %s;\n", __func__, buffer);
}
+ clock_gettime(CLOCK_THREAD_CPUTIME_ID, &after);
+ printf ("compilation took %ld
nanoseconds\n",(after.tv_sec*1000000000+after.tv_nsec)-(before.tv_sec*1000000000+before.tv_nsec));
+
return program;
}
@@ -58,24 +63,41 @@ cl_program getMultiplyProgram(unsigned i
}
cl_program getProjectionLeftjoinProgram(unsigned int approximationBits,
unsigned int offsetBits){
- const char* sourceCode = "__kernel void project (\n"
- "__global struct{int count; int base; unsigned char values[];}*
outputTail,\n"
- "__global struct{int count; int padding; int positions[];}*
inputTail,"
- "__global const struct{int count; int base; unsigned char
values[];}* approximationTail\n) {\n"
- " if(get_global_id(0) < inputTail->count){\n"
- " __global const unsigned char* approximation =
approximationTail->values;"
- " const int offset =
inputTail->positions[get_global_id(0)]*approximationBytes;\n"
- /* " int value = approximationTail->base;" */
- " for(int i = 0; i < approximationBytes; i++){\n"
- " outputTail->values[get_global_id(0)*approximationBytes+i]
= approximation[offset + i];\n"
- /* " value += approximation[offset + i] << 8*(i+1);" */
- " }\n"
- /* " 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"
- "}";
+ static cl_program programCache[4] = {};
+
+ if(programCache[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"
+ " 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"
+ /* " value += approximation[offset + i] << 8*(i+1);"
*/
+ /* " 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);
- return compileProgram(sourceCode, options);
+ programCache[approximationBits/8-offsetBits/8] =
compileProgram(sourceCode, options);
+ }
+ return programCache[approximationBits/8-offsetBits/8];
}
cl_program getUSelectProgram(int type, char* predicateOperation, char*
predicateOperation2, unsigned int approximationBits, unsigned int offsetBits,
char inputIsVoidHeaded){
@@ -108,7 +130,7 @@ cl_program getUSelectProgram(int type, c
"{\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"
+ /* " 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"
@@ -123,34 +145,41 @@ cl_program getUSelectProgram(int type, c
"}",
[1] = "__kernel void uselect (\n"
"__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"
- ",\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"
- " __global const unsigned char* approximation =
approximationTail->values;"
+ /* " __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"
+ "#else\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"
+ " value += (approximationTail->values[inputOffset + i] <<
((i+sizeof(%1$s)-%6$d - %4$d)*8));\n"
+ "#endif\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"
+ " ){\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\\n\", value, get_global_id(0), index);\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"
- " outputHead->positions[index] = 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]);" */
+ " outputTail->values[offset+i] =
approximationTail->values[inputOffset + i];\n"
" }\n"
+ "#endif\n"
" }\n"
" }\n"
"}"
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
@@ -468,7 +468,7 @@ str uselectrefine(bat *res, bat *bid, pt
const int compressedValue =
*(int*)&(compressedTail->elements[offset]) & tailApproximationMask; \
const int deCompressedValue =
compressedTail->base+ \
(compressedValue <<
tailResidualBits)
\
- +
(*(int*)&residuals[index*residualBytes] & residualMask); \
+ +
(*(int*)&residuals[i*residualBytes] & residualMask); \
{
\
positionRegion[j] =
index;
\
j+= ((deCompressedValue
comparator value1) & \
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
