aboutsummaryrefslogtreecommitdiff
path: root/sqlite3-select-opencl.c
diff options
context:
space:
mode:
Diffstat (limited to 'sqlite3-select-opencl.c')
-rw-r--r--sqlite3-select-opencl.c476
1 files changed, 476 insertions, 0 deletions
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;
+}
generated by cgit v1.2.3 (git 2.25.1) at 2024-06-26 15:53:55 +0000