Changeset: 7a0d1b7a2492 for MonetDB
URL: http://dev.monetdb.org/hg/MonetDB?cmd=changeset;node=7a0d1b7a2492
Modified Files:
monetdb5/extras/bwd/cl_program_utilities.c
monetdb5/extras/bwd/utilities.c
monetdb5/extras/bwd/utilities.h
Branch: bwd
Log Message:
* committing temporary fixes for multi-gpu support (need testing/verification)
Unterschiede (96 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
@@ -66,9 +66,9 @@ cl_program getMultiplyProgram(unsigned i
}
cl_program getProjectionLeftjoinProgram(unsigned int approximationBits,
unsigned int offsetBits){
- static cl_program programCache[4] = {};
-
- if(programCache[approximationBits/8-offsetBits/8] == NULL){
+ static cl_program programCache[4][8] = {}; //TODO: limited to four
devices
+ const int deviceForThisThread = getGPUDeviceForThisThread();
+ if(programCache[deviceForThisThread][approximationBits/8-offsetBits/8]
== NULL){
const char* sourceCode = "__kernel void project (\n"
"#if approximationBytes == 2\n"
"__global struct{int count; int base; unsigned short
values[];}* outputTail,\n"
@@ -98,9 +98,9 @@ cl_program getProjectionLeftjoinProgram(
"}";
char options[64];
snprintf(options, 64, "-D approximationBytes=%d",
approximationBits/8-offsetBits/8);
- programCache[approximationBits/8-offsetBits/8] =
compileProgram(sourceCode, options);
+
programCache[deviceForThisThread][approximationBits/8-offsetBits/8] =
compileProgram(sourceCode, options);
}
- return programCache[approximationBits/8-offsetBits/8];
+ return
programCache[deviceForThisThread][approximationBits/8-offsetBits/8];
}
cl_program getUSelectProgram(int type, char* predicateOperation, char*
predicateOperation2, unsigned int approximationBits, unsigned int offsetBits,
char inputIsVoidHeaded){
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
@@ -166,7 +166,7 @@ cl_command_queue getCommandQueue(void){
}
#pragma mark Decomposed BAT Handling
-#define MAX_DECOMPOSED_BATS 1024
+#define MAX_DECOMPOSED_BATS 4096
static DecomposedBATSlot bwdRegistry[MAX_DECOMPOSED_BATS] = {};
const char* batRegistryIndex = "bwd.batRegistryIndex";
@@ -182,7 +182,8 @@ const char* batRegistryIndex = "bwd.batR
if(0) printf("%s: bat hasn't been decomposed;\n",
__func__);
return NULL;
}
- return getDecomposedBATSlot(subject)->tailApproximation;
+ DecomposedBATSlot* slot = getDecomposedBATSlot(subject);
+ return
slot->isPersistentBAT?slot->tailApproximations[getGPUDeviceForThisThread()]:slot->tailApproximation;
}
@@ -313,12 +314,12 @@ unsigned int decomposeIntArray(const int
- const bounds subjectBounds = findBounds(subject,size);
- const int neededBits = ceil(log2(subjectBounds.max -
subjectBounds.min));
- slot->tailOffsetBits = 32-ceil(neededBits/8.0)*8;
- slot->approximationBits = approximationBits;
- slot->isPersistentBAT = 1;
- printf ("bounds: %d, %d, offset %zd bytes\n",subjectBounds.min,
subjectBounds.max, slot->tailOffsetBits/8);
+ const bounds subjectBounds = findBounds(subject,size);
+ const int neededBits = ceil(log2(subjectBounds.max -
subjectBounds.min));
+ slot->tailOffsetBits = 32-ceil(neededBits/8.0)*8;
+ slot->approximationBits = approximationBits;
+ slot->isPersistentBAT = 1;
+ printf ("bounds: %d, %d, offset %zd bytes\n",subjectBounds.min,
subjectBounds.max, slot->tailOffsetBits/8);
diff --git a/monetdb5/extras/bwd/utilities.h b/monetdb5/extras/bwd/utilities.h
--- a/monetdb5/extras/bwd/utilities.h
+++ b/monetdb5/extras/bwd/utilities.h
@@ -1,11 +1,12 @@
#ifndef _UTILITIES_H_
#define _UTILITIES_H_
-#define MAX_DECOMPOSED_BATS 1024
+#define MAX_DECOMPOSED_BATS 4096
typedef struct {
cl_mem headApproximation;
cl_mem tailApproximation;
+ cl_mem tailApproximations[4];
cl_mem tailPositions;
int isPersistentBAT;
char used;
@@ -42,7 +43,8 @@ unsigned int getNextFreeDecomposedBATSlo
size_t batTailOffsetBits(const BAT* subject) __attribute__((pure));
size_t batTailResidualBits(const BAT* subject) __attribute__((pure));
const unsigned char* batTailResiduals(const BAT*) __attribute__((pure));
- int isPersistentBAT(const BAT* subject) __attribute__((pure));
+int isPersistentBAT(const BAT* subject) __attribute__((pure));
+int getGPUDeviceForThisThread();
unsigned int decomposeIntArray(const int* subject, const size_t size, const
size_t approximationBits);
_______________________________________________
checkin-list mailing list
[email protected]
https://www.monetdb.org/mailman/listinfo/checkin-list
