pin coords

Author: qsyao

cuda_backend/kernel/utils.cu

@@ -47,7 +47,9 @@ void Handle::set_2D(size_t y, size_t x){
     checkCudaErrors(cudaMallocHost((void **)&pin_output,
                             total_size * sizeof(float)));
 
-    checkCudaErrors(cudaMallocHost((void **)&coords,
+    checkCudaErrors(cudaMalloc((void **)&coords,
+                        2 * total_size * sizeof(float)));    
+    checkCudaErrors(cudaMallocHost((void **)&pin_coords,
                         2 * total_size * sizeof(float)));
 
     dim3 threads(min(total_size, (long)512), 1, 1);
@@ -81,9 +83,11 @@ void Handle::set_3D(size_t z, size_t y, size_t x){
     checkCudaErrors(cudaMallocHost((void **)&pin_output,
                             total_size * sizeof(float)));
 
-    checkCudaErrors(cudaMallocHost((void **)&coords,
+    checkCudaErrors(cudaMalloc((void **)&coords,
                         3 * total_size * sizeof(float)));
-                
+    checkCudaErrors(cudaMallocHost((void **)&pin_coords,
+                        3 * total_size * sizeof(float)));
+
     dim3 threads(min(total_size, (long)512), 1, 1);
     dim3 blocks(total_size/512 + 1, 1, 1);
     set_coords_3D<<<blocks, threads>>>(coords, dim_z, dim_y, dim_x);
@@ -110,19 +114,14 @@ void Handle::copy_output(float* ret){
 void Handle::check_coords(float* output){
     float* pin;
     if(is_3D){
-        checkCudaErrors(cudaMallocHost((void **)&pin,
-                           3 * total_size * sizeof(float)));
-        checkCudaErrors(cudaMemcpyAsync(pin, coords, 3 * total_size * sizeof(float),
+        checkCudaErrors(cudaMemcpyAsync(pin_coords, coords, 3 * total_size * sizeof(float),
                             cudaMemcpyDeviceToHost));
-        memcpy(output, pin, 3 * total_size * sizeof(float));
+        memcpy(output, pin_coords, 3 * total_size * sizeof(float));
     }
     else{
-        checkCudaErrors(cudaMallocHost((void **)&pin,
-                           2 * total_size * sizeof(float)));
-        checkCudaErrors(cudaMemcpyAsync(pin, coords, 2 * total_size * sizeof(float),
+        checkCudaErrors(cudaMemcpyAsync(pin_coords, coords, 2 * total_size * sizeof(float),
                             cudaMemcpyDeviceToHost));
-        memcpy(output, pin, 2 * total_size * sizeof(float));       
+        memcpy(output, pin_coords, 2 * total_size * sizeof(float));       
     }
-    checkCudaErrors(cudaFreeHost(pin));
 }
 

cuda_backend/kernel/utils.cuh

@@ -27,7 +27,7 @@ public:
     void copy_output(float* ret);
 
     void check_coords(float* coords);
-    
+
     ~Handle(){
         checkCudaErrors(cudaFree(img));
         checkCudaErrors(cudaFree(output));
@@ -46,6 +46,7 @@ private:
     float* pin_output;
 
     float* coords;
+    float* pin_coords;
 };
 
 }

python_augmentation.py

@@ -24,10 +24,12 @@ def create_zero_centered_coordinate_mesh(shape):
     test = output == array_image
     print(test)
 
-
-    coords = cuda_handle.check_coords()
-    cor_coords = create_zero_centered_coordinate_mesh(array_image.shape)
-    import ipdb; ipdb.set_trace()
+    start = time.time()
+    for i in range(Iters):
+        coords = cuda_handle.check_coords()
+    end = time.time()
+    print("Get Coords for Shape:{} Cost {}ms".format(array_image.shape, \
+                                    (end - start) * 1000 / Iters))    
 
     start = time.time()
     for i in range(Iters):