Changeset: 1ac389ecb1d9 for MonetDB URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=1ac389ecb1d9 Modified Files: monetdb5/extras/bwd/cl_program_utilities.c monetdb5/extras/bwd/operations.c Branch: bwd Log Message:
* turns out we have to use atomic instructions when bitwise projecting
approximate values
Unterschiede (79 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
@@ -86,42 +86,25 @@ cl_program getProjectionLeftjoinProgram(
"__global struct{int count; int base; unsigned int
values[];}* outputTail,\n"
"__global struct{int count; int padding; unsigned int
positions[];}* inputTail,\n"
"__global const struct{int count; int base; unsigned
int values[];}* approximationTail\n) {\n"
- " if(get_global_id(0) < inputTail->count){\n"
- /* " int value = approximationTail->base;" */
- /* "#if approximationBytes == 2\n" */
- /* " outputTail->values[get_global_id(0)] =
approximationTail->values[inputTail->positions[get_global_id(0)]];\n" */
- /* "#else\n" */
- /* " __global const unsigned char* approximation =
approximationTail->values;\n" */
- /* " const int offset =
inputTail->positions[get_global_id(0)]*approximationBytes;\n" */
- /* " for(int i = 0; i < approximationBytes; i++){\n" */
- /* "
outputTail->values[get_global_id(0)*approximationBytes+i] =
approximation[offset + i];\n" */
- /* " }\n" */
- /* "#endif\n" */
+ " const size_t index = get_global_id(0);\n"
+ " if(index < inputTail->count){\n"
" const size_t inputIndex =
inputTail->positions[get_global_id(0)];"
" size_t slot =
(inputIndex*approximationBits)/targetTypeBits;\n"
" size_t offset =
(inputIndex*approximationBits)%targetTypeBits;\n"
- /* " printf(\"getting position %u\\n\",
get_global_id(0));\n" */
" __global const unsigned int* vals =
approximationTail->values;\n"
" unsigned int delta = (("
"
(((offset+approximationBits)>targetTypeBits)?((vals[slot]<<(targetTypeBits-offset))
+ vals[slot+1]>>(approximationBits-targetTypeBits+offset)):0)\n"
- /* "0" */
" +
(((offset+approximationBits)<=targetTypeBits)*(vals[slot]>>(targetTypeBits-offset-approximationBits)))"
" )&approximationMask);\n"
- /* " printf(\"projected delta #%u: %u\\n\",
get_global_id(0), delta);\n" */
-
- " const size_t index = get_global_id(0);\n"
" size_t outslot =
(index*approximationBits)/targetTypeBits;\n"
" size_t outoffset =
(index*approximationBits)%targetTypeBits;\n"
" if(outoffset+approximationBits >
8*sizeof(accessType))"
- " outputTail->values[outslot] |= (delta <<
(approximationBits-(8*sizeof(accessType)-outoffset)));"
+ " atomic_add(&(outputTail->values[outslot]),
(delta << (approximationBits-(8*sizeof(accessType)-outoffset))));"
" else"
- " outputTail->values[outslot] |= (delta <<
(8*sizeof(accessType)-outoffset-approximationBits));"
- /* " value += approximation[offset + i] << 8*(i+1);"
*/
+ " atomic_add(&(outputTail->values[outslot]),
(delta << (8*sizeof(accessType)-outoffset-approximationBits)));"
" int value =
approximationTail->base+(delta<<residualBits);"
- " if(index<10)"
+ " if(index==2114)"
" printf(\"%d->%d: offset+(%d>>x)=%d,
val[slot=%u]=%u, %d\\n\", inputIndex, get_global_id(0), delta, value, outslot,
outputTail->values[outslot],
(8*sizeof(accessType)-outoffset-approximationBits));"
-
- /* " printf(\"projected value (%d): %d (base: %d + %d
+ %d<<8, outbase: %d), approximationBytes: %d\\n\",
inputTail->positions[get_global_id(0)], value, approximationTail->base,
approximation[offset], approximation[offset + 1], outputTail->base,
approximationBytes);" */
" }\n"
"}";
char options[256];
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
@@ -481,15 +481,15 @@ static inline unsigned int refinementLoo
const int deCompressedValue = compressedTail->base+
\
(compressedValue << tailResidualBits)
\
+ residual;
\
- {
\
- const int positionIsCorrectPositive =
\
- (positionIsTruePositiveCondition);
\
- positionRegion[j] = index;
\
- j += ((deCompressedValue comparator value1) &
\
- (positionIsCorrectPositive));
\
- inputPositionsI += positionIsCorrectPositive;
\
- i++;
\
-}
\
+ {
\
+ const int positionIsCorrectPositive =
\
+ (positionIsTruePositiveCondition);
\
+ positionRegion[j] = index;
\
+ j += ((deCompressedValue comparator value1) &
\
+ (positionIsCorrectPositive));
\
+ inputPositionsI += positionIsCorrectPositive;
\
+ i++;
\
+ }
\
}
#define refineLoop(comparator, comparator2, positionIsTruePositiveCondition,
residualPosition) \
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
