Changeset: bb0fe6290853 for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=bb0fe6290853
Modified Files:
monetdb5/extras/bwd/cl_program_utilities.c
monetdb5/extras/bwd/operations.c
monetdb5/extras/bwd/utilities.h
Branch: bwd
Log Message:
* fixed a minor bug with result set space allocation for the approximation
Unterschiede (89 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
@@ -51,7 +51,7 @@ cl_program getProjectionLeftjoinProgram(
" 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);" */
+ /* " printf(\"projected value (%d): %d\\n\",
inputTail->positions[get_global_id(0)], value<<(8*(4-approximationBytes)));" */
"}";
char options[64];
snprintf(options, 64, "-D approximationBytes=%d",
approximationBits/8);
@@ -94,10 +94,13 @@ cl_program getUSelectProgram(int type, c
"{\n"
" const int index = atomic_inc(&(outputHead->count));"
" atomic_inc(&(outputTail->count));" // 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" */
" const int offset = index * %4$d;\n"
" outputHead->positions[index] = get_global_id(0);\n"
- " for(int i = 0; i < %4$d; i++)\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"
" }\n"
"}";
char* sourceCode = malloc(16384);
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
@@ -238,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
@@ -256,7 +256,8 @@ BAT* data = BATdescriptor(*bid);
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->tailApproximation =
clCreateBuffer(getCLContext(), CL_MEM_READ_WRITE,
dataCount*slot->approximationBits/8+sizeof(clTail), NULL, &err);
+ if(0) printf ("%s result tail approximation: %p\n",
__func__, slot->tailApproximation);
const int size = dataCount*sizeof(int)+sizeof(clHead);
slot->headApproximation =
clCreateBuffer(getCLContext(), CL_MEM_READ_WRITE, size, NULL, &err);
slot->tailPositions = slot->headApproximation;
@@ -288,6 +289,18 @@ BAT* data = BATdescriptor(*bid);
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));
+ if(0){
+ clFinish(getCommandQueue());
+ clTail* compressedTail;
+ size_t resultSize;
+ printf ("%s result tail result: %p\n",
__func__, batTailApproximation(result));
+
+
clGetMemObjectInfo(batTailApproximation(result), CL_MEM_SIZE, sizeof(size_t),
&resultSize
+
, NULL);
+ compressedTail = malloc(resultSize);
+ cl_int err =
clEnqueueReadBuffer(getCommandQueue(), batTailApproximation(result), CL_TRUE,
0, resultSize, compressedTail , 0, NULL, NULL);
+ if(err) printf("#%s, clEnqueueReadBuffer:
%s;\n", __func__, clError(err));
+ }
}
if(0) 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);
@@ -333,6 +346,8 @@ str uselectrefine(bat *res, bat *bid, pt
clTail* compressedTail;
{
size_t approximationSize;
+ if(0) printf ("%s approximation tail
approximation: %p\n", __func__, batTailApproximation(approximation));
+
clGetMemObjectInfo(batTailApproximation(approximation), CL_MEM_SIZE,
sizeof(size_t), &approximationSize
, NULL);
compressedTail = malloc(approximationSize);
diff --git a/monetdb5/extras/bwd/utilities.h b/monetdb5/extras/bwd/utilities.h
--- a/monetdb5/extras/bwd/utilities.h
+++ b/monetdb5/extras/bwd/utilities.h
@@ -14,7 +14,7 @@ typedef struct {
typedef struct {
int count;
int padding;
- char elements[];
+ unsigned char elements[];
} clTail;
typedef struct {
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
