Merge pull request #6 from A2R-Lab/older-gpu-compatibility

fixed compatibility issues with older GPUs

Author: EmreAdabag

README.md

@@ -6,18 +6,19 @@ Numerical experiments and the open-source solver from the paper ["MPCGPU: Real-T
 
 ```
 git clone https://github.com/A2R-Lab/MPCGPU
+cd MPCGPU
 git submodule update --init --recursive
 make build_qdldl
 make examples
 mkdir -p tmp/results
+```
+Either install the qdldl shared library by running ```cd qdldl/build && make install``` or modify the ```LD_LIBRARY_PATH``` environment variable to include the path to ```MPCGPU/qdldl/build/out```.
+
+```
 ./examples/pcg.exe
 ./examples/qdldl.exe
 ```
 
-The track_iiwa_pcg file outlines how you could setup a tracking problem with underlying linear system solver as pcg.
-
-The track_iiwa_qdldl file outlines how you could setup a tracking problem with underlying linear system solver as qdldl.
-
 ### Setting parameters
 
 You can set a bunch of parameters in `include/setting.cuh` file. You can also modify these by passing them as

include/common/merit.cuh

@@ -9,7 +9,8 @@
 template <typename T>
 size_t get_merit_smem_size(uint32_t state_size, uint32_t control_size)
 {
-    return sizeof(T) * ((4 * state_size + 2 * control_size ) + grid::EE_POS_SHARED_MEM_COUNT + max((2 * state_size + control_size), state_size + gato_plant::forwardDynamics_TempMemSize_Shared()));
+    return sizeof(T) * (6 + (2 * state_size + control_size ) + 
+                        ((int) 1.5 * state_size) + gato_plant::forwardDynamics_TempMemSize_Shared());
 }
 
 // cost compute for line search

include/dynamics/iiwa/iiwa_eepos_grid.cuh

@@ -98,7 +98,7 @@ namespace grid {
     const int NUM_EES = 1;
     const int ID_DYNAMIC_SHARED_MEM_COUNT = 770;
     const int MINV_DYNAMIC_SHARED_MEM_COUNT = 1395;
-    const int FD_DYNAMIC_SHARED_MEM_COUNT = 1444;
+    const int FD_DYNAMIC_SHARED_MEM_COUNT = 1724;
     const int ID_DU_DYNAMIC_SHARED_MEM_COUNT = 2450;
     const int FD_DU_DYNAMIC_SHARED_MEM_COUNT = 2450;
     const int ID_DU_MAX_SHARED_MEM_COUNT = 2695;

include/mpcsim.cuh

@@ -254,7 +254,7 @@ std::tuple<std::vector<toplevel_return_type>, std::vector<linsys_t>, linsys_t> s
 
 
 #if LIVE_PRINT_PATH
-        grid::end_effector_positions_kernel<T><<<1,128>>>(d_eePos, d_xs, grid::NUM_JOINTS, (grid::robotModel<T> *) d_dynmem, 1);
+        grid::end_effector_positions_kernel<T><<<1,128,144*sizeof(T)>>>(d_eePos, d_xs, grid::NUM_JOINTS, (grid::robotModel<T> *) d_dynmem, 1);
         gpuErrchk(cudaMemcpy(h_eePos, d_eePos, 6*sizeof(T), cudaMemcpyDeviceToHost));
         for (uint32_t i = 0; i < 6; i++){
             std::cout << h_eePos[i] << (i < 5 ? " " : "\n");
@@ -297,7 +297,7 @@ std::tuple<std::vector<toplevel_return_type>, std::vector<linsys_t>, linsys_t> s
         if (!shifted && time_since_timestep > shift_threshold){
 
             // record tracking error
-            grid::end_effector_positions_kernel<T><<<1,128>>>(d_eePos, d_xs, grid::NUM_JOINTS, (grid::robotModel<T> *) d_dynmem, 1);
+            grid::end_effector_positions_kernel<T><<<1,128,144*sizeof(T)>>>(d_eePos, d_xs, grid::NUM_JOINTS, (grid::robotModel<T> *) d_dynmem, 1);
             gpuErrchk(cudaMemcpy(h_eePos, d_eePos, 6*sizeof(T), cudaMemcpyDeviceToHost));
             gpuErrchk(cudaMemcpy(h_eePos_goal, d_eePos_goal, 6*sizeof(T), cudaMemcpyDeviceToHost));
             cur_tracking_error = 0.0;
@@ -401,7 +401,7 @@ std::tuple<std::vector<toplevel_return_type>, std::vector<linsys_t>, linsys_t> s
 #endif
 
 
-    grid::end_effector_positions_kernel<T><<<1,128>>>(d_eePos, d_xs, grid::NUM_JOINTS, (grid::robotModel<T> *) d_dynmem, 1);
+    grid::end_effector_positions_kernel<T><<<1,128,144*sizeof(T)>>>(d_eePos, d_xs, grid::NUM_JOINTS, (grid::robotModel<T> *) d_dynmem, 1);
     gpuErrchk(cudaMemcpy(h_eePos, d_eePos, 6*sizeof(T), cudaMemcpyDeviceToHost));
     gpuErrchk(cudaMemcpy(h_eePos_goal, d_eePos_goal, 6*sizeof(T), cudaMemcpyDeviceToHost));
     cur_tracking_error = 0.0;