Changeset: 39c77f2493ff for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=39c77f2493ff
Modified Files:
monetdb5/extras/bwd/cl_program_utilities.c
monetdb5/extras/bwd/cl_program_utilities.h
monetdb5/extras/bwd/operations.c
Branch: bwd
Log Message:
* implemented a compaction kernel for the grouping histogram (so far, the
histogram was created sparse)
Unterschiede (123 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
@@ -266,3 +266,22 @@ cl_program getGroupProgram(unsigned int
snprintf(options, 256, "-DapproximationBits=%d", approximationBits);
return compileProgram(sourceCode, options);
}
+
+cl_program getHistogramCompactionProgram(void){
+ const char* sourceCode = ""
+ "__kernel void compact (\n"
+ "__global const unsigned int* oldHistogramOIDs,\n"
+ "__global const unsigned int* oldHistogramCounts,\n"
+ "__global unsigned int* newHistogramOIDs,\n"
+ "__global unsigned int* newHistogramCounts\n"
+ ") {\n"
+ " const size_t inputIndex = get_global_id(0);\n"
+ " if(oldHistogramCounts[inputIndex] > 0){"
+ " atomic_inc(newHistogramCounts);"
+ " const unsigned int newSlot = atomic_inc(newHistogramOIDs);"
+ " newHistogramOIDs[newSlot] = oldHistogramOIDs[inputIndex];"
+ " newHistogramCounts[newSlot] =
oldHistogramCounts[inputIndex];"
+ " }"
+ "}\n";
+ return compileProgram(sourceCode, "");
+}
diff --git a/monetdb5/extras/bwd/cl_program_utilities.h
b/monetdb5/extras/bwd/cl_program_utilities.h
--- a/monetdb5/extras/bwd/cl_program_utilities.h
+++ b/monetdb5/extras/bwd/cl_program_utilities.h
@@ -11,6 +11,7 @@
cl_program getUSelectProgram(int type, char* predicateOperation, char*
predicateOperation2, unsigned int approximationBits, unsigned int offsetBits,
char inputIsVoidHeaded);
cl_program getProjectionLeftjoinProgram(unsigned int approximationBits,
unsigned int offsetBits);
cl_program getGroupProgram(unsigned int approximationBits);
+cl_program getHistogramCompactionProgram(void);
cl_program compileProgram(const char* sourceCode, char* options);
/* cl_program getProjectionSemijoinProgram(unsigned int approximationBits); */
#endif /* _CL_PROGRAM_UTILITIES_H_ */
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
@@ -1123,7 +1123,16 @@ str BWDMulticolumnGroupApproximate(Clien
return MAL_SUCCEED;// GRPmulticolumngroup(cntxt, mb, stk, pci);
}
-
+cl_mem createZeroHeadedCopy(cl_mem input){
+ size_t bufferSize;
+ cl_int err;
+ cl_mem result;
+ clGetMemObjectInfo(input, CL_MEM_SIZE, sizeof(size_t), &bufferSize,
NULL);
+ result = bwdClCreateBuffer(getCLContext(), CL_MEM_READ_WRITE,
bufferSize, NULL, &err);;
+ err = clEnqueueFillBuffer(getCommandQueue(), result, zeroIntPattern,
sizeof(int), 0, sizeof(clHead), 0, NULL, NULL);
+ if(err) printf("#%s, clEnqueueFillBuffer: %s;\n", __func__,
clError(err));
+ return result;
+}
str BWDGroupApproximate(int *rethisto, int *retbid, int *bid){
// this function should be used as a template for new BWD-operators
str result = MAL_SUCCEED;
@@ -1143,6 +1152,9 @@ str BWDGroupApproximate(int *rethisto, i
{ // create output objects
BAT* groupIDs = BATnew(b->htype, TYPE_oid, BATcount(b));
BAT* histo = BATnew(TYPE_oid,TYPE_int, 0);
+ const size_t histogramSize = pow(2.0f,
(float)batTailApproximationBits(b));
+ DecomposedBATSlot* histogramSlot;
+
{ // groupID object
const unsigned int newIndex =
getNextFreeDecomposedBATSlotIndex();
DecomposedBATSlot* groupIDSlot =
getDecomposedBATSlotForIndex(newIndex);
@@ -1155,9 +1167,8 @@ str BWDGroupApproximate(int *rethisto, i
}
{ // histogram object
const unsigned int newIndex =
getNextFreeDecomposedBATSlotIndex();
- DecomposedBATSlot* histogramSlot =
getDecomposedBATSlotForIndex(newIndex);
+ histogramSlot =
getDecomposedBATSlotForIndex(newIndex);
BATsetprop(histo, batRegistryIndex, TYPE_int,
(int[]){newIndex});
- const size_t histogramSize = pow(2.0f,
(float)batTailApproximationBits(b));
{
histogramSlot->tailPositions =
bwdClCreateBuffer(getCLContext(), CL_MEM_READ_WRITE,
sizeof(int)*histogramSize+sizeof(clHead), NULL, &err);
@@ -1170,9 +1181,11 @@ str BWDGroupApproximate(int *rethisto, i
if(err) printf("#%s,
clEnqueueFillBuffer: %s;\n", __func__, clError(err));
}
}
- { // run kernel
+ { // run group kernel
+ int i;
cl_kernel groupKernel =
clCreateKernel(getGroupProgram(batTailApproximationBits(b)), "group", &err);
- int i;
+ if(err) printf("#%s, clCreateKernel: %s;\n",
__func__, clError(err));
+
for (i = 0; i < 4; ++i) {
if((err = clSetKernelArg(groupKernel,
i, sizeof(cl_mem), &((cl_mem[]){
batTailPositions(groupIDs),
@@ -1184,6 +1197,29 @@ str BWDGroupApproximate(int *rethisto, i
if((err =
clEnqueueNDRangeKernel(getCommandQueue(), groupKernel, 1, (const size_t[]){0},
(const
size_t[]){ceil(newTailDefinition.count/((float)WORK_GROUP_SIZE))*WORK_GROUP_SIZE},
(const size_t[]){WORK_GROUP_SIZE}, 0, NULL, NULL)))
printf("#%s, clEnqueueNDRangeKernel:
%s;\n", __func__, clError(err));
if (synchronousGPU) clFinish(getCommandQueue());
+
+ }
+ { // run histogram compaction kernel
+ cl_kernel compactKernel =
clCreateKernel(getHistogramCompactionProgram(), "compact", &err);
+ int i;
+ cl_mem newHead =
createZeroHeadedCopy(batHeadApproximation(histo));
+ cl_mem newTail =
createZeroHeadedCopy(batTailApproximation(histo));
+ for (i = 0; i < 4; ++i) {
+ if((err = clSetKernelArg(compactKernel,
i, sizeof(cl_mem), &((cl_mem[]){
+
batHeadApproximation(histo),
+
batTailApproximation(histo),
+
newHead, newTail}[i]
+ ))))
printf("#%s, clSetKernelArg(%d): %s;\n", __func__, i, clError(err));
+ }
+ if((err =
clEnqueueNDRangeKernel(getCommandQueue(), compactKernel, 1, (const
size_t[]){0}, (const size_t[]){histogramSize}, (const size_t[]){2}, 0, NULL,
NULL)))
+ printf("#%s, clEnqueueNDRangeKernel:
%s;\n", __func__, clError(err));
+ if (synchronousGPU) clFinish(getCommandQueue());
+
clReleaseMemObject(histogramSlot->tailApproximation);
+ histogramSlot->tailApproximation = newTail;
+
clReleaseMemObject(histogramSlot->headApproximation);
+ histogramSlot->headApproximation = newHead;
+ histogramSlot->tailPositions =
histogramSlot->headApproximation;
+
}
BBPkeepref(*retbid = groupIDs->batCacheid);
BBPkeepref(*rethisto = histo->batCacheid);
_______________________________________________
checkin-list mailing list
checkin-list@monetdb.org
https://www.monetdb.org/mailman/listinfo/checkin-list
