reverse butter mempool and ssimu2 mempool to avoid massive vram leakage produced by python vapoursynth but also to have less impact on vram consumption inside a gpu filterchain

Author: Line-fr

src/butter/main.hpp

@@ -117,10 +117,19 @@ Plane_d getdiffmap(Plane_d* src1_d, Plane_d* src2_d, float* mem_d, int width, in
     return diffmap;
 }
 
-std::tuple<float, float, float> butterprocess(const uint8_t *dstp, int dststride, const uint8_t *srcp1[3], const uint8_t *srcp2[3], float* mem_d, float* pinned, GaussianHandle& gaussianHandle, int stride, int width, int height, float intensity_multiplier, int maxshared, hipStream_t stream){
+std::tuple<float, float, float> butterprocess(const uint8_t *dstp, int dststride, const uint8_t *srcp1[3], const uint8_t *srcp2[3], float* pinned, GaussianHandle& gaussianHandle, int stride, int width, int height, float intensity_multiplier, int maxshared, hipStream_t stream){
+    int wh = width*height;
+    const int totalscalesize = wh;
 
     //big memory allocation, we will try it multiple time if failed to save when too much threads are used
+    hipError_t erralloc;
 
+    const int totalplane = 31;
+    float* mem_d;
+    erralloc = hipMallocAsync(&mem_d, sizeof(float)*totalscalesize*(totalplane), stream); //2 base image and 6 working buffers
+    if (erralloc != hipSuccess){
+        throw VshipError(OutOfVRAM, __FILE__, __LINE__);
+    }
     //initial color planes
     Plane_d src1_d[3] = {Plane_d(mem_d, width, height, stream), Plane_d(mem_d+width*height, width, height, stream), Plane_d(mem_d+2*width*height, width, height, stream)};
     Plane_d src2_d[3] = {Plane_d(mem_d+3*width*height, width, height, stream), Plane_d(mem_d+4*width*height, width, height, stream), Plane_d(mem_d+5*width*height, width, height, stream)};
@@ -179,6 +188,8 @@ std::tuple<float, float, float> butterprocess(const uint8_t *dstp, int dststride
         throw e;
     }
 
+    hipFreeAsync(mem_d, stream);
+
     return finalres;
 }
 
@@ -187,7 +198,6 @@ typedef struct ButterData{
     VSNode *distorted;
     float intensity_multiplier;
     float** PinnedMemPool;
-    float** VRAMMemPool;
     GaussianHandle gaussianHandle;
     int maxshared;
     int diffmap;
@@ -236,9 +246,9 @@ static const VSFrame *VS_CC butterGetFrame(int n, int activationReason, void *in
         const int stream = d->streamSet->pop();
         try{
             if (d->diffmap){
-                val = butterprocess(vsapi->getWritePtr(dst, 0), vsapi->getStride(dst, 0), srcp1, srcp2, d->VRAMMemPool[stream], d->PinnedMemPool[stream], d->gaussianHandle, stride, width, height, d->intensity_multiplier, d->maxshared, d->streams[stream]);
+                val = butterprocess(vsapi->getWritePtr(dst, 0), vsapi->getStride(dst, 0), srcp1, srcp2, d->PinnedMemPool[stream], d->gaussianHandle, stride, width, height, d->intensity_multiplier, d->maxshared, d->streams[stream]);
             } else {
-                val = butterprocess(NULL, 0, srcp1, srcp2, d->VRAMMemPool[stream], d->PinnedMemPool[stream], d->gaussianHandle, stride, width, height, d->intensity_multiplier, d->maxshared, d->streams[stream]);
+                val = butterprocess(NULL, 0, srcp1, srcp2, d->PinnedMemPool[stream], d->gaussianHandle, stride, width, height, d->intensity_multiplier, d->maxshared, d->streams[stream]);
             }
         } catch (const VshipError& e){
             vsapi->setFilterError(e.getErrorMessage().c_str(), frameCtx);
@@ -272,11 +282,9 @@ static void VS_CC butterFree(void *instanceData, VSCore *core, const VSAPI *vsap
     vsapi->freeNode(d->distorted);
 
     for (int i = 0; i < d->streamnum; i++){
-        hipFree(d->VRAMMemPool[i]);
         hipHostFree(d->PinnedMemPool[i]);
         hipStreamDestroy(d->streams[i]);
     }
-    free(d->VRAMMemPool);
     free(d->PinnedMemPool);
     free(d->streams);
     d->gaussianHandle.destroy();
@@ -385,9 +393,7 @@ static void VS_CC butterCreate(const VSMap *in, VSMap *out, void *userData, VSCo
     d.streamSet = new threadSet(newstreamset);
 
     const int pinnedsize = allocsizeScore(viref->width, viref->height);
-    const int vramsize = viref->width*viref->height;
     d.PinnedMemPool = (float**)malloc(sizeof(float*)*d.streamnum);
-    d.VRAMMemPool = (float**)malloc(sizeof(float*)*d.streamnum);
     hipError_t erralloc;
     for (int i = 0; i < d.streamnum; i++){
         erralloc = hipHostMalloc(d.PinnedMemPool+i, sizeof(float)*pinnedsize);
@@ -397,13 +403,6 @@ static void VS_CC butterCreate(const VSMap *in, VSMap *out, void *userData, VSCo
             vsapi->freeNode(d.distorted);
             return;
         }
-        erralloc = hipMalloc(d.VRAMMemPool+i, sizeof(float)*31*vramsize);
-        if (erralloc != hipSuccess){
-            vsapi->mapSetError(out, VshipError(OutOfVRAM, __FILE__, __LINE__).getErrorMessage().c_str());
-            vsapi->freeNode(d.reference);
-            vsapi->freeNode(d.distorted);
-            return;
-        }
     }
 
     data = (ButterData *)malloc(sizeof(d));