Changeset: 385c42302cf9 for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=385c42302cf9
Modified Files:
monetdb5/extras/bwd/cl_program_utilities.c
monetdb5/extras/bwd/operations.c
monetdb5/extras/bwd/utilities.c
Branch: bwd
Log Message:
* bitwise selection seems to work for bytewise decompositions now
Unterschiede (199 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
@@ -158,6 +158,7 @@ cl_program getUSelectProgram(int type, c
" for(int i = 0; i < %4$d; i++)\n"
" value += (approximationTail->values[inputOffset + i] <<
((i+sizeof(%1$s)-%6$d - %4$d)*8));\n"
"#endif\n"
+ " printf(\"===== %%d (global id: %%d)\\n\", value,
get_global_id(0));"
"\n"
" if((value %2$s operand)"
" && (%5$d || value %3$s operand2)"
@@ -177,7 +178,8 @@ cl_program getUSelectProgram(int type, c
" }\n"
"}",
- [2] = "__kernel void uselect (\n" //
void-headed case
+ [2] = " __constant static const size_t targetTypeBits
= (sizeof(targetType)*8);"
+ "__kernel void uselect (\n" // void-headed case
"__global struct{int count; int padding; int positions[];}*
outputHead,\n"
"__global struct{int count; int base; unsigned char values[];}*
outputTail,\n"
"__global const struct{int count; int base; unsigned char
values[];}* approximationTail,\n"
@@ -185,15 +187,16 @@ cl_program getUSelectProgram(int type, c
"const %1$s operand2\n"
") {\n"
" if(get_global_id(0) < approximationTail->count){\n"
- " static const size_t targetTypeBits = (sizeof(targetType)*8);"
" targetType value = approximationTail->base;\n"
- " size_t slot =
(get_global_id(0)*approximationBits)/targetTypeBits;"
- " size_t offset =
(get_global_id(0)*approximationBits)%%targetTypeBits;"
- " {"
- " __global const unsigned char* vals =
approximationTail->values;"
- " value +=
((offset+approximationBits>targetTypeBits)*((vals[slot]<<(targetTypeBits-offset))
+ vals[slot+1]>>(approximationBits-targetTypeBits+offset))\n"
- " +
(offset+approximationBits>targetTypeBits)*(vals[slot]>>(targetTypeBits-offset-approximationBits)))&((1<<approximationBits)-1);"
- " }"
+ " size_t slot =
(get_global_id(0)*approximationBits)/targetTypeBits;\n"
+ " size_t offset =
(get_global_id(0)*approximationBits)%%targetTypeBits;\n"
+ " {\n"
+ " __global const unsigned char* vals =
approximationTail->values;\n"
+ " value +=
(((((offset+approximationBits)>targetTypeBits)*((vals[slot]<<(targetTypeBits-offset))
+ vals[slot+1]>>(approximationBits-targetTypeBits+offset))\n"
+ " +
((offset+approximationBits)<=targetTypeBits)*(vals[slot]>>(targetTypeBits-offset-approximationBits)))&((1<<approximationBits)-1))<<residualBits);\n"
+ " printf(\"===== %%d in slot %%d, offset %%d,
targetTypeBits: %%d, approximationBits: %%d (global id: %%d) = %%d\\n\", value,
slot, offset, targetTypeBits, approximationBits, get_global_id(0),
(((offset+approximationBits)>targetTypeBits)*((vals[slot]<<(targetTypeBits-offset))
+ vals[slot+1]>>(approximationBits-targetTypeBits+offset))\n"
+ " +
((offset+approximationBits)<=targetTypeBits)*(vals[slot]>>(targetTypeBits-offset-approximationBits))));\n"
+ " }\n"
"\n"
" if((value %2$s operand)"
" && (%5$d || value %3$s operand2)"
@@ -215,14 +218,14 @@ cl_program getUSelectProgram(int type, c
};
char* sourceCode = malloc(16384);
- static const str typeNames[] = {[TYPE_int] = "int"};
- static const size_t typeSizes[] = {[TYPE_int] = sizeof(int)};
- snprintf(sourceCode, 16384, sourceCodeTemplates[!!inputIsVoidHeaded],
typeNames[type], (approximationBits ==
8*typeSizes[type])?predicateOperation:(approximateOperation(predicateOperation)),
predicateOperation2?((approximationBits ==
8*typeSizes[type])?predicateOperation2:approximateOperation(predicateOperation2)):"==",
approximationBits/8-offsetBits/8, predicateOperation2 == NULL?1:0,
offsetBits/8);
+ const str typeNames[] = {[TYPE_int] = "int"};
+ const size_t typeSizes[] = {[TYPE_int] = sizeof(int)};
+ snprintf(sourceCode, 16384,
sourceCodeTemplates[(!!inputIsVoidHeaded)*2], typeNames[type],
(approximationBits ==
8*typeSizes[type])?predicateOperation:(approximateOperation(predicateOperation)),
predicateOperation2?((approximationBits ==
8*typeSizes[type])?predicateOperation2:approximateOperation(predicateOperation2)):"==",
approximationBits/8-offsetBits/8, predicateOperation2 == NULL?1:0,
offsetBits/8);
{
- char options[64];
+ char options[256];
cl_program program;
- snprintf(options, 64, "-D targetType=int -D
approximationBits=%d", approximationBits-offsetBits);
- printf ("%s\n", sourceCode);
+ snprintf(options, 256, "-D targetType=int -D
approximationBits=%d -D residualBits=%lu", approximationBits-offsetBits,
typeSizes[type]*8-approximationBits-offsetBits);
+ /* printf ("%s\n", sourceCode); */
program = compileProgram(sourceCode, options);
free(sourceCode);
return program;
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
@@ -376,11 +376,11 @@ static inline str uselect(bat *res, bat
{
int base = batTailOffsetValue(data);
+ int i;
const int parameters[2] = {
(((((*(int*)val)-base)>>batTailResidualBits(data))<<batTailResidualBits(data))
+ base),
val2?(((((*(int*)val2)-base)>>batTailResidualBits(data))<<batTailResidualBits(data))
+ base):0};
printf ("relaxed conditions from (%d,
%d) to (%d, %d)\n", *(int*)val, val2?(*(int*)val2):-1, parameters[0],
parameters[1]);
- int i;
for (i = 0; i < 2; ++i)
if((err =
clSetKernelArg(selectKernel, bufferI+i, sizeof(int), &(parameters[i])))) //
type specific
printf("#%s,
clSetKernelArg(%d): %s;\n", __func__, bufferI+i, clError(err));
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
@@ -160,9 +160,9 @@ cl_context getCLContext(void){
cl_uint numberOfDevices;
cl_platform_id platforms[4];
cl_uint foundPlatforms;
+ cl_int err;
if(clGetPlatformIDs(4,platforms,&foundPlatforms) != CL_SUCCESS)
printf("problem when finding the platforms");
if (!foundPlatforms) printf("didn't find any OpenCL Platforms");
- cl_int err;
if ((err = clGetDeviceIDs(*platforms,CL_DEVICE_TYPE_ALL, 4,
devices, &numberOfDevices)) != CL_SUCCESS) printf("error getting GPU device
id: %s\n", clError(err));
clContextSingleton = clCreateContext(0,numberOfDevices,
devices, NULL, NULL, &err);
initialized = 1;
@@ -202,14 +202,16 @@ const char* batRegistryIndex = "bwd.batR
return BATgetprop(subject, batRegistryIndex) != NULL;
}
- cl_mem batTailApproximation(const BAT* subject){
- PROPrec* rightTailApproximationProperty;
- if(!(rightTailApproximationProperty = BATgetprop(subject,
batRegistryIndex))){
- if(0) printf("%s: bat hasn't been decomposed;\n",
__func__);
- return NULL;
- }
+cl_mem batTailApproximation(const BAT* subject){
+ PROPrec* rightTailApproximationProperty;
+ if(!(rightTailApproximationProperty = BATgetprop(subject,
batRegistryIndex))){
+ if(0) printf("%s: bat hasn't been decomposed;\n", __func__);
+ return NULL;
+ }
+ {
DecomposedBATSlot* slot = getDecomposedBATSlot(subject);
return
slot->isPersistentBAT?slot->tailApproximations[getGPUDeviceForThisThread()]:slot->tailApproximation;
+ }
}
@@ -370,44 +372,48 @@ unsigned int decomposeIntArray(const int
*((unsigned
int*)(&approximation->elements[i*approximationBytes])) |=
((subject[i]-approximation->base) >> residualBits);
}
+ {
+ 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));
+ }
- 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] = {};
-
- 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));
- const cl_context context = getCLContext();
- char* sourceCode = "__kernel void dummy (__global const
struct{int count; int padding; char values[];}* inputTail1){}";
- 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));
- 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);
+ 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;
}
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
