Changeset: af513e091f06 for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=af513e091f06
Modified Files:
monetdb5/extras/bwd/bwd.c
monetdb5/extras/bwd/cl_program_utilities.c
monetdb5/extras/bwd/operations.c
monetdb5/extras/bwd/utilities.c
monetdb5/extras/bwd/utilities.h
Branch: bwd
Log Message:
* got something running that is halfway towards tpch6 (select count(l_shipdate)
as revenue from lineitem where l_shipdate >= date '1994-01-01' and l_shipdate <
date '2013-01-01' and l_discount between 1.0 - 0.01 and 1.0 + 0.01;)
* fixed an insane number of bugs
* unified in-gpu representation of bats (a struct with the first value
being the count)
Unterschiede (gekürzt von 675 auf 300 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
@@ -144,7 +144,9 @@ char* resolveBatToAttribute(bat bid, Cli
}
}
}
- return "unknown";
+ char* buffer = malloc(32);
+ snprintf(buffer, 32, "%d", bid);
+ return buffer;
/* BAT *bn = store_funcs.bind_col(tr, c, 0); */
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
@@ -43,10 +43,15 @@ cl_program getProjectionLeftjoinProgram(
const char* sourceCode = "__kernel void project (\n"
"__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"
+ "__global const struct{int count; int padding; char values[];}*
approximationTail\n) {\n"
+ " __global const char* approximation =
approximationTail->values;"
" const int offset =
inputTail->positions[get_global_id(0)]*approximationBytes;\n"
- " for(int i = 0; i < approximationBytes; i++)\n"
+ " int value = 0;"
+ " 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;"
+ " }\n"
+ /* " printf(\"projected value (%d): %d\\n\",
inputTail->positions[get_global_id(0)], value);" */
"}";
char options[64];
snprintf(options, 64, "-D approximationBytes=%d",
approximationBits/8);
@@ -69,33 +74,35 @@ cl_program getProjectionLeftjoinProgram(
/* } */
cl_program getUSelectProgram(int type, char* predicateOperation, char*
predicateOperation2, unsigned int approximationBits){
- const char* sourceCodeTemplate = "__kernel void uselect
(\n"
- "__global struct{int count; int padding; int
positions[];}* outputHead,\n"
- "__global struct{int count; int padding; char
values[];}* outputTail,\n"
- "__global const char* approximation"
- ",\n"
- "const %1$s operand,\n"
- "const %1$s operand2\n"
- ") {\n"
- " %1$s value = 0;\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)"
- " && (%5$d || value %3$s operand2)"
- " )"
- "{\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),
approximateOperation(predicateOperation2), approximationBits/8,
predicateOperation2 == NULL?1:0);
- cl_program program = compileProgram(sourceCode,"");
- free(sourceCode);
- return program;
+ const char* sourceCodeTemplate = "__kernel void uselect (\n"
+ "__global struct{int count; int padding; int positions[];}*
outputHead,\n"
+ "__global struct{int count; int padding; char values[];}*
outputTail,\n"
+ "__global const struct{int count; int padding; char values[];}*
approximationTail\n"
+ ",\n"
+ "const %1$s operand,\n"
+ "const %1$s operand2\n"
+ ") {\n"
+ " __global const char* approximation =
approximationTail->values;"
+ " %1$s value = 0;\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)"
+ " && (%5$d || value %3$s operand2)"
+ " )"
+ "{\n"
+ " const int index = atomic_inc(&(outputHead->count));"
+ " atomic_inc(&(outputTail->count));" // TODO: this could
probably be done more efficiently
+ " 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),
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
@@ -19,27 +19,27 @@ static const int activateWorkInProgress
#pragma mark Actual MAL Operations Implementation
-static const char* typeNames[] = {[TYPE_int] = "int", [TYPE_void] = "void",
[TYPE_oid] = "oid"};
-
-typedef struct {
- int count;
- int padding;
- char elements[];
-} clTail;
-
-typedef struct {
- int count;
- int padding;
- int positions[];
-} clHead;
+int getCount(cl_mem memoryObject){
+ cl_int err = 0;
+ int result;
+ err = clEnqueueReadBuffer(getCommandQueue(), memoryObject, CL_TRUE, 0,
sizeof(int), &result, 0, NULL, NULL);
+
+ if(err)
+ printf("#%s, clEnqueueReadBuffer: %s;\n", __func__,
clError(err));
+ return result;
+}
clHead* getPositionsColumn(cl_mem memoryObject, clHead* buffer, size_t*
bufferSize){
- cl_int err;
+ cl_int err = 0;
if(!buffer)
clGetMemObjectInfo(memoryObject, CL_MEM_SIZE, sizeof(size_t),
bufferSize, NULL);
- else
- err = clEnqueueReadBuffer(getCommandQueue(), memoryObject,
CL_TRUE, 0, *bufferSize, buffer, 0, NULL, NULL);
- if(err) printf("#%s, clEnqueueReadBuffer: %s;\n", __func__,
clError(err));
+ else{
+ assert(*bufferSize > 0);
+ err = clEnqueueReadBuffer(getCommandQueue(), memoryObject,
CL_TRUE, 0, *bufferSize, buffer, 0, NULL, NULL); }
+
+ if(err){
+ printf("#%s, clEnqueueReadBuffer: %s;\n", __func__,
clError(err));
+ }
return buffer;
}
@@ -62,8 +62,10 @@ str BWDLeftJoinApproximate(bat * res, ba
}
BAT* left = BATdescriptor(*l);
- if(BAThvoid(left) && BATtvoid(left) && BAThvoid(right) &&
right->tseqbase != oid_nil){
- // not sure if this is actually used (or even correct)
+ printf ("%s left bat (%d) is of type [%s%s, %s%s]\n", __func__,
left->batCacheid, typeNames[BAThtype(left)], BAThvoid(left)?" (void)":"",
typeNames[BATttype(left)], BATtvoid(left)?" (void)":"");
+
+ if(0 && BAThvoid(left) && BATtvoid(left) && BAThvoid(right) &&
right->tseqbase != oid_nil){
+ // the more I think about this, the more I feel this case can
never happen
BAT* result = BATnew(ATOMtype(left->htype),
ATOMtype(right->ttype), left->batCount);
BATsetcount(result, left->batCount);
bzero(Tloc(result, BUNfirst(result)),
left->batCount*sizeof(int));
@@ -88,20 +90,17 @@ str BWDLeftJoinApproximate(bat * res, ba
BBPreleaseref(left->batCacheid);
BBPreleaseref(right->batCacheid);
return MAL_SUCCEED;
- } else if(BAThvoid(left) && BAThvoid(right) && right->tseqbase
!= oid_nil){
+ } else if(BAThvoid(left) && BAThvoid(right) && right->tseqbase !=
oid_nil){
cl_mem leftColumn = batTailApproximation(left);
if(!leftColumn) leftColumn =
batHeadApproximation(BATmirror(left));
cl_mem rightColumn = batTailApproximation(right);
if(!rightColumn) rightColumn =
batHeadApproximation(BATmirror(right));
cl_program program =
getProjectionLeftjoinProgram(batTailApproximationBits(right));
-
-
cl_int err;
cl_kernel projectKernel = clCreateKernel(program, "project",
&err);
if(err) printf("#%s, clCreateKernel: %s;\n", __func__,
clError(err));
-
int headCount; // :-)
if((err = clEnqueueReadBuffer(getCommandQueue(), leftColumn,
CL_TRUE, 0, sizeof(int), &headCount , 0, NULL, NULL)))
printf("#%s, clEnqueueReadBuffer: %s;\n", __func__,
clError(err));
@@ -109,12 +108,13 @@ str BWDLeftJoinApproximate(bat * res, ba
BAT* result = BATnew(TYPE_void, ATOMtype(right->ttype), 0);
{
const unsigned int newIndex =
getNextFreeDecomposedBATSlotIndex();
- DecomposedBATSlot* slot =
getDecomposedBATSlot(newIndex);
+ DecomposedBATSlot* slot =
getDecomposedBATSlotForIndex(newIndex);
BATsetprop(result, batRegistryIndex, TYPE_int,
(int[]){newIndex});
slot->approximationBits =
batTailApproximationBits(right);
slot->residuals = NULL;
slot->tailPositions = leftColumn;
slot->tailApproximation =
clCreateBuffer(getCLContext(), CL_MEM_READ_WRITE,
headCount*slot->approximationBits/8+sizeof(clTail), NULL, &err);
+ assert(slot->tailPositions);
if(err) printf("#%s, clCreateBuffer: %s;\n", __func__,
clError(err));
clEnqueueWriteBuffer(getCommandQueue(),
slot->tailApproximation, CL_TRUE, 0, sizeof(int), &headCount, 0, NULL, NULL);
// I wonder what is faster transfering a single integer to the GPU or running a
kernel that initializes a value
@@ -162,13 +162,15 @@ str BWDLeftJoinRefine(bat * res, bat * l
return ALGleftjoin(res, l, r);
}
BAT* left = BATdescriptor(*l);
+ printf ("%s left bat (%d) is of type [%s%s, %s%s]\n", __func__,
left->batCacheid, typeNames[BAThtype(left)], BAThvoid(left)?" (void)":"",
typeNames[BATttype(left)], BATtvoid(left)?" (void)":"");
BAT* right = BATdescriptor(*r);
BAT* refinement;
BAT* approximation = BATdescriptor(*approx);
const register size_t approximationBits =
batTailApproximationBits(right);
/* BATnew(approximation->htype, approximation->ttype,
approximation->batCount); */
- if(BAThvoid(left) && BATtvoid(left) && BAThvoid(right) &&
right->tseqbase != oid_nil){
+ if(0 && BAThvoid(left) && BATtvoid(left) && BAThvoid(right) &&
right->tseqbase != oid_nil){
+ // not very likely optimization, hence disabled
if(!batTailIsDecomposed(right))
throw (MAL, "bwd.BWDLeftJoinApproximate", "bat is not
decomposed: %d", *r);
refinement = BATnew(ATOMtype(left->htype),
ATOMtype(right->ttype), left->batCount);
@@ -236,7 +238,7 @@ static inline void superverboseprintf(co
static inline str uselect(bat *res, bat *bid, ptr val, char *OP, ptr val2,
char *OP2, Client cntxt){
- BAT* data = BATdescriptor(*bid);
+BAT* data = BATdescriptor(*bid);
if(data == NULL)
return MAL_SUCCEED;
assert(ATOMstorage(BATttype(data)) == TYPE_int); // type specific
@@ -244,23 +246,28 @@ static inline str uselect(bat *res, bat
throw (MAL, "bwd problem", "while running %s, I noticed that
bat is not decomposed: %s", __func__, resolveBatToAttribute(data->batCacheid,
cntxt));
BAT* result = BATnew(BAThtype(data), TYPE_void, data->batCount);
//TODO: not sure if we actually need to allocate space here (count could be 0)
- if(data->batCount){
+ size_t dataCount = getCount(batTailApproximation(data));
+
+ if(dataCount){
{
const unsigned int newIndex =
getNextFreeDecomposedBATSlotIndex();
- DecomposedBATSlot* slot =
getDecomposedBATSlot(newIndex);
+ DecomposedBATSlot* slot =
getDecomposedBATSlotForIndex(newIndex);
BATsetprop(result, batRegistryIndex, TYPE_int,
(int[]){newIndex});
slot->approximationBits =
batTailApproximationBits(data);
slot->residuals = NULL;
cl_int err;
slot->tailApproximation =
clCreateBuffer(getCLContext(), CL_MEM_READ_WRITE,
BATcount(data)*slot->approximationBits/8+sizeof(clTail), NULL, &err);
- slot->headApproximation =
clCreateBuffer(getCLContext(), CL_MEM_READ_WRITE,
BATcount(data)*sizeof(int)+sizeof(clHead), NULL, &err);
+ const int size = dataCount*sizeof(int)+sizeof(clHead);
+ slot->headApproximation =
clCreateBuffer(getCLContext(), CL_MEM_READ_WRITE, size, NULL, &err);
slot->tailPositions = slot->headApproximation;
+ assert(slot->tailPositions);
+
clEnqueueWriteBuffer(getCommandQueue(),
slot->headApproximation, CL_TRUE, 0, sizeof(int), (int[1]){}, 0, NULL, NULL);
// I wonder what is faster transfering a single integer to the GPU or running a
kernel that initializes a value
if(err) printf("#%s, clCreateBuffer: %s;\n", __func__,
clError(err));
}
- cl_program program = getUSelectProgram(BATttype(data), OP, OP2,
batTailApproximationBits(data));
+ cl_program program =
getUSelectProgram(ATOMstorage(BATttype(data)), OP, OP2?OP2:"==",
batTailApproximationBits(data));
int err;
@@ -279,9 +286,10 @@ static inline str uselect(bat *res, bat
if((err = clSetKernelArg(selectKernel, 3+i,
sizeof(int), &(parameters[i])))) // type specific
printf("#%s, clSetKernelArg(%d): %s;\n",
__func__, 3+i, clError(err));
- err = clEnqueueNDRangeKernel(getCommandQueue(), selectKernel,
1, (const size_t[]){0}, (const size_t[]){BATcount(data)}, (const size_t[]){1},
0, NULL, NULL);
+ err = clEnqueueNDRangeKernel(getCommandQueue(), selectKernel,
1, (const size_t[]){0}, (const size_t[]){dataCount}, (const size_t[]){1}, 0,
NULL, NULL);
if(err) printf("#%s, clEnqueueNDRangeKernel: %s;\n", __func__,
clError(err));
}
+ printf ("%s result bat (%d) is of type [%s%s, %s%s]\n", __func__,
result->batCacheid, typeNames[BAThtype(result)], BAThvoid(result)?" (void)":"",
typeNames[BATttype(result)], BATtvoid(result)?" (void)":"");
BBPkeepref(*res = result->batCacheid);
BBPreleaseref(*bid);
return MAL_SUCCEED;
@@ -325,7 +333,8 @@ str uselectrefine(bat *res, bat *bid, pt
clTail* compressedTail;
{
size_t approximationSize;
-
clGetMemObjectInfo(batTailApproximation(approximation), CL_MEM_SIZE,
sizeof(size_t), &approximationSize, NULL);
+
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));
@@ -340,79 +349,79 @@ str uselectrefine(bat *res, bat *bid, pt
const unsigned int residualBytes =
batTailResidualBits(data)/8;
int i = 0, j = 0;
+
+ if(1){
#define refineLoopDoubleOperator(comparator, comparator2)
\
- 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]); \
+ const int deCompressedValue =
\
+ (compressedValue <<
batTailResidualBits(approximation)) \
+ +
(*(int*)&residuals[index*residualBytes] & residualMask); \
+ if(deCompressedValue comparator
*(int*)val &&
\
+ deCompressedValue comparator2
*(int*)val2) { \
+ positionRegion[j++] = index;
\
+ }
\
+ }
+#define refineLoopSingleOperator(comparator)
\
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
