Changeset: 841571026bad for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=841571026bad
Modified Files:
monetdb5/extras/bwd/cl_program_utilities.c
monetdb5/extras/bwd/operations.c
monetdb5/extras/bwd/utilities.c
Branch: bwd
Log Message:
* fixed multi-column grouping and improved buffer initialization
Unterschiede (204 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
@@ -56,8 +56,10 @@ cl_program compileProgram(const char* so
cl_program getZeroOutProgram(void){
const char* sourceCode = "__kernel void zeroOut (\n"
- "__global int* buffer\n){\n"
- " buffer[get_global_id(0)] = 0;"
+ "__global int* buffer,\n"
+ "const unsigned int offset)\n"
+ "{\n"
+ " buffer[get_global_id(0)+offset] = 0;"
"}";
return compileProgram(sourceCode, "");
}
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
@@ -137,7 +137,7 @@ cl_int clEnqueueFillBuffer(cl_command_qu
cl_uint num_events_in_wait_list ,
const cl_event * event_wait_list ,
cl_event * event ) {
- if(pattern_size == 4 && ((int*)pattern)[0] == 0){
+ if(pattern_size == 4 && ((int*)pattern)[0] == 0 && (offset == 0 ||
offset == 8)){
cl_int err = 0;
if(0){
int* tmpBuffer = GDKzalloc(size);
@@ -148,6 +148,7 @@ cl_int clEnqueueFillBuffer(cl_command_qu
} else {
cl_kernel kernel = clCreateKernel(getZeroOutProgram(),
"zeroOut", &err);
clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer);
+ clSetKernelArg(kernel, 1, sizeof(int),
(int[]){offset/4});
err = bwdEnqueueNDRangeKernel(getCommandQueue(),
kernel, 1, (const size_t[]){0}, (size_t[]){ ceil(size/128.0)*32}, (const
size_t[]){32}, num_events_in_wait_list, event_wait_list, event);
if(err) printf("#%s, bwdEnqueueNDRangeKernel: %s;\n",
__func__, clError(err));
err= clFinish(getCommandQueue());
@@ -200,8 +201,8 @@ str BWDLeftJoinApproximate(bat * res, ba
if((err = clEnqueueReadBuffer(getCommandQueue(), leftColumn,
CL_TRUE, 0, sizeof(int), &headCount , 0, NULL, NULL)))
printf("#%s, clEnqueueReadBuffer (left): %s;\n",
__func__, clError(err));
+ newTailDefinition = (clTail){.count = headCount}; // :-)
- newTailDefinition = (clTail){.count = headCount}; // :-)
if((err = clEnqueueReadBuffer(getCommandQueue(), rightColumn,
CL_TRUE, sizeof(int), sizeof(int), &newTailDefinition.base , 0, NULL, NULL)))
printf("#%s, clEnqueueReadBuffer (right): %s;\n",
__func__, clError(err));
@@ -220,11 +221,12 @@ str BWDLeftJoinApproximate(bat * res, ba
slot->tailApproximation =
bwdClCreateBuffer(getCLContext(), CL_MEM_READ_WRITE,
calculatedBufferSize(headCount,slot->approximationBits)+sizeof(clTail), NULL,
&err);
if(err) printf("#%s, bwdClCreateBuffer: %s;\n",
__func__, clError(err));
+ err = clEnqueueFillBuffer(getCommandQueue(),
slot->tailApproximation, zeroIntPattern, sizeof(int), sizeof(clTail),
calculatedBufferSize(headCount,slot->approximationBits), 0, NULL, NULL);
+ if(err) printf("#%s, clEnqueueFillBuffer: %s;\n",
__func__, clError(err));
+
err = clEnqueueWriteBuffer(getCommandQueue(),
slot->tailApproximation, CL_TRUE, 0, sizeof(clTail), &newTailDefinition, 0,
NULL, NULL);
if(err) printf("#%s, clEnqueueWriteBuffer: %s;\n",
__func__, clError(err));
- err = clEnqueueFillBuffer(getCommandQueue(),
slot->tailApproximation, zeroIntPattern, sizeof(int), sizeof(clTail),
calculatedBufferSize(headCount,slot->approximationBits), 0, NULL, NULL);
- if(err) printf("#%s, clEnqueueFillBuffer: %s;\n",
__func__, clError(err));
}
for (i = 0; i < 3; ++i) {
if((err = clSetKernelArg(projectKernel, i,
sizeof(cl_mem), &((cl_mem[]){
@@ -559,9 +561,9 @@ static inline str uselect(bat *res, bat
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
{ //propagate base
+ clEnqueueFillBuffer(getCommandQueue(),
slot->tailApproximation, zeroIntPattern, sizeof(int), sizeof(clTail),
tailBufferElementSize, 0, NULL, NULL);
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));
- clEnqueueFillBuffer(getCommandQueue(),
slot->tailApproximation, zeroIntPattern, sizeof(int), sizeof(clTail),
tailBufferElementSize, 0, NULL, NULL);
}
@@ -638,7 +640,6 @@ static inline unsigned int refinementLoo
const int tailApproximationMask = (1<<tailApproximationBits)-1;
const unsigned int residualMask = (1 << tailResidualBits)-1;
unsigned int i=0, j=0, inputPositionsI = 0;
-
if(1){/* use cpp | egrep -v "^#" | indent | pbcopy to expand */
const unsigned int targetTypeBits = sizeof(int)*8;
const unsigned int* vals = (unsigned int*)
compressedTail->elements;
@@ -1449,9 +1450,9 @@ str BWDGroupApproximate(int *rethisto, i
if((err = clEnqueueReadBuffer(getCommandQueue(),
batTailApproximation(b), CL_TRUE, 0, sizeof(int), &newTailDefinition.count , 0,
NULL, NULL)))
printf("#%s, clEnqueueReadBuffer (b): %s;\n", __func__,
clError(err));
{ // create output objects
- BAT* groupIDs = BATnew(b->htype, TYPE_oid, BATcount(b));
+ BAT* groupIDs = BATnew(TYPE_void, TYPE_oid,
BATcount(b));
BAT* histo = BATnew(TYPE_oid,TYPE_int, 0);
- const size_t histogramSize = pow(2.0f,
(float)batTailApproximationBits(b));
+ const size_t histogramSize = round(pow(2.0f,
(float)batTailApproximationBits(b)));
DecomposedBATSlot* histogramSlot;
{ // groupID object
@@ -1459,9 +1460,11 @@ str BWDGroupApproximate(int *rethisto, i
DecomposedBATSlot* groupIDSlot =
getDecomposedBATSlotForIndex(newIndex);
BATsetprop(groupIDs, batRegistryIndex,
TYPE_int, (int[]){newIndex});
- groupIDSlot->tailPositions =
bwdClCreateBuffer(getCLContext(), CL_MEM_READ_WRITE,
sizeof(int)*newTailDefinition.count+sizeof(clHead), NULL, &err);
- err = clEnqueueFillBuffer(getCommandQueue(),
groupIDSlot->tailPositions, zeroIntPattern, sizeof(int), 0, sizeof(clHead), 0,
NULL, NULL);
- if(err) printf("#%s, clEnqueueFillBuffer:
%s;\n", __func__, clError(err));
+ groupIDSlot->tailApproximation =
bwdClCreateBuffer(getCLContext(), CL_MEM_READ_WRITE,
sizeof(int)*newTailDefinition.count+sizeof(clHead), NULL, &err);
+ err = clEnqueueWriteBuffer(getCommandQueue(),
groupIDSlot->tailApproximation, CL_TRUE, 0, sizeof(clTail), &newTailDefinition,
0, NULL, NULL);
+ if(err) printf("#%s, clEnqueueWriteBuffer:
%s;\n", __func__, clError(err));
+ /* err = clEnqueueFillBuffer(getCommandQueue(),
groupIDSlot->tailApproximation, zeroIntPattern, sizeof(int), 0, sizeof(clHead),
0, NULL, NULL); */
+ /* if(err) printf("#%s, clEnqueueFillBuffer:
%s;\n", __func__, clError(err)); */
}
{ // histogram object
@@ -1487,7 +1490,7 @@ str BWDGroupApproximate(int *rethisto, i
for (i = 0; i < 4; ++i) {
if((err = clSetKernelArg(groupKernel,
i, sizeof(cl_mem), &((cl_mem[]){
-
batTailPositions(groupIDs),
+
batTailApproximation(groupIDs),
batHeadApproximation(histo),
batTailApproximation(histo),
batTailApproximation(b)}[i]
@@ -1554,6 +1557,7 @@ str BWDMulticolumnGroupApproximate(Clien
if((err = clEnqueueReadBuffer(getCommandQueue(),
batTailApproximation(inputs[0]), CL_TRUE, 0, sizeof(int),
&newTailDefinition.count , 0, NULL, NULL)))
printf("#%s, clEnqueueReadBuffer (b): %s;\n", __func__,
clError(err));
+ printf ("grouping %u values\n", newTailDefinition.count);
{ // create output objects
BAT* groupIDs = BATnew(TYPE_void, TYPE_oid,
BATcount(inputs[0]));
BAT* histo = BATnew(TYPE_oid,TYPE_int, 0);
@@ -1571,8 +1575,10 @@ str BWDMulticolumnGroupApproximate(Clien
groupIDSlot->tailOffsetBits = sizeof(int)*8;
groupIDSlot->tailApproximation =
bwdClCreateBuffer(getCLContext(), CL_MEM_READ_WRITE,
sizeof(int)*newTailDefinition.count+sizeof(clHead), NULL, &err);
- err = clEnqueueFillBuffer(getCommandQueue(),
groupIDSlot->tailApproximation, zeroIntPattern, sizeof(int), 0, sizeof(clHead),
0, NULL, NULL);
- if(err) printf("#%s, clEnqueueFillBuffer:
%s;\n", __func__, clError(err));
+ err = clEnqueueWriteBuffer(getCommandQueue(),
groupIDSlot->tailApproximation, CL_TRUE, 0, sizeof(clTail), &newTailDefinition,
0, NULL, NULL);
+ if(err) printf("#%s, clEnqueueWriteBuffer:
%s;\n", __func__, clError(err));
+ /* err = clEnqueueFillBuffer(getCommandQueue(),
groupIDSlot->tailApproximation, zeroIntPattern, sizeof(int), 0, sizeof(clHead),
0, NULL, NULL); */
+ /* if(err) printf("#%s, clEnqueueFillBuffer:
%s;\n", __func__, clError(err)); */
}
{ // histogram object
@@ -1683,7 +1689,7 @@ BAT* materializeBATInHostMemory(BAT* inp
assert(batTailResidualBits(input) == 0);
switch(BAThtype(input)){
case TYPE_oid:{
- register oid* outputRegion = (oid*) Tloc(result,
BUNfirst(result));
+ register oid* outputRegion = (oid*) Hloc(result,
BUNfirst(result));
unsigned int i;
clHead* positionsColumn;
@@ -1704,6 +1710,45 @@ BAT* materializeBATInHostMemory(BAT* inp
abort();
}
+ switch(BATttype(input)){
+ case TYPE_oid:{
+ register oid* outputRegion = (oid*) Tloc(result,
BUNfirst(result));
+ unsigned int i;
+
+ clTail* valueColumn;
+ size_t bufferSize;
+ getApproximateValuesColumn(batTailApproximation(input), NULL,
&bufferSize);
+ valueColumn =
getApproximateValuesColumn(batTailApproximation(input), GDKmalloc(bufferSize),
&bufferSize);
+ assert(count == valueColumn->count);
+ for (i = 0; i < count; ++i)
+ outputRegion[i] = ((unsigned
int*)valueColumn->elements)[i];
+ break;
+ }
+ case TYPE_int:{
+ register int* outputRegion = (oid*) Tloc(result,
BUNfirst(result));
+ unsigned int i;
+
+ clTail* valueColumn;
+ size_t bufferSize;
+ getApproximateValuesColumn(batTailApproximation(input), NULL,
&bufferSize);
+ valueColumn =
getApproximateValuesColumn(batTailApproximation(input), GDKmalloc(bufferSize),
&bufferSize);
+ if(count != valueColumn->count){
+ printf ("count != valueColumn->count, %u != %u\n",
count, valueColumn->count);
+ abort();
+ }
+ for (i = 0; i < count; ++i)
+ outputRegion[i] = ((unsigned
int*)valueColumn->elements)[i];
+ break;
+ }
+ case TYPE_void:{
+ BATseqbase(result, 0);
+ break;
+ }
+ default:
+ printf("unsupported tail type in materialization: %s, %s,
%s:%d", BATatoms[input->T->type].name, __func__, __FILE__, __LINE__);
+ abort();
+ }
+
BATsetcount(result, count);
return result;
}
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
@@ -109,7 +109,7 @@ char* humanreadablesize(unsigned int val
cl_mem bwdClCreateBuffer(cl_context context, cl_mem_flags flags, size_t size,
void *host_ptr, cl_int *errcode_ret){
- cl_mem result = clCreateBuffer (context, flags, ceil(size/128.0)*128,
host_ptr,errcode_ret);
+ cl_mem result = clCreateBuffer (context, flags, ceil(size/128.0)*128+8,
host_ptr,errcode_ret);
char buffer[64];
humanreadablesize(size, buffer, 64);
printf ("created cl_mem %p of size %s\n", result, buffer);
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
