rocm-systems/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMallocConcurrency.cpp

/*
Copyright (c) 2020-present Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/

/**
Testcase Scenarios :

 (TestCase 1)::
 1) Test hipMalloc() api passing zero size and confirming *ptr returning
 nullptr. Also pass nullptr to hipFree() api.
 2) Pass maximum value of size_t for hipMalloc() api and make sure appropriate
 error is returned.
 3) Check for hipMalloc() error code, passing invalid/null pointer.

 (TestCase 2)::
 4) Regress hipMalloc()/hipFree() in loop for bigger chunk of allocation
 with adequate number of iterations and later test for kernel execution on
 default gpu.
 5) Regress hipMalloc()/hipFree() in loop while allocating smaller chunks
 keeping maximum number of iterations and then run kernel code on default
 gpu, perfom data validation.

 (TestCase 3)::
 6) Check hipMalloc() api adaptability when app creates small chunks of memory
 continuously, stores it for later use and then frees it at later point
 of time.

 (TestCase 4)::
 7) Run hipMalloc() api/kernel code on same gpu parallely from parent and child
 processes, validate the results.

 (TestCase 5)::
 8) Execute hipMalloc() api simultaneously on all the gpus by spawning multiple
 child processes. Validate buffers allocated after running kernel code.

 (TestCase 6)::
 9) Multithread Scenario : Exercise hipMalloc() api parellely on all gpus from
 multiple threads and regress the api.

 (TestCases 2, 3, 4, 5, 6)::
 10) Validate memory usage with hipMemGetInfo() while regressing hipMalloc()
 api. Check for any possible memory leaks.
*/

/* HIT_START
 * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
 * TEST_NAMED: %t hipMalloc_ArgValidation  --tests 1
 * TEST_NAMED: %t hipMalloc_LoopRegression_AllocFreeCycle --tests 2
 * TEST_NAMED: %t hipMalloc_LoopRegression_AllocPool --tests 3
 * TEST_NAMED: %t hipMallocChild_Concurrency_DefaultGpu --tests 4
 * TEST_NAMED: %t hipMallocChild_Concurrency_MultiGpu --tests 5
 * TEST_NAMED: %t hipMalloc_MultiThreaded_MultiGpu --tests 6
 * HIT_END
 */

#include <sys/types.h>
#include <sys/wait.h>
#include <unistd.h>

#include <iostream>
#include <vector>
#include <limits>
#include <atomic>

#include "test_common.h"

/* Max alloc/free iterations for bigger chunks */
#define MAX_ALLOCFREE_BC (10000)

/* Buffer size for alloc/free cycles */
#define BUFF_SIZE_AF (5*1024*1024)

/* You may change it for individual test.
 * But default 100 is for quick return in Jenkin Build */
#define NUM_DIV 100

/* Max alloc/free iterations for smaller chunks */
#define MAX_ALLOCFREE_SC (5000000/NUM_DIV)

/* Max alloc and pool iterations (TBD) */
#define MAX_ALLOCPOOL_ITER (2000000/NUM_DIV)

/**
 * Validates data consitency on supplied gpu
 */
bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
  size_t Nbytes = N * sizeof(int);
  int *A_d, *B_d, *C_d;
  int *A_h, *B_h, *C_h;
  size_t prevAvl, prevTot, curAvl, curTot;
  bool TestPassed = true;

  HIPCHECK(hipSetDevice(gpu));
  HIPCHECK(hipMemGetInfo(&prevAvl, &prevTot));
  HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);

  unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);

  HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
  HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));

  hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
                  0, 0, static_cast<const int*>(A_d),
                  static_cast<const int*>(B_d), C_d, N);

  HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));

  if (!HipTest::checkVectorADD(A_h, B_h, C_h, N)) {
    printf("Validation PASSED for gpu %d from pid %d\n", gpu, getpid());
  } else {
    printf("%s : Validation FAILED for gpu %d from pid %d\n",
        __func__, gpu, getpid());
    TestPassed &= false;
  }

  HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
  HIPCHECK(hipMemGetInfo(&curAvl, &curTot));

  if (!concurOnOneGPU && (prevAvl != curAvl || prevTot != curTot)) {
    //In concurrent calls on one GPU, we cannot verify leaking in this way
    printf("%s : Memory allocation mismatch observed."
        "Possible memory leak.\n", __func__);
    TestPassed &= false;
  }

  return TestPassed;
}

/**
 * Fetches Gpu device count
 */
void getDeviceCount(int *pdevCnt) {
#ifdef __linux__
  int fd[2], val = 0;
  pid_t childpid;

  // create pipe descriptors
  pipe(fd);

  // disable visible_devices env from shell
  unsetenv("ROCR_VISIBLE_DEVICES");
  unsetenv("HIP_VISIBLE_DEVICES");

  childpid = fork();

  if (childpid > 0) {  // Parent
    close(fd[1]);
    // parent will wait to read the device cnt
    read(fd[0], &val, sizeof(val));

    // close the read-descriptor
    close(fd[0]);

    // wait for child exit
    wait(NULL);

    *pdevCnt = val;
  } else if (!childpid) {  // Child
    int devCnt = 1;
    // writing only, no need for read-descriptor
    close(fd[0]);

    HIPCHECK(hipGetDeviceCount(&devCnt));
    // send the value on the write-descriptor:
    write(fd[1], &devCnt, sizeof(devCnt));

    // close the write descriptor:
    close(fd[1]);
    exit(0);
  } else {  // failure
    *pdevCnt = 1;
    return;
  }

#else
  HIPCHECK(hipGetDeviceCount(pdevCnt));
#endif
}

/**
 * Regress memory allocation and free in loop
 */
bool regressAllocInLoop(int gpu) {
  bool TestPassed = true;
  size_t tot, avail, ptot, pavail;
  int i = 0;
  int *ptr;

  HIPCHECK(hipSetDevice(gpu));

  // Exercise allocation in loop with bigger chunks
  for (i = 0; i < MAX_ALLOCFREE_BC; i++) {
    size_t numBytes = BUFF_SIZE_AF;

    HIPCHECK(hipMemGetInfo(&pavail, &ptot));
    HIPCHECK(hipMalloc(&ptr, numBytes));
    HIPCHECK(hipMemGetInfo(&avail, &tot));

    if (pavail-avail < numBytes)  // We expect pavail-avail >= numBytes
    {
      printf("LoopAllocation %d : Memory allocation of %6.2fMB "
             "not matching with hipMemGetInfo - FAIL\n"
             "pavail=%zu, ptot=%zu, avail=%zu, tot=%zu, pavail-avail=%zu \n",
             i, numBytes/(1024.0*1024.0), pavail, ptot, avail, tot, pavail-avail);
      TestPassed &= false;
      HIPCHECK(hipFree(ptr));
      break;
    }

    HIPCHECK(hipFree(ptr));
  }

  // Exercise allocation in loop with smaller chunks and max iters
  HIPCHECK(hipMemGetInfo(&pavail, &ptot));

  for (i = 0; i < MAX_ALLOCFREE_SC; i++) {
    size_t numBytes = 16;

    HIPCHECK(hipMalloc(&ptr, numBytes));

    HIPCHECK(hipFree(ptr));
  }

  HIPCHECK(hipMemGetInfo(&avail, &tot));

  if ((pavail != avail) || (ptot != tot)) {
    printf("LoopAllocation : Memory allocation mismatch observed."
        "Possible memory leak.");
    TestPassed &= false;
  }

  return TestPassed;
}

/*
 * Thread func to regress alloc and check data consistency
 */

std::atomic<bool> g_thTestPassed(true);

void threadFunc(int gpu) {
  g_thTestPassed = g_thTestPassed & regressAllocInLoop(gpu);
  g_thTestPassed = g_thTestPassed & validateMemoryOnGPU(gpu);

  printf("thread execution status on gpu(%d) : %d\n", gpu, g_thTestPassed.load());
}

int main(int argc, char* argv[]) {
  HipTest::parseStandardArguments(argc, argv, true);

  if (p_tests == 1) {  // Arg validation
    // Test hipMalloc for zero size
    bool TestPassed = true;
    int *ptr;

    HIPCHECK(hipMalloc(&ptr, 0));

    // ptr expected to be reset to null ptr
    if (ptr) {
      printf("ArgValidation : Failed in zero size test\n");
      TestPassed &= false;
    }

    // Free null ptr
    HIPCHECK(hipFree(ptr));

    // Test hipMalloc for invalid arguments
    hipError_t ret;

    if ((ret = hipMalloc(NULL, 100)) != hipErrorInvalidValue) {
      printf("ArgValidation : Inappropritate error value returned"
          " for invalid argument. Error: '%s'(%d)\n",
          hipGetErrorString(ret), ret);
      TestPassed &= false;
    }

    // Test hipMalloc for Maximum value of size_t
    if ((ret = hipMalloc(&ptr, std::numeric_limits<std::size_t>::max()))
        != hipErrorMemoryAllocation) {
      printf("ArgValidation : Invalid error returned for max size_t."
          " Error: '%s'(%d)\n", hipGetErrorString(ret), ret);
      TestPassed &= false;
    }

    if (TestPassed) {
      passed();
    } else {
      failed("hipMalloc ArgumentValidation Failure!");
    }

  } else if (p_tests == 2) {  // Loop Regression Alloc/Free Cycle
    bool TestPassed = true;

    TestPassed &= regressAllocInLoop(0);
    TestPassed &= validateMemoryOnGPU(0);

    if (TestPassed) {
      passed();
    } else {
      failed("hipMalloc_LoopRegression_AllocFreeCycle Failure!");
    }

  } else if (p_tests == 3) {  // Loop Regression Alloc and Pool
    size_t avail, tot, pavail, ptot;
    bool TestPassed = true;
    hipError_t err;
    int *ptr;

    std::vector<int *> ptrlist;

    HIPCHECK(hipMemGetInfo(&pavail, &ptot));

    // Allocate small chunks of memory million times
    for (int i = 0; i < MAX_ALLOCPOOL_ITER; i++) {  // Iterations TBD
      if ((err = hipMalloc(&ptr, 10)) != hipSuccess) {
        HIPCHECK(hipMemGetInfo(&avail, &tot));

        printf("Loop regression pool allocation failure. "
        "Total gpu memory : %6.2fMB, Free memory %6.2fMB iter %d error '%s'\n",
        tot/(1024.0*1024.0), avail/(1024.0*1024.0), i, hipGetErrorString(err));

        TestPassed &= false;
        break;
      }

      // Store pointers allocated to emulate memory pool of app
      ptrlist.push_back(ptr);
    }

    // Free ptrs at later point of time
    for ( auto &t : ptrlist ) {
      HIPCHECK(hipFree(t));
    }

    HIPCHECK(hipMemGetInfo(&avail, &tot));

    TestPassed &= validateMemoryOnGPU(0);

    if ((pavail != avail) || (ptot != tot)) {
      printf("%s : Memory allocation mismatch observed. Possible memory leak.",
          __func__);
      TestPassed &= false;
    }

    if (TestPassed) {
      passed();
    } else {
      failed("hipMalloc_LoopRegression_AllocPool failure!");
    }

  } else if (p_tests == 4) {
    bool TestPassed = true;

#ifdef __linux__
    // Parallel execution of parent and child on gpu0
    int pid;

    if ((pid = fork()) < 0) {
      printf("Child_Concurrency_Gpu0 : fork() returned error %d.", pid);
      TestPassed &= false;

    } else if (!pid) {   // Child process
      bool TestPassedChild = true;
      TestPassedChild = validateMemoryOnGPU(0, true);

      if (TestPassedChild) {
        exit(0);  // child exit with success status
      } else {
        printf("Child_Concurrency_Gpu0 : childpid %d failed\n", getpid());
        exit(1);  // child exit with failure status
      }

    } else {  // Parent process
      int exitStatus;
      TestPassed = validateMemoryOnGPU(0, true);

      pid = wait(&exitStatus);
      if ( WEXITSTATUS(exitStatus) || ( pid < 0 ) )
        TestPassed &= false;
    }
#else
    printf("Test hipMallocChild_Concurrency_DefaultGpu skipped on non-linux\n");
#endif

    // TC scenarios specific to linux
    // are treated as pass in windows.
    if (TestPassed) {
      passed();
    } else {
      failed("hipMallocChild_Concurrency_DefaultGpu Failed!");
    }

  } else if (p_tests == 5) {
    bool TestPassed = true;
#ifdef __linux__
    // Parallel execution on multiple gpus from different child processes
    int devCnt = 1, pid = 0, cumStatus = 0;

    // Get GPU count
    getDeviceCount(&devCnt);

    // Spawn child for each GPU
    for (int gpu = 0; gpu < devCnt; gpu++) {
      if ((pid = fork()) < 0) {
         printf("Child_Concurrency_MultiGpu : fork() returned error %d\n", pid);
         failed("Test Failed!");

      } else if (!pid) {  // Child process
         bool TestPassedChild = true;
         TestPassedChild = validateMemoryOnGPU(gpu);

         if (TestPassedChild) {
            exit(0);  // child exit with success status
         } else {
            printf("Child_Concurrency_MultiGpu : childpid %d failed\n",
                getpid());
            exit(1);  // child exit with failure status
         }
      }
    }

    // Parent shall wait for child to complete
    for (int i = 0; i < devCnt; i++) {
      int pidwait = 0, exitStatus;
      pidwait = wait(&exitStatus);

      if (pidwait < 0) {
        TestPassed &= false;
        break;
      }

      cumStatus |= WEXITSTATUS(exitStatus);
    }

    // Cummulative status of all child
    if (cumStatus) {
       TestPassed &= false;
    }

#else
    printf("Test hipMallocChild_Concurrency_MultiGpu skipped on non-linux\n");
#endif


    // TC scenarios specific to linux
    // are treated as pass in windows.
    if (TestPassed) {
      passed();
    } else {
      failed("hipMallocChild_Concurrency_MultiGpu Failed!");
    }

  } else if (p_tests == 6) {  // Multithreaded multiple gpu execution
    std::vector<std::thread> threadlist;
    int devCnt = 1;

    // Get GPU count
    getDeviceCount(&devCnt);


    for (int i = 0; i < devCnt; i++) {
      threadlist.push_back(std::thread(threadFunc, i));
    }

    for (auto &t : threadlist) {
      t.join();
    }

    if (g_thTestPassed) {
      passed();
    } else {
      failed("hipMalloc_MultiThreaded_MultiGpu Failed!");
    }
  } else {
    failed("Didnt receive any valid option. Try options 1 to 6\n");
  }
}