Changeset: d3f51d4cb437 for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=d3f51d4cb437
Modified Files:
monetdb5/extras/bwd/Makefile.ag
monetdb5/extras/bwd/operations.c
Branch: bwd
Log Message:
* working on the selection (approximative selection is still too restrictive
but at leaste the approximate values find their way into the host memory)
* falling back to classic leftjoin for unimplemented cases
Unterschiede (208 Zeilen):
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
@@ -21,6 +21,7 @@ INCLUDES = \
../../../gdk \
../../../common/stream \
../../../common/options \
+ ../../modules/kernel \
$(OPENCL_CFLAGS)
MTSAFE
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
@@ -2,6 +2,8 @@
#include "monetdb_config.h"
#include "opt_statistics.h"
+#include "algebra.h"
+
#include "bwd.h"
#if defined(HAVE_OPENCL_OPENCL_H)
@@ -15,12 +17,18 @@
#pragma mark Actual MAL Operations Implementation
str BWDLeftJoinApproximate(bat * res, bat * l, bat * r){
+
ALGODEBUG THRprintf(GDKout, "#BWDfetchjoin: approximating;\n");
+ BAT* right = BATdescriptor(*r);
+ if(!batTailIsDecomposed(right)){
+ THRprintf(GDKout, "bwd.%s : bat is not decomposed: %d, no
approximation can be provided", __func__, *r);
+ BBPreleaseref(right->batCacheid);
+ return MAL_SUCCEED;
+ }
+
BAT* left = BATdescriptor(*l);
- BAT* right = BATdescriptor(*r);
if(BAThvoid(left) && BATtvoid(left) && BAThvoid(right) &&
right->tseqbase != oid_nil){
- if(!batTailIsDecomposed(right))
- throw (MAL, "bwd.BWDLeftJoinApproximate", "bat is not
decomposed: %d", *r);
+
BAT* result = BATnew(ATOMtype(left->htype),
ATOMtype(right->ttype), left->batCount);
BATsetcount(result, left->batCount);
bzero(Tloc(result, BUNfirst(result)),
left->batCount*sizeof(int));
@@ -46,11 +54,16 @@ str BWDLeftJoinApproximate(bat * res, ba
BBPreleaseref(right->batCacheid);
return MAL_SUCCEED;
} else
- throw(MAL, "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);
+ THRprintf(GDKout, "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;
};
str BWDLeftJoinRefine(bat * res, bat * l, bat * r, bat * approx){
+ if (!approx || !*approx) {
+ THRprintf(GDKout, "bwd.%s : no approximation provided, falling
back to normal leftjoin", __func__);
+ return ALGleftjoin(res, l, r);
+ }
ALGODEBUG THRprintf(GDKout, "#BWDfetchjoin: approximating;\n");
BAT* left = BATdescriptor(*l);
BAT* right = BATdescriptor(*r);
@@ -76,6 +89,32 @@ str BWDLeftJoinRefine(bat * res, bat * l
return MAL_SUCCEED;
};
+static inline size_t clDeviceAddressBytes(){
+ static size_t CL_DEVICE_ADDRESS_BYTES = 0;
+ if (!CL_DEVICE_ADDRESS_BYTES) {
+ cl_int err = clGetDeviceInfo(getDeviceID(),
CL_DEVICE_ADDRESS_BITS, sizeof(size_t), &CL_DEVICE_ADDRESS_BYTES, NULL);
+ if(err) THRprintf(GDKout, "#%s, %s;\n", __func__, clError(err));
+ CL_DEVICE_ADDRESS_BYTES /= 8;
+ }
+ return CL_DEVICE_ADDRESS_BYTES; //TODO: check
that CL_DEVICE_ADDRESS_BITS
+}
+
+typedef struct {
+ int count;
+ int padding;
+ char elements[];
+} clBAT;
+
+typedef union {
+ struct {
+ int index;
+ char value[];
+ } intIndexed;
+ struct {
+ int index;
+ char value[];
+ } longIndexed;
+} indexedValue;
str BWDThetauselectApproximate(bat *res, bat *bid, ptr val, str *OP) {
BAT* data = BATdescriptor(*bid);
@@ -93,23 +132,29 @@ str BWDThetauselectApproximate(bat *res,
slot->approximationBits =
batTailApproximationBits(data);
slot->residuals = NULL;
cl_int err;
- slot->approximation = clCreateBuffer(getCLContext(),
CL_MEM_READ_WRITE, (BATcount(data)+1)*(slot->approximationBits/8), NULL, &err);
+ slot->approximation = clCreateBuffer(getCLContext(),
CL_MEM_READ_WRITE, (BATcount(data)+1)*(clDeviceAddressBytes() +
slot->approximationBits/8)+sizeof(clBAT), NULL, &err);
+ void* zeros =
calloc((BATcount(data)+1)*(clDeviceAddressBytes() +
slot->approximationBits/8)+sizeof(clBAT), 1); //TODO: migrate buffer
initialization to GPU
+ clEnqueueWriteBuffer(getCommandQueue(),
slot->approximation, CL_TRUE, 0, (BATcount(data)+1)*(clDeviceAddressBytes() +
slot->approximationBits/8)+sizeof(clBAT), zeros, 0, NULL, NULL);
+ free(zeros);
if(err) THRprintf(GDKout, "#%s, clCreateBuffer: %s;\n",
__func__, clError(err));
}
-
-
const char* sourceCodeTemplate = "__kernel void uselect ("
- "__global void* output,"
+ "__global struct{int count; int padding; char
elements[];}* output,"
"__global const void* approximation,"
- "const %1$s operand,"
- "__global unsigned int* output_cursor"
+ "const %1$s operand"
") {"
- " printf(\"%%d\\n\", get_global_id(0));"
+ /* " printf(\"%%d:\", get_global_id(0));" */
" const %1$s value = (*(int*)(approximation +
get_global_id(0)*%3$d) << %4$d);"
" if(value %2$s operand)"
- " *(int*)(output + atomic_inc(output_cursor)*%3$d)
|= value;"
+ " {"
+ " const int slotNumber =
atomic_inc(&(output->count));"
+ " size_t* outputSlot = (size_t*)(output->elements +
slotNumber*(%3$d + sizeof(size_t)));"
+ " *outputSlot |= get_global_id(0);"
+ " outputSlot[1] |= (value >> %4$d);"
+ " printf(\"value: %%d, value in slot(%%d) %%d,
global_id: %%d\\n\", value, slotNumber, outputSlot[1], get_global_id(0));"
+ " }"
"}";
char* sourceCode = malloc(16384);
snprintf(sourceCode, 16384, sourceCodeTemplate,
typeNames[BATttype(data)], *OP, batTailApproximationBits(data)/8,
32-batTailApproximationBits(data));
@@ -134,14 +179,15 @@ str BWDThetauselectApproximate(bat *res,
err = clSetKernelArg(selectKernel, 1, sizeof(cl_mem),
(cl_mem[]){batTailApproximation(data)}); if (err) THRprintf(GDKout, "#%s,
clSetKernelArg(%d): %s;\n", __func__, 1, clError(err));
err = clSetKernelArg(selectKernel, 2, sizeof(int), val); if
(err) THRprintf(GDKout, "#%s, clSetKernelArg(%d): %s;\n", __func__, 2,
clError(err));
- cl_mem outputCursor;
- {
- outputCursor = clCreateBuffer(getCLContext(),
CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(int), (int[]){0}, &err);
- if(err) THRprintf(GDKout, "#%s, clCreateBuffer, cursor:
%s;\n", __func__, clError(err));
- }
+ /* cl_mem outputCursor; */
+ /* { */
+ /* outputCursor = clCreateBuffer(getCLContext(),
CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(int), (int[]){0}, &err); */
+ /* if(err) THRprintf(GDKout, "#%s, clCreateBuffer, cursor:
%s;\n", __func__, clError(err)); */
+ /* } */
- err = clSetKernelArg(selectKernel, 3, sizeof(cl_mem),
&outputCursor); if (err) THRprintf(GDKout, "#%s, clSetKernelArg(%d): %s;\n",
__func__, 3, clError(err));
- clEnqueueNDRangeKernel(getCommandQueue(), selectKernel, 1,
(const size_t[]){0}, (const size_t[]){BATcount(data)}, (const size_t[]){2}, 0,
NULL, NULL);
+ /* err = clSetKernelArg(selectKernel, 3, sizeof(cl_mem),
&outputCursor); if (err) THRprintf(GDKout, "#%s, clSetKernelArg(%d): %s;\n",
__func__, 3, clError(err)); */
+ err = clEnqueueNDRangeKernel(getCommandQueue(), selectKernel,
1, (const size_t[]){0}, (const size_t[]){BATcount(data)}, (const size_t[]){1},
0, NULL, NULL);
+ if(err) THRprintf(GDKout, "#%s, clEnqueueNDRangeKernel: %s;\n",
__func__, clError(err));
free(sourceCode);
BBPkeepref(*res = result->batCacheid);
}
@@ -149,6 +195,48 @@ str BWDThetauselectApproximate(bat *res,
return MAL_SUCCEED;
}
-str BWDThetauselectRefine(bat *result, bat *bid, ptr val, str *OP, bat
*approx) {
- return NULL;
+str BWDThetauselectRefine(bat *res, bat *bid, ptr val, str *OP, bat *approx) {
+ BAT* data = BATdescriptor(*bid);
+ assert(BATttype(data) == TYPE_int); // type specific
+
+ BAT* approximation = BATdescriptor(*approx);
+ assert(BATttype(approximation) == TYPE_int); // type specific
+
+ size_t approximationSize;
+ clGetMemObjectInfo(batTailApproximation(approximation), CL_MEM_SIZE,
sizeof(size_t), &approximationSize, NULL);
+
+
+ BAT* result = BATnew(BAThtype(data), BATttype(data),
data->batCount+(sizeof(clBAT)/Tsize(data)));
+ BATseqbase(result, 0);
+
+ clBAT* resultClBAT = (clBAT*) malloc(approximationSize);
+ /* Tloc(result, BUNfirst(result)); */
+
+ cl_int err = clEnqueueReadBuffer(getCommandQueue(),
batTailApproximation(approximation), CL_TRUE, 0, approximationSize,
resultClBAT, 0, NULL, NULL);
+ if(err) THRprintf(GDKout, "#%s, clEnqueueReadBuffer: %s;\n", __func__,
clError(err));
+ int* resultRegion = (int*) Tloc(result, BUNfirst(result)); // type
specific
+ oid* positionRegion = (oid*) Hloc(result, BUNfirst(result)); // type
specific
+ size_t candidateCount = resultClBAT->count;
+ int i = 0, j = 0;
+ /* const register unsigned int residualBits = 32-approximationBits; */
+ const register unsigned int approximationMask = ~((1 << (32 -
batTailApproximationBits(approximation)))-1);
+
+ while(i < candidateCount) {
+ const int offset = (clDeviceAddressBytes() +
batTailApproximationBits(approximation)/8)*i++;
+ /* printf ("offset: %d\n", offset); */
+ const int index = *(int*)&resultClBAT->elements[offset]; //
+ const int compressedValue =
*(int*)&(resultClBAT->elements[clDeviceAddressBytes()+offset]);
+ /* const int bm = ~((int)255); */
+ positionRegion[j] = index;
+ resultRegion[j++] = compressedValue << (32 -
batTailApproximationBits(approximation));
+ printf ("compressed value: %d, index: %d, value: %d\n",
compressedValue, index, resultRegion[j-1]);
+ }
+ BATsetcount(result, resultClBAT->count);
+ free(resultClBAT);
+
+
+ BBPkeepref(*res = result->batCacheid);
+ BBPreleaseref(*bid);
+ return MAL_SUCCEED;
}
+
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
