diff -urpN JohnTheRipper-unstable-jumbo/src/cuda/pwsafe.cu JohnTheRipper-unstable-jumbo_multigpu/src/cuda/pwsafe.cu
--- JohnTheRipper-unstable-jumbo/src/cuda/pwsafe.cu	2013-01-15 16:24:32.000000000 +0000
+++ JohnTheRipper-unstable-jumbo_multigpu/src/cuda/pwsafe.cu	2013-01-16 19:01:46.807046489 +0000
@@ -10,8 +10,8 @@
 #include "../cuda_pwsafe.h"
 #include "cuda_common.cuh"
 
-#define PWSAFE_IN_SIZE (KEYS_PER_CRYPT * sizeof(pwsafe_pass))
-#define PWSAFE_OUT_SIZE (KEYS_PER_CRYPT * sizeof(pwsafe_hash))
+#define PWSAFE_IN_SIZE (KEYS_PER_GPU * sizeof(pwsafe_pass))
+#define PWSAFE_OUT_SIZE (KEYS_PER_GPU * sizeof(pwsafe_hash))
 #define PWSAFE_SALT_SIZE (sizeof(pwsafe_salt))
 
 __global__ void kernel_pwsafe(pwsafe_pass * in, pwsafe_salt * salt,
@@ -149,7 +149,11 @@ __global__ void kernel_pwsafe(pwsafe_pas
 extern "C" void gpu_pwpass(pwsafe_pass * host_in, pwsafe_salt * host_salt,
     pwsafe_hash * host_out)
 {
-        pwsafe_pass *cuda_pass = NULL;  ///passwords
+	unsigned int gpu=0;
+
+//#define _DEFAULT_
+#ifdef _DEFAULT_
+ pwsafe_pass *cuda_pass = NULL;  ///passwords
         pwsafe_salt *cuda_salt = NULL;  ///salt
         pwsafe_hash *cuda_hash = NULL;  ///hashes
 
@@ -175,5 +179,58 @@ extern "C" void gpu_pwpass(pwsafe_pass *
         cudaFree(cuda_pass);
         cudaFree(cuda_salt);
         cudaFree(cuda_hash);
+#else
+
+	//int runtimeVersion=0;
+	//cudaRuntimeGetVersion(&runtimeVersion);
+	//printf("Cuda runtime: %d.%d\n",runtimeVersion/1000,(runtimeVersion%100)/10); 	
+
+
+	//unsigned int TB=THREADS*BLOCKS;
+        pwsafe_pass *cuda_pass[GPUS];  ///passwords
+        pwsafe_salt *cuda_salt[GPUS];  ///salt
+        pwsafe_hash *cuda_hash[GPUS];  ///hashes
+
+	puts("stage 0");
+        ///Aloc memory and copy data to gpus
+	for(gpu=0;gpu<GPUS;gpu++){
+		HANDLE_ERROR(cudaSetDevice(gpu));
+        	HANDLE_ERROR(cudaMalloc(&cuda_pass[gpu], PWSAFE_IN_SIZE));
+        	HANDLE_ERROR(cudaMalloc(&cuda_salt[gpu], PWSAFE_SALT_SIZE));
+        	HANDLE_ERROR(cudaMalloc(&cuda_hash[gpu], PWSAFE_OUT_SIZE));
+
+		///Somehow this memset, which is not required, speeds things up a bit
+		HANDLE_ERROR(cudaMemset(cuda_hash[gpu], 0, PWSAFE_OUT_SIZE));
+        	HANDLE_ERROR(cudaMemcpy(cuda_pass[gpu], host_in+gpu*KEYS_PER_GPU, PWSAFE_IN_SIZE, cudaMemcpyHostToDevice));
+        	HANDLE_ERROR(cudaMemcpy(cuda_salt[gpu], host_salt, PWSAFE_SALT_SIZE,
+            		cudaMemcpyHostToDevice));
+	}
+	puts("stage 1");
+	for(gpu=0;gpu<GPUS;gpu++){
+		HANDLE_ERROR(cudaSetDevice(gpu));
+		printf("gpu=%d\n",gpu);
+        	///Run kernel and wait for execution end
+		pwsafe_pass *current_pass=cuda_pass[gpu];
+		pwsafe_salt *current_salt=cuda_salt[gpu];
+		pwsafe_hash *current_hash=cuda_hash[gpu];
+
+        	kernel_pwsafe <<< BLOCKS, THREADS >>> (current_pass, current_salt,
+            		current_hash);
+
+        	//cudaThreadSynchronize();
+		HANDLE_ERROR(cudaGetLastError());
+	}
+	puts("stage 2");
+	for(gpu=0;gpu<GPUS;gpu++){
+		HANDLE_ERROR(cudaSetDevice(gpu));
+        	///Free memory and copy results back
+        	HANDLE_ERROR(cudaMemcpy(host_out+KEYS_PER_GPU*gpu, cuda_hash[gpu], 				PWSAFE_OUT_SIZE,cudaMemcpyDeviceToHost));
+        	HANDLE_ERROR(cudaFree(cuda_pass[gpu]));
+        	HANDLE_ERROR(cudaFree(cuda_salt[gpu]));
+        	HANDLE_ERROR(cudaFree(cuda_hash[gpu]));
+	}
+	puts("stage 3");
+#endif
+
 }
 
diff -urpN JohnTheRipper-unstable-jumbo/src/cuda_pwsafe_fmt.c JohnTheRipper-unstable-jumbo_multigpu/src/cuda_pwsafe_fmt.c
--- JohnTheRipper-unstable-jumbo/src/cuda_pwsafe_fmt.c	2013-01-15 16:24:32.000000000 +0000
+++ JohnTheRipper-unstable-jumbo_multigpu/src/cuda_pwsafe_fmt.c	2013-01-16 18:49:14.315023253 +0000
@@ -20,6 +20,7 @@
 #include "base64.h"
 #include "memory.h"
 #include "cuda_pwsafe.h"
+#include "cuda_common.h"
 #define FORMAT_LABEL            "pwsafe-cuda"
 #define FORMAT_NAME             "Password Safe SHA-256"
 #define ALGORITHM_NAME          "CUDA"
@@ -55,7 +56,7 @@ static void init(struct fmt_main *self)
         host_pass = mem_calloc(KEYS_PER_CRYPT * sizeof(pwsafe_pass));
         host_hash = mem_calloc(KEYS_PER_CRYPT * sizeof(pwsafe_hash));
         host_salt = mem_calloc(sizeof(pwsafe_salt));
-
+	cuda_init(cuda_gpu_id);
         atexit(done);
 }
 
diff -urpN JohnTheRipper-unstable-jumbo/src/cuda_pwsafe.h JohnTheRipper-unstable-jumbo_multigpu/src/cuda_pwsafe.h
--- JohnTheRipper-unstable-jumbo/src/cuda_pwsafe.h	2013-01-15 16:24:32.000000000 +0000
+++ JohnTheRipper-unstable-jumbo_multigpu/src/cuda_pwsafe.h	2013-01-16 18:51:19.075027207 +0000
@@ -9,9 +9,11 @@
 #define uint8_t                         unsigned char
 #define uint32_t                        unsigned int
 
-#define THREADS                         512
-#define BLOCKS                          112//14 // 112 is good for gtx460
-#define KEYS_PER_CRYPT                  THREADS*BLOCKS
+#define GPUS				2
+#define THREADS                         256
+#define BLOCKS                          128//112//14 // 112 is good for gtx460
+#define KEYS_PER_GPU			THREADS*BLOCKS
+#define KEYS_PER_CRYPT                  KEYS_PER_GPU*GPUS
 
 #define MIN(a,b) (((a)<(b))?(a):(b))
 #define MAX(a,b) (((a)>(b))?(a):(b))
diff -urpN JohnTheRipper-unstable-jumbo/src/Makefile JohnTheRipper-unstable-jumbo_multigpu/src/Makefile
--- JohnTheRipper-unstable-jumbo/src/Makefile	2013-01-15 16:24:32.000000000 +0000
+++ JohnTheRipper-unstable-jumbo_multigpu/src/Makefile	2013-01-16 03:38:43.079127119 +0000
@@ -29,7 +29,7 @@ SORT = sort
 OMPFLAGS =
 # gcc with OpenMP
 #OMPFLAGS = -fopenmp
-#OMPFLAGS = -fopenmp -msse2
+OMPFLAGS = -fopenmp -msse2
 # Sun Studio with OpenMP (set the OMP_NUM_THREADS env var at runtime)
 #OMPFLAGS = -xopenmp
 # icc with OpenMP (for make target linux-x86-64-icc)
@@ -46,7 +46,7 @@ OMPFLAGS =
 
 # Change sm_10 to sm_20 if you have Fermi support (400 series or later) or
 # sm_30 if you have Kepler (600 series or later) and your nvcc supports it.
-NVCC_FLAGS = -c -Xptxas -v -arch sm_10
+NVCC_FLAGS = -c -Xptxas -v -arch sm_30
 CUDAPATH = /usr/local/cuda/lib
 CUDA64PATH = /usr/local/cuda/lib64