✨ Make it works with texture array that are wrote at the same time

Author: davidAlgis

Plugin/PluginInteropUnityCUDA/include/Texture/texture.h

@@ -106,7 +106,7 @@ class Texture
     UNITY_INTERFACE_EXPORT void *getNativeTexturePtr() const;
 
     /**
-     * Get the pointer of _surfObjArray
+     * Get the pointer of d_surfObjArray
      * This array of surface object is necessary
      * to write or read into a texture
      */
@@ -137,9 +137,15 @@ class Texture
 
     private:
     // An array of surface object that will be of the size of texture depth
+    // This array is allocate on host side and will be copy to device memory
+    // when texture is map to it
+    cudaSurfaceObject_t *_surfObjArray;
+
+    // A device array of surface object that will be of the size of texture depth
+    // This array is allocate on device memory.
     // the surface object is the object that you can used to write into texture
     // from cuda api (eg. with surf2DWrite)
-    cudaSurfaceObject_t *_surfObjArray;
+    cudaSurfaceObject_t *d_surfObjArray;
 
     dim3 _dimBlock;
     dim3 _dimGrid;

Plugin/PluginInteropUnityCUDA/src/Texture/texture.cpp

@@ -29,6 +29,9 @@ Texture::Texture(void *textureHandle, int textureWidth, int textureHeight,
     // initialize surface object
     _surfObjArray = new cudaSurfaceObject_t[textureDepth];
 
+    CUDA_CHECK(cudaMalloc(&d_surfObjArray,
+                          _textureDepth * sizeof(cudaSurfaceObject_t)));
+
     for (int i = 0; i < textureDepth; i++)
     {
         _surfObjArray[i] = 0;
@@ -40,6 +43,7 @@ Texture::Texture(void *textureHandle, int textureWidth, int textureHeight,
 Texture::~Texture()
 {
     delete (_surfObjArray);
+    CUDA_CHECK(cudaFree(d_surfObjArray));
 }
 
 void Texture::mapTextureToSurfaceObject()
@@ -63,6 +67,9 @@ void Texture::mapTextureToSurfaceObject()
         CUDA_CHECK(cudaCreateSurfaceObject(&_surfObjArray[i], &resDesc));
         CUDA_CHECK(cudaGetLastError());
     }
+    CUDA_CHECK(cudaMemcpy(d_surfObjArray, _surfObjArray,
+                          _textureDepth * sizeof(cudaSurfaceObject_t),
+                          cudaMemcpyHostToDevice));
 }
 
 void Texture::unmapTextureToSurfaceObject()
@@ -108,7 +115,9 @@ void *Texture::getNativeTexturePtr() const
 
 cudaSurfaceObject_t *Texture::getSurfaceObjectArray() const
 {
-    return _surfObjArray;
+    // to use a complete array of surface object in a kernel,
+    // we need to use the array allocate on device memory
+    return d_surfObjArray;
 }
 
 cudaSurfaceObject_t Texture::getSurfaceObject(int indexInArray) const
@@ -121,5 +130,10 @@ cudaSurfaceObject_t Texture::getSurfaceObject(int indexInArray) const
         return 0;
     }
 
+    // to use a single surface object in a kernel
+    // we can use directly the surface object that 
+    // is on host side, because cudaSurfaceObject_t is a
+    // typename for unsigned long long which can be directly
+    // send to kernel as it's managed memory ? 
     return _surfObjArray[indexInArray];
 }

Plugin/SampleBasic/src/sample_kernels.cu

@@ -56,11 +56,7 @@ void kernelCallerWriteTextureArray(const dim3 dimGrid, const dim3 dimBlock,
                                    const float time, const int width,
                                    const int height, const int depth)
 {
-
-    writeTex<<<dimGrid, dimBlock>>>(surfObjArray[0], width, height, time);
-
-    writeTex<<<dimGrid, dimBlock>>>(surfObjArray[1], width, height, 2*time);
-    // writeTexArray<<<dimGrid, dimBlock>>>(surfObjArray, width, height, depth, time);
+    writeTexArray<<<dimGrid, dimBlock>>>(surfObjArray, width, height, depth, time);
 }
 
 void kernelCallerWriteBuffer(const dim3 dimGrid, const dim3 dimBlock,