Changeset: 8f1e2bac28ee for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=8f1e2bac28ee
Modified Files:
monetdb5/extras/bwd/cl_program_utilities.c
monetdb5/extras/bwd/cl_program_utilities.h
Branch: bwd
Log Message:
* implemented a (hard-coded) grouping of two attributes kernel
Unterschiede (64 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
@@ -267,6 +267,49 @@ cl_program getGroupProgram(unsigned int
return compileProgram(sourceCode, options);
}
+cl_program getTwoAttributeGroupProgram(unsigned int firstApproximationBits,
unsigned int secondApproximationBits){
+ const char* sourceCode = ""
+ "#define accessType unsigned int\n"
+ "__constant static const unsigned int firstApproximationMask =
((1<<firstApproximationBits)-1);\n"
+ "__constant static const unsigned int secondApproximationMask =
((1<<secondApproximationBits)-1);\n"
+ "__constant static const unsigned int targetTypeBits =
sizeof(accessType)*8;\n"
+ "__kernel void group (\n"
+ "__global unsigned int* groupIDs,\n"
+ "__global unsigned int* histogramOIDs,\n"
+ "__global unsigned int* histogramCounts,\n"
+ "__global const unsigned int* firstInputTail,\n"
+ "__global const unsigned int* secondInputTail\n"
+ ") {\n"
+ " const size_t inputIndex = get_global_id(0);\n"
+ " if(inputIndex < firstInputTail[0]){\n"
+ " size_t firstSlot =
(inputIndex*firstApproximationBits)/targetTypeBits;\n"
+ " size_t firstOffset =
(inputIndex*firstApproximationBits)%targetTypeBits;\n"
+ " __global const unsigned int* firstVals =
&(fistInputTail[2]);\n"
+ " const unsigned int firstDelta = (("
+ "
(((firstOffset+firstApproximationBits)>targetTypeBits)?(((firstVals[firstSlot]<<(firstApproximationBits-targetTypeBits+firstOffset))
+
(firstVals[firstSlot+1]>>(targetTypeBits-(firstApproximationBits-targetTypeBits+firstOffset))))&firstApproximationMask):0)\n"
+ " +
(((firstOffset+firstApproximationBits)<=targetTypeBits)*(firstVals[firstSlot]>>(targetTypeBits-firstOffset-firstApproximationBits)))\n"
+ " )&firstApproximationMask);\n"
+ "\n"
+ " size_t secondSlot =
(inputIndex*secondApproximationBits)/targetTypeBits;\n"
+ " size_t secondOffset =
(inputIndex*secondApproximationBits)%targetTypeBits;\n"
+ " __global const unsigned int* secondVals =
&(fistInputTail[2]);\n"
+ " const unsigned int secondDelta = (("
+ "
(((secondOffset+secondApproximationBits)>targetTypeBits)?(((secondVals[secondSlot]<<(secondApproximationBits-targetTypeBits+secondOffset))
+
(secondVals[secondSlot+1]>>(targetTypeBits-(secondApproximationBits-targetTypeBits+secondOffset))))&secondApproximationMask):0)\n"
+ " +
(((secondOffset+secondApproximationBits)<=targetTypeBits)*(secondVals[secondSlot]>>(targetTypeBits-secondOffset-secondApproximationBits)))\n"
+ " )&secondApproximationMask);\n"
+ "\n"
+ " const unsigned int delta =
firstDelta<<firstApproximationBits + secondDelta;"
+ " const uint oldValue = atomic_cmpxchg(histogramOIDs + delta +
2, 0, inputIndex);\n"
+ " const uint groupID = oldValue?oldValue:inputIndex;\n"
+ " groupIDs[inputIndex+2] = groupID;\n"
+ " atomic_inc(histogramCounts+delta+2);\n"
+ " }\n"
+ "}\n";
+ char options[256];
+ snprintf(options, 256, "-DfirstApproximationBits=%d
-DsecondApproximationBits=%d", firstApproximationBits, secondApproximationBits);
+ return compileProgram(sourceCode, options);
+}
+
cl_program getHistogramCompactionProgram(void){
const char* sourceCode = ""
"__kernel void compact (\n"
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 getTwoAttributeGroupProgram(unsigned int firstApproximationBits,
unsigned int secondApproximationBits);
cl_program getHistogramCompactionProgram(void);
cl_program compileProgram(const char* sourceCode, char* options);
/* cl_program getProjectionSemijoinProgram(unsigned int approximationBits); */
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
