tAdded memtransfer and error checks - sphere - GPU-based 3D discrete element method algorithm with optional fluid coupling
 (HTM) git clone git://src.adamsgaard.dk/sphere
 (DIR) Log
 (DIR) Files
 (DIR) Refs
 (DIR) LICENSE
       ---
 (DIR) commit b6d92242b72333dd94520921dc43198ef420e907
 (DIR) parent c7f78b18b797dbca3702981e823f42c64a9bb18e
 (HTM) Author: Anders Damsgaard Christensen <adc@geo.au.dk>
       Date:   Fri,  9 Nov 2012 22:16:49 +0100
       
       Added memtransfer and error checks
       
       Diffstat:
         M src/raytracer.cuh                   |      11 ++++++++---
       
       1 file changed, 8 insertions(+), 3 deletions(-)
       ---
 (DIR) diff --git a/src/raytracer.cuh b/src/raytracer.cuh
       t@@ -414,6 +414,7 @@ __host__ void DEM::render(
          cudaThreadSynchronize();
        
          Float* linarr;     // Linear array to use for color visualization
       +  Float* dev_linarr; // Device linear array to use for color visualization
          std::string desc;  // Description of parameter visualized
          std::string unit;  // Unit of parameter values visualized
          unsigned int i;
       t@@ -427,7 +428,7 @@ __host__ void DEM::render(
          } else {
            
            if (method == 1) { // Visualize pressure
       -      linarr = dev_p;
       +      dev_linarr = dev_p;
              desc = "Pressure";
              unit = "Pa";
        
       t@@ -488,8 +489,11 @@ __host__ void DEM::render(
              << maxval << "] " << unit << endl;
        
            // Copy linarr to dev_linarr if required
       -    if (transfer == 1)
       -      cudaMemcpy(dev_linarr, &linarr, np*sizeof(Float), cudaMemcpyDeviceToHost);
       +    if (transfer == 1) {
       +      cudaMalloc((void**)&dev_linarr, np*sizeof(Float));
       +      cudaMemcpy(dev_linarr, &linarr, np*sizeof(Float), cudaMemcpyHostToDevice);
       +      checkForCudaErrors("Error during cudaMalloc or cudaMemcpy of linear array");
       +    }
        
            // Start raytracing kernel
            rayIntersectSpheresColormap<<< blocksPerGrid, threadsPerBlock >>>(
       t@@ -512,6 +516,7 @@ __host__ void DEM::render(
          // Free dynamically allocated global device memory
          rt_freeGlobalDeviceMemory();
          delete[] linarr;
       +  cudaFree(dev_linarr);
        
          //cudaPrintfDisplay(stdout, true);