Changeset: 1745879174b2 for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=1745879174b2
Modified Files:
monetdb5/extras/bwd/cl_program_utilities.c
monetdb5/extras/bwd/cl_program_utilities.h
monetdb5/extras/bwd/operations.c
Branch: bwd
Log Message:
* still working on projections
Unterschiede (162 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
@@ -23,9 +23,35 @@ char* approximateOperation(char* exactOp
static inline void superverboseprintf(const char * format, ... ){};
-/* type = BATttype(data) */
-/* predicateOperation = *OP */
-/* approximationBits = batTailApproximationBits(data) */
+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 getUSelectProgram(int type, char* predicateOperation, unsigned int
approximationBits){
const char* sourceCodeTemplate = "__kernel void uselect
(\n"
"__global struct{int count; int padding; int
positions[];}* outputHead,"
@@ -38,7 +64,7 @@ cl_program getUSelectProgram(int type, c
" for(int i = 0; i < %3$d; i++)\n"
" value += (approximation[inputOffset + i] <<
((i+sizeof(%1$s) - %3$d)*8));\n"
"\n"
- " if(1 || value %2$s operand){\n"
+ " if(value %2$s operand){\n"
" const int index =
atomic_inc(&(outputHead->count));"
" const int offset = index * %3$d;\n"
" outputHead->positions[index] = get_global_id(0);\n"
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
@@ -9,5 +9,5 @@
cl_program getUSelectProgram(int type, char* predicateOperation, unsigned int
approximationBits);
-
+cl_program getProjectionLeftjoinProgram(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
@@ -19,6 +19,20 @@ static const int activateWorkInProgress
#pragma mark Actual MAL Operations Implementation
+
+
+typedef struct {
+ int count;
+ int padding;
+ char elements[];
+} clTail;
+
+typedef struct {
+ int count;
+ int padding;
+ int positions[];
+} clHead;
+
str BWDLeftJoinApproximate(bat * res, bat * l, bat * r){
ALGODEBUG printf("#BWDfetchjoin: approximating;\n");
@@ -57,14 +71,55 @@ str BWDLeftJoinApproximate(bat * res, ba
BBPreleaseref(right->batCacheid);
return MAL_SUCCEED;
} else if(activateWorkInProgress && BAThvoid(left) &&
BAThvoid(right) && right->tseqbase != oid_nil){
- BAT* result = BATnew(ATOMtype(left->htype),
ATOMtype(right->ttype), 0);
cl_mem leftColumn = batTailApproximation(left);
if(!leftColumn) leftColumn =
batHeadApproximation(BATmirror(left));
cl_mem rightColumn = batTailApproximation(right);
if(!rightColumn) rightColumn =
batHeadApproximation(BATmirror(right));
+ cl_program program =
getProjectionLeftjoinProgram(batTailApproximationBits(right));
+
+
+
+ cl_int err;
+ cl_kernel projectKernel = clCreateKernel(program, "project",
&err);
+ if(err) THRprintf(GDKout, "#%s, clCreateKernel: %s;\n",
__func__, clError(err));
+
+
+ int headCount; // :-)
+ {
+ cl_int err = clEnqueueReadBuffer(getCommandQueue(),
leftColumn, CL_TRUE, 0, sizeof(int), &headCount , 0, NULL, NULL);
+ if(err) THRprintf(GDKout, "#%s, clEnqueueReadBuffer:
%s;\n", __func__, clError(err));
+ }
+
+ BAT* result = BATnew(TYPE_void, ATOMtype(right->ttype), 0);
+ {
+ const unsigned int newIndex =
getNextFreeDecomposedBATSlotIndex();
+ DecomposedBATSlot* slot =
getDecomposedBATSlot(newIndex);
+ BATsetprop(result, batRegistryIndex, TYPE_int,
(int[]){newIndex});
+ slot->approximationBits =
batTailApproximationBits(right);
+ slot->residuals = NULL;
+ cl_int err;
+ slot->tailApproximation =
clCreateBuffer(getCLContext(), CL_MEM_READ_WRITE,
headCount*slot->approximationBits/8+sizeof(clTail), NULL, &err);
+
+ clEnqueueWriteBuffer(getCommandQueue(),
slot->tailApproximation, CL_TRUE, 0, sizeof(int), &headCount, 0, NULL, NULL);
// I wonder what is faster transfering a single integer to the GPU or running a
kernel that initializes a value
+ if(err) printf("#%s, clCreateBuffer: %s;\n", __func__,
clError(err));
+
+ }
+
+ for (int i = 0; i < 3; ++i) {
+ if((err = clSetKernelArg(projectKernel, i,
sizeof(cl_mem), &((cl_mem[]){
+
batTailApproximation(result), leftColumn, rightColumn}[i]
+ )))) THRprintf(GDKout,
"#%s, clSetKernelArg(%d): %s;\n", __func__, 0, clError(err));
+ }
+
+
+
+
+
+ if((err = clEnqueueNDRangeKernel(getCommandQueue(),
projectKernel, 1, (const size_t[]){0}, (const size_t[]){headCount}, (const
size_t[]){1}, 0, NULL, NULL)))
+ THRprintf(GDKout, "#%s, clEnqueueNDRangeKernel: %s;\n",
__func__, clError(err));
+
- throw (MAL, "unimplemented", "%s (BAThvoid(left) &&
BAThvoid(right) && right->tseqbase != oid_nil), leftColumn: %p", __func__,
leftColumn);
} else
printf("bwd.BWDLeftJoinApproximate, %s (BAThvoid(left): %d,
BATtvoid(left): %d, BAThvoid(right): %d, right->tseqbase: %ld)", "this case
isn't implemented yet", BAThvoid(left), BATtvoid(left), BATtvoid(right),
left->tseqbase);
return MAL_SUCCEED;
@@ -111,18 +166,6 @@ static inline size_t clDeviceAddressByte
return CL_DEVICE_ADDRESS_BYTES; //TODO: check
that CL_DEVICE_ADDRESS_BITS
}
-typedef struct {
- int count;
- int padding;
- char elements[];
-} clTail;
-
-typedef struct {
- int count;
- int padding;
- int positions[];
-} clHead;
-
static inline void superverboseprintf(const char * format, ... ){};
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
