Changeset: a4111b6c0f42 for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=a4111b6c0f42
Modified Files:
monetdb5/extras/bwd/91_bwd.sql
monetdb5/extras/bwd/Makefile.ag
monetdb5/extras/bwd/bwd.c
monetdb5/extras/bwd/bwd.h
monetdb5/extras/bwd/cl_program_utilities.c
monetdb5/extras/bwd/cl_program_utilities.h
monetdb5/extras/bwd/operations.c
monetdb5/extras/bwd/opt_bwd.mal
monetdb5/extras/bwd/optimizer.c
monetdb5/extras/bwd/utilities.c
Branch: bwd
Log Message:
* working on tpch q6
* implemented stubs for all of the operators
* the query runs but doesn't really produce useful results
Unterschiede (gekürzt von 990 auf 300 Zeilen):
diff --git a/monetdb5/extras/bwd/91_bwd.sql b/monetdb5/extras/bwd/91_bwd.sql
--- a/monetdb5/extras/bwd/91_bwd.sql
+++ b/monetdb5/extras/bwd/91_bwd.sql
@@ -1,6 +1,14 @@
+create function bwdecompose(col decimal(7,2), bits integer)
+returns varchar(4096) external name bwd.decompose;
+
+create function bwdecomposedate(col date, bits integer)
+returns varchar(4096) external name bwd.decompose;
+
+
create function bwdecompose(col integer, bits integer)
returns varchar(4096) external name bwd.decompose;
+
create function bwdevices()
returns table(platform int, device int, deviceDescription string,
platformDescription string) external name bwd.deviceInfo;
diff --git a/monetdb5/extras/bwd/Makefile.ag b/monetdb5/extras/bwd/Makefile.ag
--- a/monetdb5/extras/bwd/Makefile.ag
+++ b/monetdb5/extras/bwd/Makefile.ag
@@ -22,6 +22,14 @@ INCLUDES = \
../../../common/stream \
../../../common/options \
../../modules/kernel \
+ ../../modules/mal \
+ ../../modules/atoms \
+ ../../../sql/server \
+ ../../../sql/include \
+ ../../../sql/common \
+ ../../../sql/storage \
+ ../../../sql/backends/monet5 \
+ ../../../clients/mapilib \
$(OPENCL_CFLAGS)
MTSAFE
diff --git a/monetdb5/extras/bwd/bwd.c b/monetdb5/extras/bwd/bwd.c
--- a/monetdb5/extras/bwd/bwd.c
+++ b/monetdb5/extras/bwd/bwd.c
@@ -4,6 +4,11 @@
#include "opt_statistics.h"
#include "bwd.h"
+#include "sql_mvc.h"
+
+#include "sql.h"
+
+
#if defined(HAVE_OPENCL_OPENCL_H)
#include <OpenCL/opencl.h>
#elif defined(HAVE_CL_CL_H)
@@ -93,11 +98,11 @@ str deviceInfo(bat * resPlatform, bat *
-str bwdecompose(bat * res, bat * subjectBAT, int* approximationBits){
+str bwdecompose(bat * res, bat * subjectBAT, int* approximationBits, Client
cntxt){
BAT* subject = BATdescriptor(*subjectBAT);
if (VIEWtparent(subject))
subject = BATdescriptor(abs(VIEWtparent(subject)));
- printf ("decomposing bat tail %d into %d and %d bits\n",
subject->batCacheid, approximationBits[0], Tsize(subject)*8-*approximationBits);
+ printf ("decomposing %s into %d and %d bits\n",
resolveBatToAttribute(subject->batCacheid, cntxt), approximationBits[0],
Tsize(subject)*8-*approximationBits);
BATsetprop(subject, batRegistryIndex, TYPE_int,
(int[]){decomposeIntArray((int*)Tloc(subject, BUNfirst(subject)),
subject->batCount, *approximationBits)});
BAT* result = BATnew(TYPE_void, TYPE_str, 0);
char buffer[4096];
@@ -109,6 +114,41 @@ str bwdecompose(bat * res, bat * subject
+char* resolveBatToAttribute(bat bid, Client cntxt){
+ mvc *m = NULL;
+ str msg = getSQLContext(cntxt, NULL, &m, NULL);
+ sql_trans *tr = m->session->tr;
+ node *nsch, *ntab, *ncol;
+
+ for( nsch= tr->schemas.set->h; nsch; nsch= nsch->next){
+ sql_base *b= nsch->data;
+ sql_schema *s= (sql_schema*) nsch->data;
+ if ( isalpha((int)b->name[0]) )
+
+ if (s->tables.set) for(ntab= (s)->tables.set->h ;ntab;
ntab= ntab->next){
+ sql_base *bt= ntab->data;
+ sql_table *t= (sql_table*) bt;
+ if (isTable(t))
+ if (t->columns.set) for (ncol=
(t)->columns.set->h; ncol; ncol= ncol->next){
+ sql_base *bc =
ncol->data;
+ sql_column *c=
(sql_column *) ncol->data;
+ BAT *bn =
store_funcs.bind_col(tr, c, 0);
+ lng sz;
+
+
if(bn->batCacheid == bid){
+ char*
buffer = malloc(strlen(bc->name) + strlen(bt->name)+15);
+
snprintf(buffer, strlen(bc->name) + strlen(bt->name)+15, "%s.%s (bat %d)",
bt->name, bc->name, bid);
+ return
buffer;
+
+ }
+ }
+ }
+ }
+ return "unknown";
+
+ /* BAT *bn = store_funcs.bind_col(tr, c, 0); */
+
+}
str
diff --git a/monetdb5/extras/bwd/bwd.h b/monetdb5/extras/bwd/bwd.h
--- a/monetdb5/extras/bwd/bwd.h
+++ b/monetdb5/extras/bwd/bwd.h
@@ -7,7 +7,7 @@
extern str OPTBWD(Client cntxt, MalBlkPtr mb, MalStkPtr stk, InstrPtr p);
str deviceInfo(bat * resPlatform, bat * resDevice, bat * resDesc, bat *
resKernels);
str bwd_result_column_wrap(Client cntxt, MalBlkPtr mb, MalStkPtr stk, InstrPtr
pci);
-
+char* resolveBatToAttribute(bat bid, Client cntxt);
#endif /* _BWD_H_ */
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
@@ -23,49 +23,69 @@ char* approximateOperation(char* exactOp
static inline void superverboseprintf(const char * format, ... ){};
-cl_program getProjectionLeftjoinProgram(unsigned int approximationBits){
- const char* sourceCodeTemplate = "__kernel void project
(\n"
- "__global struct{int count; int padding; char
values[];}* outputTail,\n"
- "__global struct{int count; int padding; int
positions[];}* inputTail,"
- "__global const char* approximation\n"
- ") {\n"
- " const int offset =
inputTail->positions[get_global_id(0)]*%1$d;\n"
- " for(int i = 0; i < %1$d; i++){\n"
- "
outputTail->values[get_global_id(0)*%1$d+i] = approximation[offset + i];\n"
- "}"
- "}";
- char* sourceCode = malloc(16384);
- snprintf(sourceCode, 16384, sourceCodeTemplate,
approximationBits/8);
- superverboseprintf("#%s, sourceCode: %s;\n", __func__,
sourceCode);
- cl_int err;
- cl_program program = clCreateProgramWithSource(getCLContext(),
1, (const char*[]){sourceCode}, (size_t[]){strlen(sourceCode)}, &err);
- if(err) THRprintf(GDKout, "#%s, clCreateProgramWithSource:
%s;\n", __func__, clError(err));
- err = clBuildProgram(program, 1, (const
cl_device_id[]){getDeviceID()}, "", NULL, NULL);
- if(err) {
- printf("#%s, clBuildProgram: %s;\n", __func__,
clError(err));
- size_t bufferSize;
- char* buffer;
- clGetProgramBuildInfo(program, getDeviceID(),
CL_PROGRAM_BUILD_LOG, 0, NULL, &bufferSize);
- clGetProgramBuildInfo(program, getDeviceID(),
CL_PROGRAM_BUILD_LOG, bufferSize, (buffer = alloca(bufferSize)), NULL);
- printf("#%s, clBuildProgram log: %s;\n", __func__,
buffer);
- }
- free(sourceCode);
- return program;
+cl_program compileProgram(const char* sourceCode, char* options){
+ cl_int err;
+ cl_program program = clCreateProgramWithSource(getCLContext(), 1,
(const char*[]){sourceCode}, (size_t[]){strlen(sourceCode)}, &err);
+ if(err) THRprintf(GDKout, "#%s, clCreateProgramWithSource: %s;\n",
__func__, clError(err));
+ err = clBuildProgram(program, 1, (const cl_device_id[]){getDeviceID()},
"", NULL, NULL);
+ if(err) {
+ printf("#%s, clBuildProgram: %s;\n", __func__, clError(err));
+ size_t bufferSize;
+ char* buffer;
+ clGetProgramBuildInfo(program, getDeviceID(),
CL_PROGRAM_BUILD_LOG, 0, NULL, &bufferSize);
+ clGetProgramBuildInfo(program, getDeviceID(),
CL_PROGRAM_BUILD_LOG, bufferSize, (buffer = alloca(bufferSize)), NULL);
+ printf("#%s, clBuildProgram log: %s;\n", __func__, buffer);
+ }
+ return program;
}
-cl_program getUSelectProgram(int type, char* predicateOperation, unsigned int
approximationBits){
+cl_program getProjectionLeftjoinProgram(unsigned int approximationBits){
+ const char* sourceCode = "__kernel void project (\n"
+ "__global struct{int count; int padding; char values[];}*
outputTail,\n"
+ "__global struct{int count; int padding; int positions[];}*
inputTail,"
+ "__global const char* approximation\n) {\n"
+ " const int offset =
inputTail->positions[get_global_id(0)]*%1$d;\n"
+ " for(int i = 0; i < approximationBytes; i++)\n"
+ " outputTail->values[get_global_id(0)*%1$d+i] =
approximation[offset + i];\n"
+ "}";
+ char options[64];
+ snprintf(options, 64, "-D approximationBytes=%d",
approximationBits/8);
+ return compileProgram(sourceCode, options);
+}
+
+/* cl_program getProjectionSemijoinProgram(unsigned int approximationBits){ */
+/* const char* sourceCode = "__kernel void project (\n" */
+/* "__global struct{int count; int padding; int positions[];}*
outputHead,\n" */
+/* "__global struct{int count; int padding; char values[];}*
outputTail,\n" */
+/* "__global struct{int count; int padding; int positions[];}*
inputHead){\n" */
+/* /\* "__global const char* approximation\n) {\n" *\/ */
+/* " const int offset =
inputHead->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" */
+/* "}"; */
+/* char options[64]; */
+/* snprintf(options, 1024, "-D approximationBytes=%d",
approximationBits/8); */
+/* return compileProgram(sourceCode, options); */
+/* } */
+
+cl_program getUSelectProgram(int type, char* predicateOperation, char*
predicateOperation2, unsigned int approximationBits){
const char* sourceCodeTemplate = "__kernel void uselect
(\n"
- "__global struct{int count; int padding; int
positions[];}* outputHead,"
+ "__global struct{int count; int padding; int
positions[];}* outputHead,\n"
"__global struct{int count; int padding; char
values[];}* outputTail,\n"
- "__global const char* approximation,\n"
- "const %1$s operand\n"
+ "__global const char* approximation"
+ ",\n"
+ "const %1$s operand,\n"
+ "const %1$s operand2\n"
") {\n"
" %1$s value = 0;\n"
" const size_t inputOffset = get_global_id(0)*%3$d;\n"
" for(int i = 0; i < %3$d; i++)\n"
" value += (approximation[inputOffset + i] <<
((i+sizeof(%1$s) - %3$d)*8));\n"
"\n"
- " if(value %2$s operand){\n"
+ " if((value %2$s operand)"
+ " && (%4$d || value %2$s operand2)"
+ " )"
+ "{\n"
" const int index =
atomic_inc(&(outputHead->count));"
" const int offset = index * %3$d;\n"
" outputHead->positions[index] = get_global_id(0);\n"
@@ -74,20 +94,8 @@ cl_program getUSelectProgram(int type, c
" }\n"
"}";
char* sourceCode = malloc(16384);
- snprintf(sourceCode, 16384, sourceCodeTemplate,
(str[]){[TYPE_int] = "int"}[type], approximateOperation(predicateOperation),
approximationBits/8, 32-approximationBits);
- superverboseprintf("#%s, sourceCode: %s;\n", __func__,
sourceCode);
- cl_int err;
- cl_program program = clCreateProgramWithSource(getCLContext(),
1, (const char*[]){sourceCode}, (size_t[]){strlen(sourceCode)}, &err);
- if(err) THRprintf(GDKout, "#%s, clCreateProgramWithSource:
%s;\n", __func__, clError(err));
- err = clBuildProgram(program, 1, (const
cl_device_id[]){getDeviceID()}, "", NULL, NULL);
- if(err) {
- printf("#%s, clBuildProgram: %s;\n", __func__,
clError(err));
- size_t bufferSize;
- char* buffer;
- clGetProgramBuildInfo(program, getDeviceID(),
CL_PROGRAM_BUILD_LOG, 0, NULL, &bufferSize);
- clGetProgramBuildInfo(program, getDeviceID(),
CL_PROGRAM_BUILD_LOG, bufferSize, (buffer = alloca(bufferSize)), NULL);
- printf("#%s, clBuildProgram log: %s;\n", __func__,
buffer);
- }
+ snprintf(sourceCode, 16384, sourceCodeTemplate,
(str[]){[TYPE_int] = "int"}[type], approximateOperation(predicateOperation),
approximationBits/8, predicateOperation2 == NULL?1:0);
+ cl_program program = compileProgram(sourceCode,"");
free(sourceCode);
return program;
}
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
@@ -8,6 +8,7 @@
#endif
-cl_program getUSelectProgram(int type, char* predicateOperation, unsigned int
approximationBits);
+cl_program getUSelectProgram(int type, char* predicateOperation, char*
predicateOperation2, unsigned int approximationBits);
cl_program getProjectionLeftjoinProgram(unsigned int approximationBits);
+/* 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
@@ -39,7 +39,7 @@ clHead* getPositionsColumn(cl_mem memory
clGetMemObjectInfo(memoryObject, CL_MEM_SIZE, sizeof(size_t),
bufferSize, NULL);
else
err = clEnqueueReadBuffer(getCommandQueue(), memoryObject,
CL_TRUE, 0, *bufferSize, buffer, 0, NULL, NULL);
- if(err) THRprintf(GDKout, "#%s, clEnqueueReadBuffer: %s;\n", __func__,
clError(err));
+ if(err) printf("#%s, clEnqueueReadBuffer: %s;\n", __func__,
clError(err));
return buffer;
}
@@ -49,13 +49,11 @@ clTail* getApproximateValuesColumn(cl_me
clGetMemObjectInfo(memoryObject, CL_MEM_SIZE, sizeof(size_t),
bufferSize, NULL);
else
err = clEnqueueReadBuffer(getCommandQueue(), memoryObject,
CL_TRUE, 0, *bufferSize, buffer, 0, NULL, NULL);
- if(err) THRprintf(GDKout, "#%s, clEnqueueReadBuffer: %s;\n", __func__,
clError(err));
+ if(err) printf("#%s, clEnqueueReadBuffer: %s;\n", __func__,
clError(err));
return buffer;
}
str BWDLeftJoinApproximate(bat * res, bat * l, bat * r){
-
- ALGODEBUG printf("#BWDfetchjoin: approximating;\n");
BAT* right = BATdescriptor(*r);
if(!batTailIsDecomposed(right)){
printf("bwd.%s : bat is not decomposed: %d, no approximation
can be provided\n", __func__, *r);
@@ -65,7 +63,7 @@ str BWDLeftJoinApproximate(bat * res, ba
BAT* left = BATdescriptor(*l);
if(BAThvoid(left) && BATtvoid(left) && BAThvoid(right) &&
right->tseqbase != oid_nil){
-
+ // not sure if this is actually used (or even correct)
BAT* result = BATnew(ATOMtype(left->htype),
ATOMtype(right->ttype), left->batCount);
BATsetcount(result, left->batCount);
bzero(Tloc(result, BUNfirst(result)),
left->batCount*sizeof(int));
@@ -101,14 +99,12 @@ str BWDLeftJoinApproximate(bat * res, ba
cl_int err;
cl_kernel projectKernel = clCreateKernel(program, "project",
&err);
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
