Changeset: a65a0bb38045 for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=a65a0bb38045
Modified Files:
monetdb5/extras/bwd/operations.c
monetdb5/extras/bwd/optimizer.c
monetdb5/extras/bwd/utilities.c
monetdb5/extras/bwd/utilities.h
Branch: bwd
Log Message:
* using fully decomposed bat-storage on the GPU as well
* this should make it easier to emulate the behaviour of monetdb operators
(markT, I'm looking at you)
* still not doing an approximation of a leftjoin
Unterschiede (gekürzt von 324 auf 300 Zeilen):
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
@@ -103,18 +103,13 @@ typedef struct {
int count;
int padding;
char elements[];
-} clBAT;
+} clTail;
-typedef union {
- struct {
- int index;
- char value[];
- } intIndexed;
- struct {
- int index;
- char value[];
- } longIndexed;
-} indexedValue;
+typedef struct {
+ int count;
+ int padding;
+ int positions[];
+} clHead;
char* approximateOperation(char* exactOperation){
if (exactOperation[1] == '\0')
@@ -131,9 +126,7 @@ str BWDThetauselectApproximate(bat *res,
if(!batTailIsDecomposed(data))
throw (MAL, "bwd problem", "while running %s, I noticed that
bat is not decomposed: %d", __func__, data->batCacheid);
{
- const char* typeNames[] = {[TYPE_int] = "int"};
- BAT* result = BATnew(BAThtype(data), BATttype(data),
data->batCount);
-
+ BAT* result = BATnew(BAThtype(data), TYPE_void, data->batCount);
{
const unsigned int newIndex =
getNextFreeDecomposedBATSlotIndex();
DecomposedBATSlot* slot =
getDecomposedBATSlot(newIndex);
@@ -141,16 +134,16 @@ 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)*(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);
+ slot->tailApproximation =
clCreateBuffer(getCLContext(), CL_MEM_READ_WRITE,
BATcount(data)*slot->approximationBits/8+sizeof(clTail), NULL, &err);
+ slot->headApproximation =
clCreateBuffer(getCLContext(), CL_MEM_READ_WRITE,
BATcount(data)*sizeof(int)+sizeof(clHead), NULL, &err);
+ clEnqueueWriteBuffer(getCommandQueue(),
slot->headApproximation, CL_TRUE, 0, sizeof(int), (int[1]){}, 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));
}
const char* sourceCodeTemplate = "__kernel void uselect (\n"
- "__global struct{int count; int padding; char
elements[];}* output,\n"
+ "__global struct{int count; int padding; int
positions[];}* outputHead,"
+ "__global struct{int count; int padding; char
values[];}* outputTail,\n"
"__global const char* approximation,\n"
"const %1$s operand\n"
") {\n"
@@ -159,35 +152,36 @@ str BWDThetauselectApproximate(bat *res,
" 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"
- " const size_t offset = atomic_inc(&(output->count))
* (sizeof(size_t) + %3$d);\n"
- " for(int i = 0; i < sizeof(size_t); i++)\n"
- " output->elements[offset+i] = (get_global_id(0)
>> (i*8));\n"
+ " if(1 || 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"
" for(int i = 0; i < %3$d; i++)\n"
- " output->elements[offset+sizeof(size_t)+i] =
approximation[inputOffset + i];\n"
+ " outputTail->values[offset+i] =
approximation[inputOffset + i];\n"
" }\n"
"}";
char* sourceCode = malloc(16384);
- snprintf(sourceCode, 16384, sourceCodeTemplate,
typeNames[BATttype(data)], approximateOperation(*OP),
batTailApproximationBits(data)/8, 32-batTailApproximationBits(data));
+ snprintf(sourceCode, 16384, sourceCodeTemplate,
(str[]){[TYPE_int] = "int"}[BATttype(data)], approximateOperation(*OP),
batTailApproximationBits(data)/8, 32-batTailApproximationBits(data));
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) {
- THRprintf(GDKout, "#%s, clBuildProgram: %s;\n",
__func__, clError(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);
- THRprintf(GDKout, "#%s, clBuildProgram log: %s;\n",
__func__, buffer);
+ printf("#%s, clBuildProgram log: %s;\n", __func__,
buffer);
}
cl_kernel selectKernel = clCreateKernel(program, "uselect",
&err);
if(err) THRprintf(GDKout, "#%s, clCreateKernel: %s;\n",
__func__, clError(err));
- err = clSetKernelArg(selectKernel, 0, sizeof(cl_mem),
(cl_mem[]){batTailApproximation(result)}); if (err) THRprintf(GDKout, "#%s,
clSetKernelArg(%d): %s;\n", __func__, 0, clError(err));
- 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),
(int[]){((*(int*)val)>>batTailResidualBits(data))<<batTailResidualBits(data)});
if (err) THRprintf(GDKout, "#%s, clSetKernelArg(%d): %s;\n", __func__, 2,
clError(err)); // type specific
+ err = clSetKernelArg(selectKernel, 0, sizeof(cl_mem),
(cl_mem[]){batHeadApproximation(result)}); if (err) THRprintf(GDKout, "#%s,
clSetKernelArg(%d): %s;\n", __func__, 0, clError(err));
+ err = clSetKernelArg(selectKernel, 1, sizeof(cl_mem),
(cl_mem[]){batTailApproximation(result)}); if (err) THRprintf(GDKout, "#%s,
clSetKernelArg(%d): %s;\n", __func__, 0, clError(err));
+ err = clSetKernelArg(selectKernel, 2, sizeof(cl_mem),
(cl_mem[]){batTailApproximation(data)}); if (err) THRprintf(GDKout, "#%s,
clSetKernelArg(%d): %s;\n", __func__, 1, clError(err));
+ err = clSetKernelArg(selectKernel, 3, sizeof(int),
(int[]){((*(int*)val)>>batTailResidualBits(data))<<batTailResidualBits(data)});
if (err) THRprintf(GDKout, "#%s, clSetKernelArg(%d): %s;\n", __func__, 2,
clError(err)); // type specific
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));
@@ -203,21 +197,31 @@ str BWDThetauselectRefine(bat *res, bat
assert(BATttype(data) == TYPE_int); // type specific
BAT* approximation = BATdescriptor(*approx);
- assert(BATttype(approximation) == TYPE_int); // type specific
+ assert(BATttype(data) == TYPE_int); // type specific
- size_t approximationSize;
- clGetMemObjectInfo(batTailApproximation(approximation), CL_MEM_SIZE,
sizeof(size_t), &approximationSize, NULL);
+ clHead* compressedHead;
+ {
+ size_t approximationSize;
+ clGetMemObjectInfo(batHeadApproximation(approximation),
CL_MEM_SIZE, sizeof(size_t), &approximationSize, NULL);
+ compressedHead = malloc(approximationSize);
+ cl_int err = clEnqueueReadBuffer(getCommandQueue(),
batHeadApproximation(approximation), CL_TRUE, 0, approximationSize,
compressedHead , 0, NULL, NULL);
+ if(err) THRprintf(GDKout, "#%s, clEnqueueReadBuffer: %s;\n",
__func__, clError(err));
+ }
+ clTail* compressedTail;
+ {
+ size_t approximationSize;
+ clGetMemObjectInfo(batTailApproximation(approximation),
CL_MEM_SIZE, sizeof(size_t), &approximationSize, NULL);
+ compressedTail = malloc(approximationSize);
+ cl_int err = clEnqueueReadBuffer(getCommandQueue(),
batTailApproximation(approximation), CL_TRUE, 0, approximationSize,
compressedTail , 0, NULL, NULL);
+ if(err) THRprintf(GDKout, "#%s, clEnqueueReadBuffer: %s;\n",
__func__, clError(err));
+ }
+
+ const size_t candidateCount = compressedHead->count;
+ BAT* result = BATnew(BAThtype(data), TYPE_void, candidateCount);
- BAT* result = BATnew(BAThtype(data), BATttype(data),
data->batCount+(sizeof(clBAT)/Tsize(data)));
- BATseqbase(result, 0);
-
- clBAT* resultClBAT = (clBAT*) malloc(approximationSize);
- 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
+ /* int* resultRegion = (int*) Tloc(result, BUNfirst(result)); // type
specific */
oid* positionRegion = (oid*) Hloc(result, BUNfirst(result)); // type
specific
- size_t candidateCount = resultClBAT->count;
const unsigned int approximationMask = ~((1 << (32 -
batTailApproximationBits(approximation)))-1);
const unsigned char* residuals = batTailResiduals(data);
const unsigned int residualMask = (1 << batTailResidualBits(data))-1;
@@ -225,13 +229,12 @@ str BWDThetauselectRefine(bat *res, bat
int i = 0, j = 0;
#define refineLoop(comparator) while(i < candidateCount) {
\
- const int offset = (clDeviceAddressBytes() +
batTailApproximationBits(approximation)/8)*i++; \
- const int index = *(int*)&resultClBAT->elements[offset];
\
- const int compressedValue =
*(int*)&(resultClBAT->elements[clDeviceAddressBytes()+offset]); \
+ const int index = compressedHead->positions[i];
\
+ const int offset =
(batTailApproximationBits(approximation)/8)*i++; \
+ const int compressedValue =
*(int*)&(compressedTail->elements[offset]); \
const int deCompressedValue = (compressedValue <<
batTailResidualBits(approximation)) + (*(int*)&residuals[index*residualBytes] &
residualMask); \
if(deCompressedValue comparator *(int*)val) {
\
- positionRegion[j] = index; \
- resultRegion[j++] = deCompressedValue; \
+ positionRegion[j++] = index; \
}\
}
switch (*OP[0]){
@@ -254,7 +257,9 @@ str BWDThetauselectRefine(bat *res, bat
}
#undef refineLoop
BATsetcount(result, j);
- free(resultClBAT);
+ BATseqbase(BATmirror(result), oid_nil);
+ free(compressedHead);
+ free(compressedTail);
BBPkeepref(*res = result->batCacheid);
diff --git a/monetdb5/extras/bwd/optimizer.c b/monetdb5/extras/bwd/optimizer.c
--- a/monetdb5/extras/bwd/optimizer.c
+++ b/monetdb5/extras/bwd/optimizer.c
@@ -1,5 +1,7 @@
#include "bwd.h"
#include "opt_prelude.h"
+#include "mal_type.h"
+#include "opt_statistics.h"
#pragma mark Utility Functions
@@ -21,25 +23,48 @@ static inline void rename_function(Instr
};
+
static inline int OPTBWDImplementation(Client client, MalBlkPtr malBlock,
MalStkPtr stack, InstrPtr instruction){
InstrPtr* oldProgram = malBlock->stmt;
size_t instructionCount = malBlock->stop;
-
+ size_t initialNumberOfVariables = malBlock->vtop;
+ int* approximations = alloca(initialNumberOfVariables*sizeof(int));
+ for (int i = 0; i < initialNumberOfVariables; ++i)
+ approximations[i] = 0;
+ for (int i = 0; i < instructionCount; ++i)
+ if(oldProgram[i]->modname == sqlRef && oldProgram[i]->fcnname
== bindRef)
+ approximations[getDestVar(oldProgram[i])] =
getDestVar(oldProgram[i]);
newMalBlkStmt(malBlock, malBlock->ssize + 1);
- int i;
- for (i = 0; i < instructionCount; ++i) {
+ for (int i = 0; i < instructionCount; ++i) {
+ ////////////////////////////// propagate approximations
//////////////////////////////
+ if((oldProgram[i]->modname == batRef && oldProgram[i]->fcnname
== reverseRef) ||
+ (oldProgram[i]->modname == algebraRef &&
oldProgram[i]->fcnname == markTRef)){
+ approximations[getDestVar(oldProgram[i])] =
newTmpVariable(malBlock, getDestType(malBlock, oldProgram[i]));
+ InstrPtr approximationReverse =
copyInstruction(oldProgram[i]);
+ setDestVar(approximationReverse,
approximations[getDestVar(oldProgram[i])]);
+ setArg(approximationReverse, 1,
approximations[getArg(oldProgram[i],1)]);
+ pushInstruction(malBlock, approximationReverse);
+ }
+
+
+ ////////////////////////////// rewrite operators
//////////////////////////////
if(match_function(oldProgram[i], "algebra", "thetauselect")
|| match_function(oldProgram[i], "algebra",
"leftjoin")) {
- int approximateResultVariable =
newTmpVariable(malBlock, getDestType(malBlock, oldProgram[i]));
+ approximations[getDestVar(oldProgram[i])] =
newTmpVariable(malBlock, getDestType(malBlock, oldProgram[i]));
{
InstrPtr approximateLeftjoin =
newInstruction(malBlock, ASSIGNsymbol);
char functionNameBuffer[512];
snprintf(functionNameBuffer, 512,
"%sapproximate", oldProgram[i]->fcnname);
rename_function(approximateLeftjoin, "bwd",
functionNameBuffer, (int[]){0}, client);
- setDestVar(approximateLeftjoin,
approximateResultVariable);
+ setDestVar(approximateLeftjoin,
approximations[getDestVar(oldProgram[i])]);
int j;
- for (j = 1; j < oldProgram[i]->argc; ++j)
- approximateLeftjoin =
pushArgument(malBlock, approximateLeftjoin, getArg(oldProgram[i], j));
+ for (j = 1; j < oldProgram[i]->argc; ++j) {
+ if(isaBatType(getArgType(malBlock,
oldProgram[i], j)))
+ approximateLeftjoin =
pushArgument(malBlock, approximateLeftjoin,
approximations[getArg(oldProgram[i], j)]);
+ else
+ approximateLeftjoin =
pushArgument(malBlock, approximateLeftjoin, getArg(oldProgram[i], j));
+ }
+
pushInstruction(malBlock, approximateLeftjoin);
}
{
@@ -51,12 +76,13 @@ static inline int OPTBWDImplementation(C
int j;
for (j = 1; j < oldProgram[i]->argc; ++j)
refineLeftjoin = pushArgument(malBlock,
refineLeftjoin, getArg(oldProgram[i], j));
- refineLeftjoin = pushArgument(malBlock,
refineLeftjoin, approximateResultVariable);
+ refineLeftjoin = pushArgument(malBlock,
refineLeftjoin, approximations[getDestVar(oldProgram[i])]);
pushInstruction(malBlock, refineLeftjoin);
}
} else {
pushInstruction(malBlock, oldProgram[i]);
}
+
}
printf ("rewriting plan for bwd\n");
return 0;
@@ -114,7 +140,7 @@ str OPTBWD(Client cntxt, MalBlkPtr mb, M
}
DEBUGoptimizers
mnstr_printf(cntxt->fdout,"#opt_reduce: " LLFMT " ms\n",t);
- QOTupdateStatistics("BWD",actions,t);
+ /* QOTupdateStatistics("BWD",actions,t); */
addtoMalBlkHistory(mb,"BWD");
return msg;
}
diff --git a/monetdb5/extras/bwd/utilities.c b/monetdb5/extras/bwd/utilities.c
--- a/monetdb5/extras/bwd/utilities.c
+++ b/monetdb5/extras/bwd/utilities.c
@@ -135,9 +135,17 @@ const cl_mem batTailApproximation(const
printf("#batTailApproximation: bat hasn't been
decomposed;\n");
return NULL;
}
- return
getDecomposedBATSlot(rightTailApproximationProperty->v.val.ival)->approximation;
+ return
getDecomposedBATSlot(rightTailApproximationProperty->v.val.ival)->tailApproximation;
}
+const cl_mem batHeadApproximation(const BAT* subject){
+ PROPrec* rightTailApproximationProperty;
+ if(!(rightTailApproximationProperty = BATgetprop(subject,
batRegistryIndex))){
+ printf("#batTailApproximation: bat hasn't been
decomposed;\n");
+ return NULL;
+ }
+ return
getDecomposedBATSlot(rightTailApproximationProperty->v.val.ival)->headApproximation;
+}
const size_t batTailApproximationBits(const BAT* subject){
PROPrec* rightTailApproximationProperty;
if(!(rightTailApproximationProperty = BATgetprop(subject,
batRegistryIndex))){
@@ -193,7 +201,7 @@ const unsigned int decomposeIntArray(con
}
{
cl_int err;
- slot->approximation = clCreateBuffer(getCLContext(),
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, (size+1)*(slot->approximationBits/8),
approximation, &err);
+ slot->tailApproximation = clCreateBuffer(getCLContext(),
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, (size+1)*(slot->approximationBits/8),
approximation, &err);
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
