add gpu_factor for load_balance, adjust main logic

gpma_bfs_demo.cu

@@ -38,11 +38,8 @@ int main(int argc, char **argv) {
     char *data_path = argv[1];
     int bfs_start_node = std::atoi(argv[2]);
 
-#if CUDA_SM >= 60
-    // heap size limit is KNOWN to be required at SM_75(Tesla T4),SM_61(Tesla P4), and KNOWN to be forbidden at SM_50(GEForce 750).
-    cudaDeviceSetLimit(cudaLimitMallocHeapSize, 1024ll * 1024 * 1024);
+    cudaDeviceSetLimit(cudaLimitMallocHeapSize, 1024ll * 1024 * 700);
     cudaDeviceSetLimit(cudaLimitDevRuntimeSyncDepth, 5);
-#endif
 
     thrust::host_vector<int> host_x;
     thrust::host_vector<int> host_y;
@@ -64,7 +61,8 @@ int main(int argc, char **argv) {
     int step = half / num_slide;
 
     LOG_TIME("before init_csr_gpma")
-    GPMA_multidev<1, 1> gpma(node_size);
+    constexpr size_t cpu_count = 4;
+    GPMA_multidev<cpu_count-1, 1> gpma(node_size);
     cudaDeviceSynchronize();
 
     LOG_TIME("before update_gpma 1")

multidev.cuh

@@ -33,6 +33,7 @@ namespace gpma_impl {
             : mapKeyToSlot(hashSize, (size_t)(-1)) {}
 
         // void init(const CpuArrT &ptrs_cpu, const GpuArrT &ptrs_gpu) {}
+        static constexpr size_t gpu_factor = 7; // 1 GPU is equals to 7 CPU.
 
         // Given KEY, returns the ID(offset from zero) of device, which is responsible to this KEY.
         [[gnu::always_inline]] size_t select_device(const KEY_TYPE &k) {
@@ -41,7 +42,8 @@ namespace gpma_impl {
             auto dev_id = mapKeyToSlot.get(hashKey);
             if(dev_id == (size_t)(-1)) {
                 // appoint a device for a new hash.
-                dev_id = hashKey % (cpu_instances + gpu_instances);
+                dev_id = hashKey % (cpu_instances + gpu_instances * gpu_factor);
+                dev_id = (dev_id > cpu_instances) ? ( cpu_instances + (dev_id-cpu_instances)/gpu_factor ) : dev_id;
                 // Add link: hashKey => dev_id
                 return mapKeyToSlot.set(hashKey, dev_id);
             }