Changeset: 4e4b80672001 for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=4e4b80672001
Modified Files:
monetdb5/extras/bwd/utilities.c
Branch: bwd
Log Message:
* doing real BITwise decomposition now but not necessarily appropriate
processing :-|
Unterschiede (184 Zeilen):
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
@@ -333,101 +333,101 @@ char* humanreadablesize(unsigned int val
}
-unsigned int decomposeIntArray(const int* subject, const size_t size, const
size_t approximationBits){
- if (approximationBits % 8 > 0 || approximationBits <= 0) {
- printf("#decomposeTail: number of bits for approximation has to
be a multiple of 8, is %ld;\n", approximationBits);
- return -1;
- } else {
- const bounds subjectBounds = findBounds(subject,size);
- const int neededBits = ceil(log2(subjectBounds.max -
subjectBounds.min));
- const unsigned int newIndex =
getNextFreeDecomposedBATSlotIndex();
- DecomposedBATSlot* slot =
getDecomposedBATSlotForIndex(newIndex);
- const unsigned int residualBytes = (32-approximationBits)/8;
- const unsigned int approximationBytes =
approximationBits/8-(32-ceil(neededBits/8.0)*8)/8;
- const unsigned int residualBits = 32-approximationBits;
- const unsigned int residualMask = (1 << residualBits)-1;
- clTail* approximation = calloc((size+1)*approximationBytes+
sizeof(clTail), sizeof(char));
- unsigned int i;
- static cl_program program = NULL;
- cl_int err;
+unsigned int decomposeIntArray(const int* subject, const size_t size, const
size_t baseBitsPlusApproximationBits){
+ const bounds subjectBounds = findBounds(subject,size);
+ const int neededBits = ceil(log2(subjectBounds.max -
subjectBounds.min));
+ const unsigned int newIndex = getNextFreeDecomposedBATSlotIndex();
+ DecomposedBATSlot* slot = getDecomposedBATSlotForIndex(newIndex);
+ const unsigned int residualBits = 32-baseBitsPlusApproximationBits;
+ const unsigned int residualMask = (1 << residualBits)-1;
+ const unsigned int approximationBits = neededBits-residualBits;
+ const size_t approximationSize = ceil(size*approximationBits/8.0) +
sizeof(long) + sizeof(clTail);
+ clTail* approximation = calloc(approximationSize, sizeof(char));
+ unsigned int i;
+ static cl_program program = NULL;
+ cl_int err;
- slot->tailOffsetBits = 32-ceil(neededBits/8.0)*8;
+ slot->tailOffsetBits = 32-neededBits;
slot->approximationBits = approximationBits;
slot->isPersistentBAT = 1;
- printf ("bounds: %d, %d, offset %zd bytes\n",subjectBounds.min,
subjectBounds.max, slot->tailOffsetBits/8);
+ printf ("bounds: %d, %d, offset %zd bits\n",subjectBounds.min,
subjectBounds.max, slot->tailOffsetBits);
- printf ("using %s for approximation and %s for residuals \n",
humanreadablesize(approximationBytes*(size+1)+ sizeof(clTail), alloca(64), 64),
humanreadablesize(residualBytes*(size+1), alloca(64), 64));
- slot->residuals = calloc((size+1)*residualBytes, sizeof(char));
- approximation->base = subjectBounds.min;
- slot->tailOffsetValue = approximation->base;
- approximation->count = size;
+ printf ("using %s for approximation and %s for residuals \n",
humanreadablesize(approximationSize, alloca(64), 64),
humanreadablesize(ceil(size*residualBits/8.0) + 8, alloca(64), 64));
+ slot->residuals = calloc(ceil(size*residualBits/8.0) + 8, sizeof(char));
+ approximation->base = subjectBounds.min;
+ slot->tailOffsetValue = approximation->base;
+ approximation->count = size;
- for (i = 0; i < size; ++i) {
- { // approximation
- const size_t slotI =
(i*(approximationBytes*8))/32;
- const unsigned int offset =
(i*(approximationBytes*8))%32;
- if(offset+(approximationBytes*8) > 32) {
- ((unsigned
int*)approximation->elements)[slotI] |= (((subject[i]-approximation->base) >>
residualBits)>>((approximationBytes*8)-(32-offset)));
- ((unsigned
int*)approximation->elements)[slotI+1] |= (((subject[i]-approximation->base) >>
residualBits)<<(32-((approximationBytes*8)-(32-offset))));
- } else {
- ((unsigned
int*)approximation->elements)[slotI] |= (((subject[i]-approximation->base) >>
residualBits)<<(32-offset-(approximationBytes*8)));
- }
- }
- { // residual
- const size_t slotI = (i*residualBits)/32;
- const unsigned int offset = (i*residualBits)%32;
- ((unsigned int*)slot->residuals)[slotI] |=
(((subject[i]-approximation->base)&residualMask)<<(32-offset-residualBits));
+ for (i = 0; i < size; ++i) {
+ { // approximation
+ const size_t slotI = (i*approximationBits)/32;
+ const unsigned int offset = (i*approximationBits)%32;
+ if(offset+approximationBits > 32) {
+ ((unsigned int*)approximation->elements)[slotI]
|= (((subject[i]-approximation->base) >>
residualBits)>>(approximationBits-(32-offset)));
+ ((unsigned
int*)approximation->elements)[slotI+1] |= (((subject[i]-approximation->base) >>
residualBits)<<(32-(approximationBits-(32-offset))));
+ } else {
+ ((unsigned int*)approximation->elements)[slotI]
|= (((subject[i]-approximation->base) >>
residualBits)<<(32-offset-approximationBits));
}
}
+ { // residual
+ const size_t slotI = (i*residualBits)/32;
+ const unsigned int offset = (i*residualBits)%32;
+ if(offset+residualBits < 32) {
+ ((unsigned int*)slot->residuals)[slotI] |=
(((subject[i]-approximation->base)&residualMask)<<(32-offset-residualBits));
+ } else {
+ ((unsigned int*)slot->residuals)[slotI] |=
(((subject[i]-approximation->base)&residualMask)>>(residualBits-(32-offset)));
+ ((unsigned int*)slot->residuals)[slotI+1] |=
(((subject[i]-approximation->base)&residualMask)<<(32-(residualBits-(32-offset))));
+ }
+ }
+ }
- {
- cl_uint numberOfDevices;
- cl_device_id devices[4] = {};
- const cl_context context = getCLContext();
- const char* sourceCode = "__kernel void dummy (__global
const struct{int count; int padding; char values[];}* inputTail1){}";
- {
- cl_platform_id platforms[4];
- cl_uint foundPlatforms;
-
if(clGetPlatformIDs(4,platforms,&foundPlatforms) != CL_SUCCESS) printf("problem
when finding the platforms");
- if (!foundPlatforms) printf("didn't find any
OpenCL Platforms");
+ {
+ cl_uint numberOfDevices;
+ cl_device_id devices[4] = {};
+ const cl_context context = getCLContext();
+ const char* sourceCode = "__kernel void dummy (__global const
struct{int count; int padding; char values[];}* inputTail1){}";
+ {
+ cl_platform_id platforms[4];
+ cl_uint foundPlatforms;
+ if(clGetPlatformIDs(4,platforms,&foundPlatforms) !=
CL_SUCCESS) printf("problem when finding the platforms");
+ if (!foundPlatforms) printf("didn't find any OpenCL
Platforms");
- if ((err =
clGetDeviceIDs(*platforms,CL_DEVICE_TYPE_ALL, 4, devices, &numberOfDevices)) !=
CL_SUCCESS) printf("error getting GPU device id: %s\n", clError(err));
+ if ((err =
clGetDeviceIDs(*platforms,CL_DEVICE_TYPE_ALL, 4, devices, &numberOfDevices)) !=
CL_SUCCESS) printf("error getting GPU device id: %s\n", clError(err));
+ }
+
+ program = clCreateProgramWithSource(context, 1, (const
char*[]){sourceCode}, (size_t[]){strlen(sourceCode)}, &err);
+ if(err) printf("#%s, clCreateProgramWithSource: %s;\n",
__func__, clError(err));
+ err = clBuildProgram(program, numberOfDevices, devices, "",
NULL, NULL);
+ if(err) printf("#%s, clBuildProgram: %s;\n", __func__,
clError(err));
+
+
+
+ {
+ uint device;
+ for (device = 0; device < numberOfDevices; ++device) {
+ cl_command_queue queue =
clCreateCommandQueue(context,devices[device],0,&err);
+ cl_kernel dummyKernel = clCreateKernel(program,
"dummy", &err);
+
+ slot->tailApproximations[device] =
bwdClCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
approximationSize, approximation, &err);
+ if(err) printf("#%s, clCreateBuffer: %s;\n",
__func__, clError(err));
+
+ if(err) printf("#%s, clCreateKernel: %s;\n",
__func__, clError(err));
+ err = clSetKernelArg(dummyKernel, 0,
sizeof(cl_mem), (cl_mem[]){slot->tailApproximations[device]});
+ if(err) printf("#%s, clSetKernelArg: %s;\n",
__func__, clError(err));
+ if((err = clEnqueueNDRangeKernel(queue,
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));
+ clFinish(queue);
+
}
-
- program = clCreateProgramWithSource(context, 1, (const
char*[]){sourceCode}, (size_t[]){strlen(sourceCode)}, &err);
- if(err) printf("#%s, clCreateProgramWithSource: %s;\n",
__func__, clError(err));
- err = clBuildProgram(program, numberOfDevices, devices,
"", NULL, NULL);
- if(err) printf("#%s, clBuildProgram: %s;\n", __func__,
clError(err));
-
-
-
- {
- uint device;
- for (device = 0; device < numberOfDevices;
++device) {
- cl_command_queue queue =
clCreateCommandQueue(context,devices[device],0,&err);
- cl_kernel dummyKernel =
clCreateKernel(program, "dummy", &err);
-
- slot->tailApproximations[device] =
bwdClCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
size*approximationBytes+sizeof(size_t)+sizeof(clTail), approximation, &err);
- if(err) printf("#%s, clCreateBuffer:
%s;\n", __func__, clError(err));
-
- if(err) printf("#%s, clCreateKernel:
%s;\n", __func__, clError(err));
- err = clSetKernelArg(dummyKernel, 0,
sizeof(cl_mem), (cl_mem[]){slot->tailApproximations[device]});
- if(err) printf("#%s, clSetKernelArg:
%s;\n", __func__, clError(err));
- if((err = clEnqueueNDRangeKernel(queue,
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));
- clFinish(queue);
+ }
- }
- }
-
- }
- free(approximation);
- return newIndex;
}
+ free(approximation);
+ return newIndex;
}
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
