Changeset: cd8a9702e032 for MonetDB URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=cd8a9702e032 Modified Files: monetdb5/extras/bwd/cl_program_utilities.c monetdb5/extras/bwd/cl_program_utilities.h monetdb5/extras/bwd/operations.c monetdb5/extras/bwd/utilities.c Branch: bwd Log Message:
* lots of performance tweaking (by now single threaded GPU/CPU is about twice
as fast as single threaded CPU and about twice as slow as parallel CPU)
Unterschiede (gekürzt von 384 auf 300 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
@@ -62,6 +62,7 @@ cl_program getProjectionLeftjoinProgram(
"__global struct{int count; int padding; char values[];}*
outputTail,\n"
"__global struct{int count; int padding; int positions[];}*
inputTail,"
"__global const struct{int count; int padding; char values[];}*
approximationTail\n) {\n"
+ " if(get_global_id(0) < inputTail->count){\n"
" __global const char* approximation =
approximationTail->values;"
" const int offset =
inputTail->positions[get_global_id(0)]*approximationBytes;\n"
" int value = 0;"
@@ -69,6 +70,7 @@ cl_program getProjectionLeftjoinProgram(
" outputTail->values[get_global_id(0)*approximationBytes+i]
= approximation[offset + i];\n"
" value += approximation[offset + i] << 8*i;"
" }\n"
+ " }\n"
/* " printf(\"projected value (%d): %d\\n\",
inputTail->positions[get_global_id(0)], value<<(8*(4-approximationBytes)));" */
"}";
char options[64];
@@ -101,6 +103,7 @@ cl_program getUSelectProgram(int type, c
"const %1$s operand,\n"
"const %1$s operand2\n"
") {\n"
+ " if(get_global_id(0) < approximationTail->count){\n"
" __global const char* approximation =
approximationTail->values;"
" %1$s value = approximationTail->base;\n"
" const size_t inputOffset = get_global_id(0)*%4$d;\n"
@@ -113,8 +116,8 @@ cl_program getUSelectProgram(int type, c
" && (%5$d || value %3$s operand2)"
" )"
"{\n"
- " const int index = atomic_inc(&(outputHead->count));"
- " atomic_inc(&(outputTail->count));" // TODO: this could
probably be done more efficiently
+ " 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\\n\", value, get_global_id(0), index);\n" */
" const int offset = index * %4$d;\n"
" outputHead->positions[index] = get_global_id(0);\n"
@@ -123,6 +126,7 @@ cl_program getUSelectProgram(int type, c
/* " printf(\"set byte %%d to %%d\\n\",
offset+i,outputTail->values[offset+i]);" */
" }\n"
" }\n"
+ " }\n"
"}";
char* sourceCode = malloc(16384);
snprintf(sourceCode, 16384, sourceCodeTemplate, (str[]){[TYPE_int] =
"int"}[type], approximateOperation(predicateOperation),
approximateOperation(predicateOperation2), approximationBits/8-offsetBits/8,
predicateOperation2 == NULL?1:0, offsetBits/8);
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
@@ -10,5 +10,6 @@
cl_program getUSelectProgram(int type, char* predicateOperation, char*
predicateOperation2, unsigned int approximationBits, unsigned int offsetBits);
cl_program getProjectionLeftjoinProgram(unsigned int approximationBits,
unsigned int offsetBits);
+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
@@ -16,6 +16,7 @@
#include "cl_program_utilities.h"
static const int activateWorkInProgress = 1;
+static const int synchronousGPU = 1;
#pragma mark Actual MAL Operations Implementation
@@ -133,8 +134,9 @@ str BWDLeftJoinApproximate(bat * res, ba
)))) printf("#%s,
clSetKernelArg(%d): %s;\n", __func__, 0, clError(err));
}
- if((err = clEnqueueNDRangeKernel(getCommandQueue(),
projectKernel, 1, (const size_t[]){0}, (const size_t[]){headCount}, (const
size_t[]){1}, 0, NULL, NULL)))
+ if((err = clEnqueueNDRangeKernel(getCommandQueue(),
projectKernel, 1, (const size_t[]){0}, (const
size_t[]){ceil(headCount/16.0)*16}, (const size_t[]){16}, 0, NULL, NULL)))
printf("#%s, clEnqueueNDRangeKernel: %s;\n", __func__,
clError(err));
+ if (synchronousGPU) clFinish(getCommandQueue());
BBPkeepref((*res = result->batCacheid));
BBPreleaseref(left->batCacheid);
@@ -294,10 +296,10 @@ static inline str uselect(bat *res, bat
if((err = clSetKernelArg(selectKernel, 3+i,
sizeof(int), &(parameters[i])))) // type specific
printf("#%s, clSetKernelArg(%d): %s;\n",
__func__, 3+i, clError(err));
- err = clEnqueueNDRangeKernel(getCommandQueue(), selectKernel,
1, (const size_t[]){0}, (const size_t[]){dataCount}, (const size_t[]){1}, 0,
NULL, NULL);
+ err = clEnqueueNDRangeKernel(getCommandQueue(), selectKernel,
1, (const size_t[]){0}, (const size_t[]){ceil(dataCount/16.0)*16}, (const
size_t[]){16}, 0, NULL, NULL);
if(err) printf("#%s, clEnqueueNDRangeKernel: %s;\n", __func__,
clError(err));
+ if (synchronousGPU) clFinish(getCommandQueue());
if(0){
- clFinish(getCommandQueue());
clTail* compressedTail;
size_t resultSize;
printf ("%s result tail result: %p\n",
__func__, batTailApproximation(result));
@@ -329,6 +331,7 @@ str uselectrefine(bat *res, bat *bid, pt
if(BATcount(data) == 0){
result = BATnew(BAThtype(data), TYPE_void, 0);
} else {
+
assert(ATOMstorage(BATttype(data)) == TYPE_int); // type
specific
BAT* approximation = BATdescriptor(*approx);
@@ -336,7 +339,6 @@ str uselectrefine(bat *res, bat *bid, pt
if(!headApproximation){
result = BATnew(BAThtype(data), TYPE_void, 0);
} else {
-
clHead* compressedHead;
{
size_t approximationSize;
@@ -347,45 +349,55 @@ str uselectrefine(bat *res, bat *bid, pt
}
const size_t candidateCount = compressedHead->count;
result = BATnew(BAThtype(data), TYPE_void,
candidateCount);
+ oid* positionRegion = (oid*) Hloc(result,
BUNfirst(result)); // type specific
+
+
+ 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;
+ for (i = 0; i < candidateCount; ++i)
+ positionRegion[i] = headPositions[i];
+ BATsetcount(result, candidateCount);
+ } else {
- clTail* compressedTail;
- {
- size_t approximationSize;
- if(0) printf ("%s approximation tail
approximation: %p\n", __func__, batTailApproximation(approximation));
+ 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);
- cl_int err =
clEnqueueReadBuffer(getCommandQueue(), batTailApproximation(approximation),
CL_TRUE, 0, approximationSize, compressedTail , 0, NULL, NULL);
- if(err) printf("#%s, clEnqueueReadBuffer:
%s;\n", __func__, clError(err));
- }
+
clGetMemObjectInfo(batTailApproximation(approximation), CL_MEM_SIZE,
sizeof(size_t), &approximationSize
+
, NULL);
+ compressedTail =
malloc(approximationSize);
+ cl_int err =
clEnqueueReadBuffer(getCommandQueue(), batTailApproximation(approximation),
CL_TRUE, 0, approximationSize, compressedTail , 0, NULL, NULL);
+ if(err) printf("#%s,
clEnqueueReadBuffer: %s;\n", __func__, clError(err));
+ }
- /* int* resultRegion = (int*) Tloc(result,
BUNfirst(result)); // type specific */
- oid* positionRegion = (oid*) Hloc(result,
BUNfirst(result)); // type specific
- const unsigned int approximationMask = ~((1 << (32 -
batTailApproximationBits(approximation)))-1);
- const unsigned char* residuals = batTailResiduals(data);
- const unsigned int residualMask = (1 <<
batTailResidualBits(data))-1;
- const unsigned int residualBytes =
batTailResidualBits(data)/8;
- int i = 0, j = 0;
+ /* int* resultRegion = (int*) Tloc(result,
BUNfirst(result)); // type specific */
+ const unsigned int approximationMask = ~((1 <<
(32 - batTailApproximationBits(approximation)))-1);
+ const unsigned char* residuals =
batTailResiduals(data);
+ const unsigned int residualMask = (1 <<
batTailResidualBits(data))-1;
+ const unsigned int residualBytes =
batTailResidualBits(data)/8;
+ int i = 0, j = 0;
- /* struct timespec before, after; */
- /* clock_gettime(CLOCK_THREAD_CPUTIME_ID, &before); */
- if(1){
- const int tailApproximationBytes =
batTailApproximationBits(approximation)/8;
- const int tailResidualBits =
batTailResidualBits(approximation);
- const int value1 = *(int*)val;
- const int value2 = val2?*(int*)val2:0;
+ /* struct timespec before, after; */
+ /* clock_gettime(CLOCK_THREAD_CPUTIME_ID,
&before); */
+ if(1){
+ const int tailApproximationBytes =
batTailApproximationBits(approximation)/8 - batTailOffsetBits(approximation)/8;
+ const int tailApproximationMask =
(1<<(tailApproximationBytes*8))-1;
+ const int tailResidualBits =
batTailResidualBits(approximation);
+ const int value1 = *(int*)val;
+ const int value2 = val2?*(int*)val2:0;
#define refineLoopDoubleOperator(comparator, comparator2)
\
while(i < candidateCount) {
\
const int index =
compressedHead->positions[i]; \
const int offset =
tailApproximationBytes*i++; \
- const int compressedValue =
compressedTail->base+*(int*)&(compressedTail->elements[offset]); \
- const int deCompressedValue =
\
- (compressedValue <<
tailResidualBits)
\
- +
(*(int*)&residuals[index*residualBytes] & residualMask); \
+ const int compressedValue =
*(int*)&(compressedTail->elements[offset]) & tailApproximationMask; \
+ const int deCompressedValue =
compressedTail->base+ \
+ (compressedValue <<
tailResidualBits)
\
+ +
(*(int*)&residuals[index*residualBytes] & residualMask); \
{
\
positionRegion[j] =
index;
\
j+= ((deCompressedValue
comparator value1) & \
@@ -394,72 +406,73 @@ str uselectrefine(bat *res, bat *bid, pt
}
#define refineLoopSingleOperator(comparator)
\
while(i < candidateCount) {
\
- const int index = compressedHead->positions[i];
\
- const int offset =
(tailApproximationBytes)*i++;
\
- const int compressedValue =
compressedTail->base+*(int*)&(compressedTail->elements[offset]); \
- const int deCompressedValue =
\
- (compressedValue << tailResidualBits)
\
- +
(*(int*)&residuals[index*residualBytes] & residualMask); \
- if(deCompressedValue comparator value1) {
\
- positionRegion[j++] = index;
\
- }
\
- }
-#define refineLoop(comparator, comparator2) \
- if(comparator2 == NULL) {
\
- refineLoopSingleOperator(comparator);
\
- } else {
\
- switch (comparator2[0]){
\
- case '<':
\
- switch((comparator2)[1]){
\
- case '\0':
\
- refineLoopDoubleOperator(comparator, <);
\
- case '=':
\
- refineLoopDoubleOperator(comparator, <=);
\
- }
\
- case '>':
\
- switch((comparator2)[1]){
\
- case '\0':
\
- refineLoopDoubleOperator(comparator, >);
\
- case '=':
\
- refineLoopDoubleOperator(comparator, <=);
\
- }
\
- case '=':
\
- refineLoopDoubleOperator(comparator, ==);
\
- }
\
- }
+ const int index =
compressedHead->positions[i];
\
+ const int offset =
(tailApproximationBytes)*i++;
\
+ const int compressedValue =
*(int*)&(compressedTail->elements[offset]) & tailApproximationMask; \
+ const int deCompressedValue =
compressedTail->base+ \
+ (compressedValue <<
tailResidualBits)
\
+ +
(*(int*)&residuals[index*residualBytes] & residualMask); \
+ if(deCompressedValue comparator value1)
{
\
+ positionRegion[j++] = index;
\
+ }
\
+ }
+#define refineLoop(comparator, comparator2)
\
+ if(comparator2 == NULL) {
\
+
refineLoopSingleOperator(comparator); \
+ } else {
\
+ switch (comparator2[0]){
\
+ case '<':
\
+
switch((comparator2)[1]){
\
+ case '\0':
\
+
refineLoopDoubleOperator(comparator, <); \
+ case '=':
\
+
refineLoopDoubleOperator(comparator, <=); \
+ }
\
+ case '>':
\
+
switch((comparator2)[1]){
\
+ case '\0':
\
+
refineLoopDoubleOperator(comparator, >); \
+ case '=':
\
+
refineLoopDoubleOperator(comparator, <=); \
+ }
\
+ case '=':
\
+
refineLoopDoubleOperator(comparator, ==); \
+ }
\
+ }
switch (OP[0]){
- case '<':
- switch(OP[1]){
- case '\0':
- refineLoop(<, OP2);
- case '=':
- refineLoop(<=, OP2);
- }
- case '>':
- switch(OP[1]){
- case '\0':
- refineLoop(>, OP2);
- case '=':
- refineLoop(>=, OP2);
- }
- case '=':
- refineLoop(==, OP2);
- }
+ case '<':
+ switch(OP[1]){
+ case '\0':
+ refineLoop(<, OP2);
+ case '=':
+ refineLoop(<=, OP2);
+ }
+ case '>':
+ switch(OP[1]){
+ case '\0':
+ refineLoop(>, OP2);
+ case '=':
+ refineLoop(>=, OP2);
+ }
+ case '=':
+ refineLoop(==, OP2);
+ }
#undef refineLoop
#undef refineLoopSingleOperator
#undef refineLoopDoubleOperator
+ }
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
