Changeset: 839c37ad6fe5 for MonetDB URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=839c37ad6fe5 Modified Files: monetdb5/extras/bwd/cl_program_utilities.c monetdb5/extras/bwd/operations.c Branch: bwd Log Message:
* fixing bugs for tpch6
Unterschiede (299 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
@@ -27,7 +27,7 @@ cl_program compileProgram(const char* so
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));
- err = clBuildProgram(program, 1, (const cl_device_id[]){getDeviceID()},
"", NULL, NULL);
+ err = clBuildProgram(program, 1, (const cl_device_id[]){getDeviceID()},
options, NULL, NULL);
if(err) {
printf("#%s, clBuildProgram: %s;\n", __func__, clError(err));
size_t bufferSize;
@@ -44,9 +44,9 @@ cl_program getProjectionLeftjoinProgram(
"__global struct{int count; int padding; char values[];}*
outputTail,\n"
"__global struct{int count; int padding; int positions[];}*
inputTail,"
"__global const char* approximation\n) {\n"
- " const int offset =
inputTail->positions[get_global_id(0)]*%1$d;\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)*%1$d+i] =
approximation[offset + i];\n"
+ " outputTail->values[get_global_id(0)*approximationBytes+i]
= approximation[offset + i];\n"
"}";
char options[64];
snprintf(options, 64, "-D approximationBytes=%d",
approximationBits/8);
@@ -78,23 +78,23 @@ cl_program getUSelectProgram(int type, c
"const %1$s operand2\n"
") {\n"
" %1$s value = 0;\n"
- " const size_t inputOffset = get_global_id(0)*%3$d;\n"
- " for(int i = 0; i < %3$d; i++)\n"
- " value += (approximation[inputOffset + i] <<
((i+sizeof(%1$s) - %3$d)*8));\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) - %4$d)*8));\n"
"\n"
" if((value %2$s operand)"
- " && (%4$d || value %2$s operand2)"
+ " && (%5$d || value %3$s operand2)"
" )"
"{\n"
- " const int index =
atomic_inc(&(outputHead->count));"
- " const int offset = index * %3$d;\n"
- " outputHead->positions[index] = get_global_id(0);\n"
- " for(int i = 0; i < %3$d; i++)\n"
- " outputTail->values[offset+i] =
approximation[inputOffset + i];\n"
- " }\n"
- "}";
+ " const int index =
atomic_inc(&(outputHead->count));"
+ " 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"
+ " }\n"
+ "}";
char* sourceCode = malloc(16384);
- snprintf(sourceCode, 16384, sourceCodeTemplate,
(str[]){[TYPE_int] = "int"}[type], approximateOperation(predicateOperation),
approximationBits/8, predicateOperation2 == NULL?1:0);
+ snprintf(sourceCode, 16384, sourceCodeTemplate,
(str[]){[TYPE_int] = "int"}[type], approximateOperation(predicateOperation),
approximateOperation(predicateOperation2), approximationBits/8,
predicateOperation2 == NULL?1:0);
cl_program program = compileProgram(sourceCode,"");
free(sourceCode);
return program;
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
@@ -275,14 +275,10 @@ static inline str uselect(bat *res, bat
const int parameters[2] = {
((*(int*)val)>>batTailResidualBits(data))<<batTailResidualBits(data),
((*(int*)val2)>>batTailResidualBits(data))<<batTailResidualBits(data)};
- for (int i = 0; i < 2; ++i) {
- printf ("setting parameter %d to %d (accurate value:
%d)\n", 3+i, parameters[i], (int[2]){(*(int*)val2),(*(int*)val2)}[i]);
+ for (int i = 0; i < 2; ++i)
if((err = clSetKernelArg(selectKernel, 3+i,
sizeof(int), &(parameters[i])))) // type specific
printf("#%s, clSetKernelArg(%d): %s;\n",
__func__, 3+i, clError(err));
-
- }
- printf ("global work size: %lu\n", BATcount(data));
err = clEnqueueNDRangeKernel(getCommandQueue(), selectKernel,
1, (const size_t[]){0}, (const size_t[]){BATcount(data)}, (const size_t[]){1},
0, NULL, NULL);
if(err) printf("#%s, clEnqueueNDRangeKernel: %s;\n", __func__,
clError(err));
}
@@ -305,57 +301,62 @@ str uselectrefine(bat *res, bat *bid, pt
if(BATcount(data) == 0){
result = BATnew(BAThtype(data), TYPE_void, 0);
} else {
- assert(BATttype(data) == TYPE_int); // type specific
+ assert(ATOMstorage(BATttype(data)) == TYPE_int); // type
specific
BAT* approximation = BATdescriptor(*approx);
- assert(BATttype(data) == TYPE_int); // type specific
+ cl_mem headApproximation = batHeadApproximation(approximation);
+ if(!headApproximation){
+ result = BATnew(BAThtype(data), TYPE_void, 0);
+ } else {
- clHead* compressedHead;
- {
- size_t approximationSize;
- clGetMemObjectInfo(batHeadApproximation(approximation),
CL_MEM_SIZE, sizeof(size_t), &approximationSize, NULL);
- compressedHead = malloc(approximationSize);
- cl_int err = clEnqueueReadBuffer(getCommandQueue(),
batHeadApproximation(approximation), CL_TRUE, 0, approximationSize,
compressedHead , 0, NULL, NULL);
- if(err) printf("#%s, clEnqueueReadBuffer: %s;\n",
__func__, clError(err));
- }
- const size_t candidateCount = compressedHead->count;
- result = BATnew(BAThtype(data), TYPE_void, candidateCount);
+ clHead* compressedHead;
+ {
+ size_t approximationSize;
+ clGetMemObjectInfo(headApproximation,
CL_MEM_SIZE, sizeof(size_t), &approximationSize, NULL);
+ compressedHead = malloc(approximationSize);
+ cl_int err =
clEnqueueReadBuffer(getCommandQueue(), batHeadApproximation(approximation),
CL_TRUE, 0, approximationSize, compressedHead , 0, NULL, NULL);
+ if(err) printf("#%s, clEnqueueReadBuffer:
%s;\n", __func__, clError(err));
+ }
+ const size_t candidateCount = compressedHead->count;
+ result = BATnew(BAThtype(data), TYPE_void,
candidateCount);
- clTail* compressedTail;
- {
- size_t approximationSize;
- clGetMemObjectInfo(batTailApproximation(approximation),
CL_MEM_SIZE, sizeof(size_t), &approximationSize, NULL);
- compressedTail = malloc(approximationSize);
- cl_int err = clEnqueueReadBuffer(getCommandQueue(),
batTailApproximation(approximation), CL_TRUE, 0, approximationSize,
compressedTail , 0, NULL, NULL);
- if(err) printf("#%s, clEnqueueReadBuffer: %s;\n",
__func__, clError(err));
- }
+ clTail* compressedTail;
+ {
+ size_t approximationSize;
+
clGetMemObjectInfo(batTailApproximation(approximation), CL_MEM_SIZE,
sizeof(size_t), &approximationSize, NULL);
+ compressedTail = malloc(approximationSize);
+ cl_int err =
clEnqueueReadBuffer(getCommandQueue(), batTailApproximation(approximation),
CL_TRUE, 0, approximationSize, compressedTail , 0, NULL, NULL);
+ if(err) printf("#%s, clEnqueueReadBuffer:
%s;\n", __func__, clError(err));
+ }
- /* int* resultRegion = (int*) Tloc(result, BUNfirst(result));
// type specific */
- oid* positionRegion = (oid*) Hloc(result, BUNfirst(result)); //
type specific
- const unsigned int approximationMask = ~((1 << (32 -
batTailApproximationBits(approximation)))-1);
- const unsigned char* residuals = batTailResiduals(data);
- const unsigned int residualMask = (1 <<
batTailResidualBits(data))-1;
- const unsigned int residualBytes = batTailResidualBits(data)/8;
- int i = 0, j = 0;
+ /* int* resultRegion = (int*) Tloc(result,
BUNfirst(result)); // type specific */
+ oid* positionRegion = (oid*) Hloc(result,
BUNfirst(result)); // type specific
+ const unsigned int approximationMask = ~((1 << (32 -
batTailApproximationBits(approximation)))-1);
+ const unsigned char* residuals = batTailResiduals(data);
+ const unsigned int residualMask = (1 <<
batTailResidualBits(data))-1;
+ const unsigned int residualBytes =
batTailResidualBits(data)/8;
+ int i = 0, j = 0;
#define refineLoopDoubleOperator(comparator, comparator2)
\
- while(i < candidateCount) {
\
- const int index = compressedHead->positions[i];
\
- const int offset =
(batTailApproximationBits(approximation)/8)*i++; \
- const int compressedValue =
*(int*)&(compressedTail->elements[offset]); \
- const int deCompressedValue =
\
- (compressedValue <<
batTailResidualBits(approximation)) \
- + (*(int*)&residuals[index*residualBytes] &
residualMask); \
- if(deCompressedValue comparator *(int*)val &&
\
- deCompressedValue comparator2 *(int*)val) {
\
- positionRegion[j++] = index;
\
- }
\
- }
+ while(i < candidateCount) {
\
+ const int index = compressedHead->positions[i];
\
+ const int offset =
(batTailApproximationBits(approximation)/8)*i++; \
+ const int compressedValue =
*(int*)&(compressedTail->elements[offset]); \
+ const int deCompressedValue =
\
+ (compressedValue <<
batTailResidualBits(approximation)) \
+ +
(*(int*)&residuals[index*residualBytes] & residualMask); \
+ printf ("compressedValue: %d,
decompressedValue: %d\n", \
+
compressedValue, deCompressedValue);
\
+ if(deCompressedValue comparator *(int*)val &&
\
+ deCompressedValue comparator2
*(int*)val2) {
\
+ positionRegion[j++] = index;
\
+ }
\
+ }
#define refineLoopSingleOperator(comparator)
\
- while(i < candidateCount) {
\
+ while(i < candidateCount) {
\
const int index = compressedHead->positions[i];
\
const int offset =
(batTailApproximationBits(approximation)/8)*i++; \
const int compressedValue =
*(int*)&(compressedTail->elements[offset]); \
@@ -363,59 +364,62 @@ str uselectrefine(bat *res, bat *bid, pt
(compressedValue <<
batTailResidualBits(approximation)) \
+ (*(int*)&residuals[index*residualBytes] &
residualMask); \
if(deCompressedValue comparator *(int*)val) {
\
- positionRegion[j++] = index;
\
- }
\
+ positionRegion[j++] = index;
\
+ }
\
}
-#define refineLoop(comparator, comparator2)
\
- if(comparator2 == NULL) {
\
- refineLoopSingleOperator(comparator);
\
- } else {
\
- switch (comparator2[0]){
\
- case '<':
\
- switch((comparator2)[1]){
\
- case '\0':
\
- refineLoopDoubleOperator(comparator,
<); \
- case '=':
\
- refineLoopDoubleOperator(comparator,
<=); \
- }
\
- case '>':
\
- switch((comparator2)[1]){
\
- case '\0':
\
- refineLoopDoubleOperator(comparator,
>); \
- case '=':
\
- refineLoopDoubleOperator(comparator,
<=); \
- }
\
- case '=':
\
- refineLoopDoubleOperator(comparator, ==);
\
- }
\
- }
+#define refineLoop(comparator, comparator2)
\
+ if(comparator2 == NULL) {
\
+ refineLoopSingleOperator(comparator);
\
+ } else {
\
+ switch (comparator2[0]){
\
+ case '<':
\
+ switch((comparator2)[1]){
\
+ case '\0':
\
+
refineLoopDoubleOperator(comparator, <); \
+ case '=':
\
+
refineLoopDoubleOperator(comparator, <=); \
+ }
\
+ case '>':
\
+ switch((comparator2)[1]){
\
+ case '\0':
\
+
refineLoopDoubleOperator(comparator, >); \
+ case '=':
\
+
refineLoopDoubleOperator(comparator, <=); \
+ }
\
+ case '=':
\
+ refineLoopDoubleOperator(comparator,
==); \
+ }
\
+ }
- switch (OP[0]){
- case '<':
- switch(OP[1]){
- case '\0':
- refineLoop(<, OP2);
- case '=':
- refineLoop(<=, OP2);
+ switch (OP[0]){
+ case '<':
+ switch(OP[1]){
+ case '\0':
+ refineLoop(<, OP2);
+ case '=':
+ refineLoop(<=, OP2);
+ }
+ case '>':
+ switch(OP[1]){
+ case '\0':
+ refineLoop(>, OP2);
+ case '=':
+ refineLoop(>=, OP2);
+ }
+ case '=':
+ refineLoop(==, OP2);
}
- case '>':
- switch(OP[1]){
- case '\0':
- refineLoop(>, OP2);
- case '=':
- refineLoop(>=, OP2);
- }
- case '=':
- refineLoop(==, OP2);
- }
#undef refineLoop
#undef refineLoopSingleOperator
#undef refineLoopDoubleOperator
- BATsetcount(result, j);
- BATseqbase(BATmirror(result), oid_nil);
- free(compressedHead);
- free(compressedTail);
+ BATsetcount(result, j);
+ BATseqbase(BATmirror(result), oid_nil);
+ free(compressedHead);
+ free(compressedTail);
+ }
+ BBPreleaseref(approximation->batCacheid);
+
}
BBPkeepref(*res = result->batCacheid);
@@ -484,6 +488,7 @@ str BWDSemijoinApproximate(int *res, int
slot->approximationBits =
batTailApproximationBits(left);
slot->residuals = NULL;
slot->tailPositions = positionColumn;
+ slot->headApproximation = positionColumn;
slot->tailApproximation =
clCreateBuffer(getCLContext(), CL_MEM_READ_WRITE,
headCount*slot->approximationBits/8+sizeof(clTail), NULL, &err);
if(err) printf("#%s, clCreateBuffer: %s;\n", __func__,
clError(err));
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
