Changeset: bc8ca326a689 for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=bc8ca326a689
Modified Files:
monetdb5/extras/bwd/cl_program_utilities.c
monetdb5/extras/bwd/operations.c
monetdb5/extras/bwd/utilities.c
Branch: bwd
Log Message:
* working on the sinterklaas case now
* introduced a max on intermediate result size (this is a hack and should be
generalized by, e.g., processing & transfering in chunks)
Unterschiede (237 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
@@ -120,16 +120,17 @@ 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"
" outputTail->values[offset+i] = approximation[inputOffset
+ i];\n"
/* " printf(\"set byte %%d to %%d\\n\",
offset+i,outputTail->values[offset+i]);" */
" }\n"
- " } else {\n"
- " printf(\"not selected value %%d (base: %%d, outbase: %%d)
from slot %%d, head value: %%d\\n\", value, approximationTail->base,
outputTail->base, get_global_id(0),
approximationHead->positions[get_global_id(0)]);\n"
- " }"
+ " } \n"
+ /* "else {\n" */
+ /* " printf(\"not selected value %%d (base: %%d, outbase:
%%d) from slot %%d, head value: %%d\\n\", value, approximationTail->base,
outputTail->base, get_global_id(0),
approximationHead->positions[get_global_id(0)]);\n" */
+ /* " }" */
" }\n"
"}",
[1] = "__kernel void uselect (\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
@@ -19,7 +19,9 @@
static const int activateWorkInProgress = 1;
static const int synchronousGPU = 1;
-#define WORK_GROUP_SIZE 1
+static const int eagerBufferLoading = 1;
+#define WORK_GROUP_SIZE 16
+#define MAX_INTERMEDIATE_RESULT_SIZE 16777216
#pragma mark Actual MAL Operations Implementation
@@ -33,10 +35,29 @@ int getCount(cl_mem memoryObject){
return result;
}
+clTail getTailHeader(cl_mem memoryObject){
+ cl_int err = 0;
+ clTail result;
+ err = clEnqueueReadBuffer(getCommandQueue(), memoryObject, CL_TRUE, 0,
sizeof(clHead), &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 = 0;
- if(!buffer)
- clGetMemObjectInfo(memoryObject, CL_MEM_SIZE, sizeof(size_t),
bufferSize, NULL);
+ if(!buffer){
+ if(eagerBufferLoading){
+ clGetMemObjectInfo(memoryObject, CL_MEM_SIZE,
sizeof(size_t), bufferSize, NULL);
+ } else {
+ int count;
+ err = clEnqueueReadBuffer(getCommandQueue(),
memoryObject, CL_TRUE, 0, sizeof(int), &count, 0, NULL, NULL);
+ printf ("positioncount: %ld\n", count);
+ *bufferSize = sizeof(int) * count;
+
+ }
+ }
else{
assert(*bufferSize > 0);
err = clEnqueueReadBuffer(getCommandQueue(), memoryObject,
CL_TRUE, 0, *bufferSize, buffer, 0, NULL, NULL); }
@@ -48,7 +69,7 @@ clHead* getPositionsColumn(cl_mem memory
}
clTail* getApproximateValuesColumn(cl_mem memoryObject, clTail* buffer,
size_t* bufferSize){
- cl_int err;
+ cl_int err = 0;
if(!buffer)
clGetMemObjectInfo(memoryObject, CL_MEM_SIZE, sizeof(size_t),
bufferSize, NULL);
else
@@ -259,7 +280,8 @@ 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)
- size_t dataCount = getCount(batTailApproximation(data));
+ const clTail dataHeader = getTailHeader(batTailApproximation(data));
+ const int dataCount = dataHeader.count;
if(dataCount){
{
@@ -270,15 +292,21 @@ static inline str uselect(bat *res, bat
slot->tailOffsetBits = batTailOffsetBits(data);
slot->residuals = NULL;
cl_int err;
- slot->tailApproximation =
clCreateBuffer(getCLContext(), CL_MEM_READ_WRITE,
dataCount*slot->approximationBits/8+sizeof(clTail), NULL, &err);
+ slot->tailApproximation =
clCreateBuffer(getCLContext(), CL_MEM_READ_WRITE,
MIN(dataCount,MAX_INTERMEDIATE_RESULT_SIZE)*slot->approximationBits/8+sizeof(clTail),
NULL, &err);
+ if(err) printf("#%s, clCreateBuffer: %s;\n", __func__,
clError(err));
if(0) printf ("%s result tail approximation: %p\n",
__func__, slot->tailApproximation);
- const int size = dataCount*sizeof(int)+sizeof(clHead);
+ const int size =
MIN(dataCount,MAX_INTERMEDIATE_RESULT_SIZE)*sizeof(int)+sizeof(clHead);
slot->headApproximation =
clCreateBuffer(getCLContext(), CL_MEM_READ_WRITE, size, NULL, &err);
+ if(err) printf("#%s, clCreateBuffer: %s;\n", __func__,
clError(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));
+ { //propagate base
+ if((err =
clEnqueueWriteBuffer(getCommandQueue(), slot->tailApproximation, CL_TRUE, 0,
sizeof(clTail), (clTail[1]){{.count = 0, .base=dataHeader.base}}, 0, NULL,
NULL)))
+ printf("#%s, clEnqueueWriteBuffer:
%s;\n", __func__, clError(err));
+ }
+
}
@@ -310,7 +338,8 @@ static inline str uselect(bat *res, bat
printf("#%s, clSetKernelArg(%d): %s;\n",
__func__, bufferI+i, clError(err));
err = clEnqueueNDRangeKernel(getCommandQueue(), selectKernel,
1, (const size_t[]){0}, (const
size_t[]){ceil(dataCount/((float)WORK_GROUP_SIZE))*WORK_GROUP_SIZE}, (const
size_t[]){WORK_GROUP_SIZE}, 0, NULL, NULL);
- if(err) printf("#%s, clEnqueueNDRangeKernel: %s;\n", __func__,
clError(err));
+ if(err)
+ printf("#%s, clEnqueueNDRangeKernel: %s;\n", __func__,
clError(err));
if (synchronousGPU) clFinish(getCommandQueue());
if(0){
clTail* compressedTail;
@@ -367,9 +396,12 @@ str uselectrefine(bat *res, bat *bid, pt
if(batTailResidualBits(data) == 0){ // if we don't have
any CPU-resident residual, we know that the approximation is accurate, no need
for refinement
const int* headPositions =
compressedHead->positions;
+ {
int i;
+#pragma omp parallel for
for (i = 0; i < candidateCount; ++i)
positionRegion[i] = headPositions[i];
+ }
BATsetcount(result, candidateCount);
BATseqbase(BATmirror(result), oid_nil);
@@ -622,44 +654,56 @@ str BWDSemijoinRefine(int *res, int *lid
refinement = BATnew(TYPE_void, BATttype(left), right->batCount);
BATseqbase(refinement, oid_nil);
+ struct timespec before, after;
+ clock_gettime(CLOCK_THREAD_CPUTIME_ID, &before);
size_t bufferSize;
getPositionsColumn(batTailPositions(approximation), NULL,
&bufferSize);
clHead* supersetPositionsColumn =
getPositionsColumn(batTailPositions(approximation), malloc(bufferSize),
&bufferSize);
unsigned char* filteredResiduals = NULL;
+
+ clock_gettime(CLOCK_THREAD_CPUTIME_ID, &after);
+ printf ("clHead* supersetPositionsColumn =
getPositionsColumn(batTailPositions(approximation), malloc(bufferSize),
&bufferSize); took %ld
nanoseconds\n",(after.tv_sec*1000000000+after.tv_nsec)-(before.tv_sec*1000000000+before.tv_nsec));
+
+
getApproximateValuesColumn(batTailApproximation(approximation),
NULL, &bufferSize);
+
clTail* supersetApproximateValuesColumn =
getApproximateValuesColumn(batTailApproximation(approximation),
malloc(bufferSize), &bufferSize);
- if(0 && supersetPositionsColumn->count == BATcount(right)){ //
the approximation was correct, no refinement necessary
- /* clEnqueueReadBuffer(getCommandQueue(),
batTailApproximation(approximation), CL_TRUE, sizeof(clTail),
BATcount(right)*sizeof(int), Tloc(refinement, BUNfirst(refinement)), 0, NULL,
NULL); */
-
- BATsetcount(refinement, supersetPositionsColumn->count);
- } else {
-
-
- size_t refinementCount = 0;
+ {
if(supersetPositionsColumn){
{ // if the approximation is empty, the memory
objects are NULL
int* refinementRegion = (int*)
Tloc(refinement, BUNfirst(refinement));
const unsigned char* residuals =
batTailResiduals(left);
const unsigned int residualBytes =
Tsize(left)-approximationBits/8;
+ if(residualBytes > 0){
+ filteredResiduals =
calloc((supersetPositionsColumn->count)*residualBytes + sizeof(size_t), 1);
+
+ int i,j;
+ size_t refinementCount = 0;
- filteredResiduals =
calloc((supersetPositionsColumn->count)*residualBytes + sizeof(size_t), 1);
- const size_t leftHMin = left->hseqbase,
leftHMax = left->hseqbase + BATcount(left);
-
- int i,j;
- for ( i = j = 0; j < BATcount(right);
++j) {
-
while(supersetPositionsColumn->positions[i] != ((oid*)right->H->heap.base)[j])
- i++;
-
refinementRegion[refinementCount++] = decompressIntValue(i, approximationBits,
offsetBits, supersetApproximateValuesColumn, residuals,
supersetPositionsColumn->positions[i]);
+ for ( i = j = 0; j <
BATcount(right); ++j) {
+
while(supersetPositionsColumn->positions[i] != ((oid*)right->H->heap.base)[j])
+ i++;
+
refinementRegion[refinementCount++] = decompressIntValue(i, approximationBits,
offsetBits, supersetApproximateValuesColumn, residuals,
supersetPositionsColumn->positions[i]);
+ }
+ BATsetcount(refinement, refinementCount);
+
+ } else {
+ const int count =
BATcount(right);
+ memcpy(refinementRegion,
supersetApproximateValuesColumn->elements, count*sizeof(int));
+/* int i; */
+/* #pragma omp parallel for */
+/* for (i = 0; i < count; ++i) { */
+/* refinementRegion[i] =
((int*)supersetApproximateValuesColumn->elements)[i]; */
+/* } */
+ BATsetcount(refinement, count);
+
}
}
}else {
printf ("uh oh, %s ran into an unimplemented
case: no approximation provided\n", __func__);
}
-
- BATsetcount(refinement, refinementCount);
-
}
BATseqbase(BATmirror(refinement), left->tseqbase);
diff --git a/monetdb5/extras/bwd/utilities.c b/monetdb5/extras/bwd/utilities.c
--- a/monetdb5/extras/bwd/utilities.c
+++ b/monetdb5/extras/bwd/utilities.c
@@ -243,7 +243,10 @@ const bounds findBounds(const int* subje
}
-
+char* humanreadablesize(unsigned int value, char* buffer, size_t bufferSize){
+ snprintf(buffer, bufferSize, "%d %s",
(value<1024)?value:((value<(1024*1024))?(value/1024):(value/(1024*1024))),
(value<1024)?"byte":((value<(1024*1024)?"kb":"Mb")));
+ return buffer;
+}
const unsigned int decomposeIntArray(const int* subject, const size_t size,
const size_t approximationBits){
@@ -266,6 +269,7 @@ const unsigned int decomposeIntArray(con
const register unsigned int residualBytes = (32-approximationBits)/8;
const register unsigned int approximationBytes =
(slot->approximationBits/8)-(slot->tailOffsetBits/8);
+ printf ("using %s for approximation and %s for residuals \n",
humanreadablesize(approximationBytes*(size+1)+ sizeof(clTail), (char[64]){},
64), humanreadablesize(residualBytes*(size+1), (char[]){64}, 64));
slot->residuals = calloc((size+1)*residualBytes, sizeof(char));
clTail* approximation = calloc((size+1)*approximationBytes+
sizeof(clTail), sizeof(char));
approximation->base = subjectBounds.min;
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
