Add sqlite3-select-opencl.c which attempts to emit an OpenCL kernelcl-codegen

   in OpenCL source code that could perform the query.
   The emitter is agressive to bail out if it runs into something it
   doesn't understand from the sqlite parse tree.

Small modifications to opencl-sqlite.h to pick up a new datatype
   or two in support of this new function.

3 files changed, 507 insertions, 0 deletions

diff --git a/opencl-sqlite.h b/opencl-sqlite.h
index e7c2f1d..68a9ce3 100644
--- a/opencl-sqlite.h
+++ b/opencl-sqlite.h
@@ -163,5 +163,33 @@ int opencl_collect_timestamp(struct timespec *timestamp);
 
 void opencl_print_time_interval(struct timespec *start, struct timespec *end);
 
+/***************************************************************************
+ * CodeGen
+ */
+
+// Open CL memory type
+#define OPENCL_PRIVATE_MEMORY 1
+#define OPENCL_LOCAL_MEMORY 2
+#define OPENCL_CONSTANT_MEMORY 3
+#define OPENCL_GLOBAL_MEMORY 4
+
+// OpenCL data type
+#define SQCL_INT 1
+#define SQCL_INT2 2
+#define SQCL_INT3 3
+#define SQCL_INT4 4
+#define SQCL_FLOAT 5
+#define SQCL_FLOAT2 6
+#define SQCL_FLOAT3 7
+#define SQCL_FLOAT4 8
+#define SQCL_CHAR 10
+#define SQCL_CHAR2 11
+#define SQCL_CHAR3 12
+#define SQCL_CHAR4 13
+#define SQCL_CHAR8 14
+#define SQCL_CHAR16 15
+#define SQCL_INTPTR 16
+#define SQCL_FLOATPTR 17
+#define SQCL_CHARPTR 18
 
 #endif /* OPENCL_SQLITE_H */
diff --git a/sqlite3-select-opencl.c b/sqlite3-select-opencl.c
new file mode 100644
index 0000000..11ff14f
--- /dev/null
+++ b/sqlite3-select-opencl.c
@@ -0,0 +1,476 @@
+
+/* this module is included into the mass that has sqlite3.c before it
+ * as parse.h which is output by lemon is extremely important so that
+ * we can make use of the token definitions TK_*
+ * These define a variety of operators
+ */
+
+int stringToType(char *typeByString) {
+	if (strcmp("INTEGER",typeByString) == 0) 
+		return SQCL_INT;
+	else if (strcmp("REAL",typeByString) == 0)
+		return SQCL_FLOAT;
+/*	else if (strcmp("TEXT",typeByString) == 0)
+		return OPENCL_DOUBLE;
+	else if (strcmp("BLOB",typeByString) == 0)
+		return OPENCL_; */
+
+	return -1; // eh what?
+}
+
+/*
+** Generate code for the SELECT statement given in the p argument.  
+**
+** The results are distributed in various ways depending on the
+** contents of the SelectDest structure pointed to by argument pDest
+** as follows:
+**
+**     pDest->eDest    Result
+**     ------------    -------------------------------------------
+**     SRT_Output      Generate a row of output (using the OP_ResultRow
+**                     opcode) for each row in the result set.
+**
+**     SRT_Mem         Only valid if the result is a single column.
+**                     Store the first column of the first result row
+**                     in register pDest->iSDParm then abandon the rest
+**                     of the query.  This destination implies "LIMIT 1".
+**
+**     SRT_Set         The result must be a single column.  Store each
+**                     row of result as the key in table pDest->iSDParm. 
+**                     Apply the affinity pDest->affSdst before storing
+**                     results.  Used to implement "IN (SELECT ...)".
+**
+**     SRT_Union       Store results as a key in a temporary table 
+**                     identified by pDest->iSDParm.
+**
+**     SRT_Except      Remove results from the temporary table pDest->iSDParm.
+**
+**     SRT_Table       Store results in temporary table pDest->iSDParm.
+**                     This is like SRT_EphemTab except that the table
+**                     is assumed to already be open.
+**
+**     SRT_EphemTab    Create an temporary table pDest->iSDParm and store
+**                     the result there. The cursor is left open after
+**                     returning.  This is like SRT_Table except that
+**                     this destination uses OP_OpenEphemeral to create
+**                     the table first.
+**
+**     SRT_Coroutine   Generate a co-routine that returns a new row of
+**                     results each time it is invoked.  The entry point
+**                     of the co-routine is stored in register pDest->iSDParm.
+**
+**     SRT_Exists      Store a 1 in memory cell pDest->iSDParm if the result
+**                     set is not empty.
+**
+**     SRT_Discard     Throw the results away.  This is used by SELECT
+**                     statements within triggers whose only purpose is
+**                     the side-effects of functions.
+**
+** This routine returns the number of errors.  If any errors are
+** encountered, then an appropriate error message is left in
+** pParse->zErrMsg.
+**
+** This routine does NOT free the Select structure passed in.  The
+** calling function needs to do that.
+* 103174
+*/
+
+/* 
+pNew->pSrc = pSrc;
+  pNew->pWhere = pWhere;
+    pNew->pGroupBy = pGroupBy;
+	  pNew->pHaving = pHaving;
+	    pNew->pOrderBy = pOrderBy;
+		  pNew->selFlags = selFlags;
+		    pNew->op = TK_SELECT;
+			  pNew->pLimit = pLimit;
+			    pNew->pOffset = pOffset;
+				  assert( pOffset==0 || pLimit!=0 );
+				    pNew->addrOpenEphm[0] = -1;
+					  pNew->addrOpenEphm[1] = -1;
+					    pNew->addrOpenEphm[2] = -1;
+						*/
+/* For now output to a file */
+
+struct sqcl_variable {
+	int data_type;
+	int memory_type;
+	int columnNumber;
+	int used;
+	char name[20];
+};
+
+typedef struct sqcl_variable sqcl_variable;
+
+struct fnentrypoint {
+	sqcl_variable *totalRows;
+	sqcl_variable *parameters;
+	sqcl_variable *temps;
+	int parameter_count;
+	char name[20];
+	sqcl_variable *return_parameter;
+};
+
+typedef struct fnentrypoint fnentrypoint;
+
+// variable store
+
+int emit_cl_cl_memory_type(FILE *clOutputFile, int mem_type) {
+	switch(mem_type) {
+		case OPENCL_GLOBAL_MEMORY:
+			fprintf(clOutputFile,"__global ");
+			break;
+		case OPENCL_LOCAL_MEMORY:
+			fprintf(clOutputFile,"__local ");
+			break;
+		case OPENCL_CONSTANT_MEMORY:
+			fprintf(clOutputFile,"__constant ");
+			break;
+		case OPENCL_PRIVATE_MEMORY:
+		default:
+			fprintf(clOutputFile,"__private ");
+			break;
+	}
+	return 0;
+}
+
+int emit_cl_data_type(FILE *clOutputFile, int var_type) {
+	switch(var_type) {
+		case SQCL_INT:
+			fprintf(clOutputFile,"int ");
+			break;
+		case SQCL_INT4:
+			fprintf(clOutputFile,"int4 ");
+			break;
+		case SQCL_INTPTR:
+			fprintf(clOutputFile,"int * ");
+			break;
+		case SQCL_FLOAT:
+			fprintf(clOutputFile,"float ");
+			break;
+		case SQCL_FLOAT4:
+			fprintf(clOutputFile,"float4 ");
+			break;
+		case SQCL_FLOATPTR:
+			fprintf(clOutputFile,"float * ");
+			break;
+		default:
+			break;
+	}
+
+	return 0;
+}
+
+int emit_cl_vector_load(FILE *clOutputFile, sqcl_variable *in, sqcl_variable *intmp) {
+
+	fprintf(clOutputFile, "%s = vload4(0, %s + offset);\n", intmp->name, in->name );
+
+	return 0;
+}
+int emit_cl_vector_result_store(FILE *clOutputFile, sqcl_variable *tmp_result, sqcl_variable * result) {
+
+	fprintf(clOutputFile, "vstore4(%s, 0, %s + offset);\n", tmp_result->name, result->name);
+	return 0;
+}
+
+int descend_expr_tree(FILE *clOutputFile, Expr *scout) {
+	int err;
+	char buffer[64];
+	switch (scout-> op) {
+		case TK_AND:
+			fprintf(clOutputFile,"(");
+			sprintf(buffer," && ");
+			break;
+		case TK_OR:
+			fprintf(clOutputFile,"(");
+			sprintf(buffer," || ");
+			break;
+		case TK_NE:
+			fprintf(clOutputFile,"(");
+			sprintf(buffer," != ");
+			break;
+		case TK_EQ:
+			fprintf(clOutputFile,"(");
+			sprintf(buffer," == ");
+			break;
+		case TK_GT:
+			fprintf(clOutputFile,"(");
+			sprintf(buffer," > ");
+			break;
+		case TK_GE:
+			fprintf(clOutputFile,"(");
+			sprintf(buffer," >= ");
+			break;
+		case TK_LT:
+			fprintf(clOutputFile,"(");
+			sprintf(buffer," < ");
+			break;
+		case TK_LE:
+			fprintf(clOutputFile,"(");
+			sprintf(buffer," <= ");
+			break;
+		case TK_INTEGER:
+			sprintf(buffer," %d ", scout->u.iValue);
+			break;
+//		case TK_FLOAT: 
+//			fprintf(clOutputFile," %f ", scou);
+//			break;
+		case TK_COLUMN:
+			sprintf(buffer," %s ", scout->u.zToken);	
+			break;
+		default:
+			return -1;
+			break;
+	}
+	// process
+	if (scout->pLeft) {
+		err=descend_expr_tree(clOutputFile, scout->pLeft);
+		if (err) {
+			return err;
+		}
+	}
+
+	// emit
+	fprintf(clOutputFile, "%s",buffer);
+
+	if (scout->pRight) {
+		err = descend_expr_tree(clOutputFile, scout->pRight);
+		if (err) {
+			return err;
+		}
+	}
+
+	switch (scout-> op) {
+		case TK_AND:
+			fprintf(clOutputFile,")");
+			break;
+		case TK_OR:
+			fprintf(clOutputFile,")");
+			break;
+		case TK_NE:
+			fprintf(clOutputFile,")");
+			break;
+		case TK_EQ:
+			fprintf(clOutputFile,")");
+			break;
+		case TK_GT:
+			fprintf(clOutputFile,")");
+			break;
+		case TK_GE:
+			fprintf(clOutputFile,")");
+			break;
+		case TK_LT:
+			fprintf(clOutputFile,")");
+			break;
+		case TK_LE:
+			fprintf(clOutputFile,")");
+			break;
+		case TK_INTEGER:
+//		case TK_FLOAT: 
+//			fprintf(clOutputFile," %f ", scou);
+//			break;
+		case TK_COLUMN:
+		default:
+			return -1;
+			break;
+	}
+	
+	return 0;
+}
+
+int parse_and_emit_cl_operation(FILE *clOutputFile, Select *s,fnentrypoint *fn) {
+
+	Expr *scout;
+	int err;
+
+	fprintf(clOutputFile, "%s = ", fn->temps[fn->parameter_count].name);
+
+	// walk WHERE expression 
+	scout = s->pWhere;
+	err = descend_expr_tree(clOutputFile, scout);
+	fprintf(clOutputFile, ";\n");
+	
+	return err;
+}
+
+int emit_cl_fn(FILE *clOutputFile, fnentrypoint *fn, Select *s, SelectDest *sd) {
+	int i, err;
+
+	fprintf(clOutputFile, "__kernel void %s (", fn->name);
+
+	emit_cl_cl_memory_type(clOutputFile, fn->totalRows->memory_type);
+	emit_cl_data_type(clOutputFile, fn->totalRows->data_type);
+
+	fprintf(clOutputFile,"%s,", fn->totalRows->name);
+
+	sqcl_variable *q, *p = fn->parameters;
+
+	for (i=0; i < fn->parameter_count; i++) {
+
+		if (p->used) {
+			emit_cl_cl_memory_type(clOutputFile, p->memory_type);
+			emit_cl_data_type(clOutputFile, p->data_type);
+
+			fprintf(clOutputFile,"%s", p->name);
+		}
+		p++;
+		fprintf(clOutputFile,", ");
+	}
+
+	// last param is the resultMask
+	emit_cl_cl_memory_type(clOutputFile, fn->return_parameter->memory_type);
+	emit_cl_data_type(clOutputFile, fn->return_parameter->data_type);
+
+	fprintf(clOutputFile,"%s", fn->return_parameter->name);
+
+	fprintf(clOutputFile, ") {\n");
+
+
+	// emit boiler plate
+	fprintf(clOutputFile, "\n");
+
+ 	// the last temp is used for per row results
+	p = fn->temps;
+	for(i=0; i<=fn->parameter_count; i++ ) {
+		if (p->used) {
+			emit_cl_cl_memory_type(clOutputFile, p->memory_type);
+			emit_cl_data_type(clOutputFile, p->data_type);
+
+			fprintf(clOutputFile,"%s;\n", p->name);
+		}
+		p++;
+	}
+
+	// emit global id info
+	fprintf(clOutputFile,"int i = get_global_id(0);\n");
+	fprintf(clOutputFile,"size_t offset = i * (totalRows/workUnits);\n");
+
+	// start while
+	fprintf(clOutputFile,"do {\n");
+
+	// emit vector loads
+	q = fn->temps;
+	p = fn->parameters;
+
+	for(i=0; i<fn->parameter_count; i++ ) {
+		emit_cl_vector_load(clOutputFile, p, q);
+		p++;
+		q++;
+	}
+
+	// emit where test
+	err = parse_and_emit_cl_operation(clOutputFile, s, fn);
+	if (err)
+		return err;
+
+	// emit vector store
+	emit_cl_vector_result_store(clOutputFile, &(fn->temps[fn->parameter_count]), fn->return_parameter);
+
+	// emit decrement && vector walk
+	fprintf(clOutputFile,"\toffset+=4\n\ttotalRows--;\n");
+
+	// emit end while
+	fprintf(clOutputFile,"} while(totalRows);\n");
+
+	fprintf(clOutputFile, "}\n");
+
+	return 0;
+}
+
+int sqlite3OpenCLSelect(Parse *p, Select *s, SelectDest *sd) {
+
+	int i, err;
+	FILE *clOutputFile;
+	fnentrypoint mainFn;
+
+	if (sqlite3AuthCheck(p, SQLITE_SELECT, 0, 0, 0) ) return 1;
+
+	strcpy(mainFn.name, "x2_entry\0");
+
+	mainFn.parameters=NULL;
+	mainFn.parameter_count=0;
+	mainFn.return_parameter = NULL;	
+	mainFn.totalRows = malloc (sizeof(sqcl_variable));
+	mainFn.totalRows->data_type = SQCL_INT;
+	mainFn.totalRows->memory_type = OPENCL_GLOBAL_MEMORY;
+	mainFn.totalRows->columnNumber = -1;
+	mainFn.totalRows->used = 1;
+	strncpy(mainFn.totalRows->name, "totalRows\0",10);
+
+	// nExpr appears to be the number of vars used in the select
+	mainFn.parameters = malloc ((sizeof(sqcl_variable)) * (s->pEList->nExpr));
+	mainFn.temps = malloc ((sizeof(sqcl_variable)) * (s->pEList->nExpr +1) );
+
+	// last temp is for collecting results
+	mainFn.temps[s->pEList->nExpr].data_type = SQCL_INT4;
+	mainFn.temps[s->pEList->nExpr].memory_type = OPENCL_PRIVATE_MEMORY;
+	mainFn.temps[s->pEList->nExpr].columnNumber = -1;
+	mainFn.temps[s->pEList->nExpr].used = 1;
+	strncpy(mainFn.temps[s->pEList->nExpr].name, "_cl_r\0",6);
+
+
+	// return_parameter is a data structure that's passed in as a global
+	// that the results of the vector operation is stored into
+	mainFn.return_parameter = malloc (sizeof(sqcl_variable));
+	mainFn.return_parameter->data_type = SQCL_INTPTR;
+	mainFn.return_parameter->memory_type = OPENCL_GLOBAL_MEMORY;
+	mainFn.return_parameter->columnNumber = -1;
+	mainFn.return_parameter->used = 1;
+	strncpy(mainFn.return_parameter->name, "_cl_resultMask\0",15);
+
+	for (i = 0; i < s->pEList->nExpr; i++) {
+		if (s->pEList->a[i].pExpr->op == TK_COLUMN) {
+			mainFn.parameters[i].columnNumber = s->pEList->a[i].pExpr->iColumn;
+			mainFn.temps[i].columnNumber = s->pEList->a[i].pExpr->iColumn;
+			if (mainFn.parameters[i].columnNumber == -1 ) {
+				mainFn.parameters[i].columnNumber = 0;
+				mainFn.temps[i].columnNumber = 0;
+			}
+
+			mainFn.parameters[i].data_type = stringToType(s->pEList->a[i].pExpr->pTab->aCol[mainFn.parameters[i].columnNumber].zType);
+			switch (mainFn.parameters[i].data_type) {
+				case SQCL_INT:
+					mainFn.parameters[i].data_type = SQCL_INTPTR;
+					mainFn.temps[i].data_type = SQCL_INT4;
+					break;
+				case SQCL_FLOAT:
+					mainFn.parameters[i].data_type = SQCL_FLOATPTR;
+					mainFn.temps[i].data_type = SQCL_FLOAT4;
+					break;
+				case -1:
+					// data type we can't handle yet. bail!
+					goto cleanAndFail;
+			}
+			mainFn.parameters[i].memory_type = OPENCL_GLOBAL_MEMORY;
+			mainFn.temps[i].memory_type = OPENCL_PRIVATE_MEMORY;
+			strncpy(mainFn.parameters[i].name, s->pEList->a[i].pExpr->u.zToken,20);
+			sprintf(mainFn.temps[i].name, "_cl_v%d\0", i);
+			mainFn.parameters[i].name[19]='\0';
+			mainFn.parameter_count++;
+			mainFn.parameters[i].used = 1;
+			mainFn.temps[i].used = 1;
+		} else {
+			goto cleanAndFail;
+		}
+	}
+
+	clOutputFile=fopen("machinegen-opencl/foo.cl", "w");
+
+	err = emit_cl_fn(clOutputFile, &mainFn, s, sd);
+
+	fclose(clOutputFile);
+
+	free(mainFn.parameters);
+	free(mainFn.temps);
+	free(mainFn.return_parameter);
+	free(mainFn.totalRows);
+
+	return err;
+
+cleanAndFail:
+	free(mainFn.parameters);
+	free(mainFn.temps);
+	free(mainFn.return_parameter);
+	free(mainFn.totalRows);
+	return -1;
+}
diff --git a/sqlite3.c b/sqlite3.c
index bf81f04..91e0cbb 100644
--- a/sqlite3.c
+++ b/sqlite3.c
@@ -12230,6 +12230,7 @@ SQLITE_PRIVATE Index *sqlite3CreateIndex(Parse*,Token*,Token*,SrcList*,ExprList*
                           Expr*, int, int);
 SQLITE_PRIVATE void sqlite3DropIndex(Parse*, SrcList*, int);
 SQLITE_PRIVATE int sqlite3Select(Parse*, Select*, SelectDest*);
+SQLITE_PRIVATE int sqlite3OpenCLSelect(Parse*, Select*, SelectDest*);
 SQLITE_PRIVATE Select *sqlite3SelectNew(Parse*,ExprList*,SrcList*,Expr*,ExprList*,
                          Expr*,ExprList*,u16,Expr*,Expr*);
 SQLITE_PRIVATE void sqlite3SelectDelete(sqlite3*, Select*);
@@ -116405,6 +116406,7 @@ static void yy_reduce(
 {
   SelectDest dest = {SRT_Output, 0, 0, 0, 0};
   sqlite3Select(pParse, yymsp[0].minor.yy387, &dest);
+  sqlite3OpenCLSelect(pParse, yymsp[0].minor.yy387, &dest);
   sqlite3ExplainBegin(pParse->pVdbe);
   sqlite3ExplainSelect(pParse->pVdbe, yymsp[0].minor.yy387);
   sqlite3ExplainFinish(pParse->pVdbe);
@@ -144991,3 +144993,4 @@ SQLITE_PRIVATE void sqlite3Fts3IcuTokenizerModule(
 
 #include "vm-opencl.c"
 #include "transfer.c"
+#include "sqlite3-select-opencl.c"