Changeset: 1f2999828fb6 for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=1f2999828fb6
Modified Files:
monetdb5/extras/bwd/cl_program_utilities.c
monetdb5/extras/bwd/operations.c
monetdb5/extras/bwd/utilities.c
Branch: bwd
Log Message:
* residuals for non-zero based values were calculated wrongly, fixed it
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
@@ -148,15 +148,15 @@ cl_program getUSelectProgram(int type, c
" for(int i = 0; i < %4$d; i++)\n"
" value += (approximation[inputOffset + i] <<
((i+sizeof(%1$s)-%6$d - %4$d)*8));\n"
"\n"
- /* " if(get_global_id(0) < 10)\n" */
- /* " printf(\"%%d: %%d, base: %%d\\n\", get_global_id(0),
value, approximationTail->base);\n" */
+ " if(get_global_id(0) < 10)\n"
+ " printf(\"%%d: %%d, base: %%d\\n\", get_global_id(0),
value, approximationTail->base);\n"
" if((value %2$s operand)"
" && (%5$d || value %3$s operand2)"
" )"
"{\n"
" 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" */
+ " 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"
" for(int i = 0; i < %4$d; i++){\n"
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
@@ -20,7 +20,11 @@
static const int activateWorkInProgress = 1;
static const int synchronousGPU = 1;
static const int eagerBufferLoading = 1;
-#define WORK_GROUP_SIZE 16
+#ifdef __APPLE__
+static const int WORK_GROUP_SIZE=1;
+#else
+static const int WORK_GROUP_SIZE=16;
+#endif
#define MAX_INTERMEDIATE_RESULT_SIZE 16777216
#pragma mark Actual MAL Operations Implementation
@@ -455,6 +459,7 @@ str uselectrefine(bat *res, bat *bid, pt
const int deCompressedValue =
compressedTail->base+ \
(compressedValue <<
tailResidualBits)
\
+
(*(int*)&residuals[index*residualBytes] & residualMask); \
+ printf ("compr: %d, decompr:
%d\n", compressedValue, deCompressedValue); \
{
\
positionRegion[j] =
index;
\
j+= ((deCompressedValue
comparator value1) & \
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
@@ -279,7 +279,7 @@ const unsigned int decomposeIntArray(con
const register unsigned int residualMask = (1 << residualBits)-1;
int i;
for (i = 0; i < size; ++i){
- *((unsigned int*)(slot->residuals + i*residualBytes)) |=
(subject[i]&residualMask);
+ *((unsigned int*)(slot->residuals + i*residualBytes)) |=
((subject[i]-approximation->base)&residualMask);
*((unsigned
int*)(&approximation->elements[i*approximationBytes])) |=
((subject[i]-approximation->base) >> residualBits);
}
{
@@ -295,6 +295,20 @@ const unsigned int decomposeIntArray(con
program = compileProgram("__kernel void dummy (__global const
struct{int count; int padding; char values[];}* inputTail1){}","");
cl_kernel dummyKernel = clCreateKernel(program, "dummy", &err);
+ if(0)
+ for (int i = 0; i < size; ++i) {
+ int approximationValue = 0, residualValue = 0;
+ for (int j = 0; j < approximationBytes ; ++j){
+ printf ("%d\n",
approximation->elements[i*approximationBytes+j]);
+ approximationValue =
approximation->elements[i*approximationBytes+j]<<(j*8);
+ }
+
+ for (int j = 0; j < residualBytes; ++j)
+ residualValue =
slot->residuals[i*residualBytes+j]<<(j*8);
+ printf ("base: %d, approximation: %d, refinement:
%d\n", approximation->base,
approximation->base+(approximationValue<<(residualBytes*8)),
approximation->base+(approximationValue<<(residualBytes*8)) + residualValue);
+
+ }
+
clSetKernelArg(dummyKernel, 0, sizeof(cl_mem),
(cl_mem[]){slot->tailApproximation});
if((err = clEnqueueNDRangeKernel(getCommandQueue(), dummyKernel, 1,
(const size_t[]){0}, (const size_t[]){1}, (const size_t[]){1}, 0, NULL, NULL)))
printf("#%s, clEnqueueNDRangeKernel: %s;\n", __func__,
clError(err));
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
