✨ Add a function in texture.cpp to map texture array to an array of surface object

Author: davidAlgis

Plugin/PluginInteropUnityCUDA/include/Texture/texture.h

@@ -80,6 +80,8 @@ class Texture
     UNITY_INTERFACE_EXPORT void unMapTextureToSurfaceObject(
         cudaSurfaceObject_t &inputSurfObj);
 
+    UNITY_INTERFACE_EXPORT cudaSurfaceObject_t* mapTextureArrayToSurfaceObject();
+
     UNITY_INTERFACE_EXPORT cudaTextureObject_t mapTextureToTextureObject(
         int indexInArray = 0);
 

Plugin/PluginInteropUnityCUDA/src/Texture/texture.cpp

@@ -46,6 +46,33 @@ cudaSurfaceObject_t Texture::mapTextureToSurfaceObject(int indexInArray)
     return inputSurfObj;
 }
 
+cudaSurfaceObject_t* Texture::mapTextureArrayToSurfaceObject()
+{
+    cudaSurfaceObject_t* surf = new cudaSurfaceObject_t[_textureDepth];
+        Log::log().debugLog("map");
+        // map the resource to cuda
+        CUDA_CHECK(cudaGraphicsMapResources(1, &_pGraphicsResource));
+    for(int i=0; i<_textureDepth;i++)
+    {
+        // cuda array on which the resources will be sended
+        cudaArray *arrayPtr;
+        // https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__INTEROP.html#group__CUDART__INTEROP_1g0dd6b5f024dfdcff5c28a08ef9958031
+        CUDA_CHECK(cudaGraphicsSubResourceGetMappedArray(
+            &arrayPtr, _pGraphicsResource, i, 0));
+
+        // Wrap the cudaArray in a surface object
+        cudaResourceDesc resDesc;
+        memset(&resDesc, 0, sizeof(resDesc));
+        resDesc.resType = cudaResourceTypeArray;
+        resDesc.res.array.array = arrayPtr;
+        surf[i] = 0;
+        CUDA_CHECK(cudaCreateSurfaceObject(&surf[i], &resDesc));
+        CUDA_CHECK(cudaGetLastError());
+    }
+
+    return surf;
+}
+
 cudaTextureObject_t Texture::mapTextureToTextureObject(int indexInArray)
 {
     // map the resource to cuda

Plugin/SampleBasic/include/action_sample_texture_array.h

@@ -15,7 +15,7 @@ namespace SampleBasic {
 
 	private:
 		Texture* _texture;
-		cudaSurfaceObject_t _surf;
+		cudaSurfaceObject_t* _surf;
 	};
 } // namespace SampleBasic
 

Plugin/SampleBasic/src/action_sample_texture_array.cpp

@@ -10,21 +10,29 @@ ActionSampleTextureArray::ActionSampleTextureArray(void *texturePtr, int width,
     : Action()
 {
     _texture = CreateTextureInterop(texturePtr, width, height, depth);
-    _surf = 0;
+    _surf = new cudaSurfaceObject_t[depth];
+    for(int i=0; i<depth; i++)
+    {
+        _surf[i] = 0;
+    }
 }
 
 inline int ActionSampleTextureArray::Start()
 {
     _texture->registerTextureInCUDA();
-    _surf = _texture->mapTextureToSurfaceObject(1);
+    _surf = _texture->mapTextureArrayToSurfaceObject();
     return 0;
 }
 
 int ActionSampleTextureArray::Update()
 {
     kernelCallerWriteTextureArray(
         _texture->getDimGrid(), _texture->getDimBlock(),
-                             _surf, GetTime(), _texture->getWidth(),
+                             _surf[0], GetTime(), _texture->getWidth(),
+                             _texture->getHeight(), _texture->getDepth());
+    kernelCallerWriteTextureArray(
+        _texture->getDimGrid(), _texture->getDimBlock(),
+                             _surf[1], 2*GetTime(), _texture->getWidth(),
                              _texture->getHeight(), _texture->getDepth());
 
     cudaDeviceSynchronize();
@@ -33,8 +41,9 @@ int ActionSampleTextureArray::Update()
 
 inline int ActionSampleTextureArray::OnDestroy()
 {
-    _texture->unMapTextureToSurfaceObject(_surf);
+    _texture->unMapTextureToSurfaceObject(_surf[0]);
     _texture->unRegisterTextureInCUDA();
+    delete(_surf);
     return 0;
 }