MonetDB: bwd - * committing temporary fixes for multi-gpu suppor...

[Holger Pirk]

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
checkin-list@monetdb.org
https://www.monetdb.org/mailman/listinfo/checkin-list