diff --git a/projects/clr/hipamd/tests/src/runtimeApi/module/hipModuleLoadMultProcessOnMultGPU.cpp b/projects/clr/hipamd/tests/src/runtimeApi/module/hipModuleLoadMultProcessOnMultGPU.cpp
new file mode 100644
index 0000000000..7e8f591323
--- /dev/null
+++ b/projects/clr/hipamd/tests/src/runtimeApi/module/hipModuleLoadMultProcessOnMultGPU.cpp
@@ -0,0 +1,312 @@
+/*
+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 WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS 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 IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/* HIT_START
+ * BUILD_CMD: kernel_composite_test.code %hc --genco %S/kernel_composite_test.cpp -o kernel_composite_test.code
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
+ * TEST: %t --tests 0x1
+ * TEST: %t --tests 0x2
+ * TEST: %t --tests 0x3
+ * HIT_END
+ */
+#include <stdio.h>
+#include <stdlib.h>
+
+#ifdef __linux__
+#include <unistd.h>
+#include <sys/wait.h>
+#endif
+#include <iostream>
+#include <fstream>
+#include <cstddef>
+#include <vector>
+#include "test_common.h"
+
+#define TEST_ITERATIONS 1000
+#define CODEOBJ_FILE "kernel_composite_test.code"
+#define CODEOBJ_GLOB_KERNEL1 "testWeightedCopy"
+#define CODEOBJ_GLOB_KERNEL2 "getAvg"
+#define BLOCKSPERCULDULD 6
+#define THREADSPERBLOCKLDULD 256
+
+unsigned int globTestID = 0;
+
+/**
+ * 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
+}
+
+/**
+ * Validates hipModuleLoadUnload if globTestID = 1
+ * Validates hipModuleLoadDataUnload if globTestID = 2
+ * Validates hipModuleLoadDataExUnload if globTestID = 3
+ */
+bool testhipModuleLoadUnloadFunc(const std::vector<char>& buffer) {
+  size_t N = 16*16;
+  size_t Nbytes = N * sizeof(int);
+  int *A_d, *B_d;
+  int *A_h, *B_h;
+  unsigned blocks = HipTest::setNumBlocks(BLOCKSPERCULDULD,
+                             THREADSPERBLOCKLDULD, N);
+  int deviceid;
+  hipGetDevice(&deviceid);
+  printf("pid = %u deviceid = %d\n", getpid(), deviceid);
+  // allocate host and device buffer
+  HIPCHECK(hipMalloc(&A_d, Nbytes));
+  HIPCHECK(hipMalloc(&B_d, Nbytes));
+
+  A_h = reinterpret_cast<int *>(malloc(Nbytes));
+  if (NULL == A_h) {
+    failed("Failed to allocate using malloc");
+  }
+  B_h = reinterpret_cast<int *>(malloc(Nbytes));
+  if (NULL == B_h) {
+    failed("Failed to allocate using malloc");
+  }
+  // set host buffers
+  for (int idx = 0; idx < N; idx++) {
+    A_h[idx] = deviceid;
+  }
+  // Copy buffer from host to device
+
+  HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
+
+  hipModule_t Module;
+  hipFunction_t Function;
+  if (1 == globTestID) {
+    HIPCHECK(hipModuleLoad(&Module, CODEOBJ_FILE));
+  } else if (2 == globTestID) {
+    HIPCHECK(hipModuleLoadData(&Module, &buffer[0]));
+  } else if (3 == globTestID) {
+    HIPCHECK(hipModuleLoadDataEx(&Module,
+            &buffer[0], 0, nullptr, nullptr));
+  }
+  HIPCHECK(hipModuleGetFunction(&Function, Module,
+                               CODEOBJ_GLOB_KERNEL1));
+  float deviceGlobalFloatH = 3.14;
+  int   deviceGlobalInt1H = 100*deviceid;
+  int   deviceGlobalInt2H = 50*deviceid;
+  short deviceGlobalShortH = 25*deviceid;
+  char  deviceGlobalCharH = 13*deviceid;
+  hipDeviceptr_t deviceGlobal;
+  size_t deviceGlobalSize;
+  HIPCHECK(hipModuleGetGlobal(&deviceGlobal,
+           &deviceGlobalSize,
+           Module, "deviceGlobalFloat"));
+  HIPCHECK(hipMemcpyHtoD(hipDeviceptr_t(deviceGlobal),
+           &deviceGlobalFloatH,
+           deviceGlobalSize));
+  HIPCHECK(hipModuleGetGlobal(&deviceGlobal,
+           &deviceGlobalSize,
+           Module, "deviceGlobalInt1"));
+  HIPCHECK(hipMemcpyHtoD(hipDeviceptr_t(deviceGlobal),
+           &deviceGlobalInt1H,
+           deviceGlobalSize));
+  HIPCHECK(hipModuleGetGlobal(&deviceGlobal,
+           &deviceGlobalSize,
+           Module,
+           "deviceGlobalInt2"));
+  HIPCHECK(hipMemcpyHtoD(hipDeviceptr_t(deviceGlobal),
+           &deviceGlobalInt2H, deviceGlobalSize));
+  HIPCHECK(hipModuleGetGlobal(&deviceGlobal,
+           &deviceGlobalSize,
+           Module, "deviceGlobalShort"));
+  HIPCHECK(hipMemcpyHtoD(hipDeviceptr_t(deviceGlobal),
+           &deviceGlobalShortH, deviceGlobalSize));
+  HIPCHECK(hipModuleGetGlobal(&deviceGlobal,
+           &deviceGlobalSize, Module, "deviceGlobalChar"));
+  HIPCHECK(hipMemcpyHtoD(hipDeviceptr_t(deviceGlobal),
+           &deviceGlobalCharH, deviceGlobalSize));
+  // Launch Function kernel function
+
+  hipStream_t stream;
+  HIPCHECK(hipStreamCreate(&stream));
+
+  struct {
+    void* _Ad;
+    void* _Bd;
+  } args;
+  args._Ad = reinterpret_cast<void*>(A_d);
+  args._Bd = reinterpret_cast<void*>(B_d);
+  size_t size = sizeof(args);
+
+  void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
+                    HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
+                    HIP_LAUNCH_PARAM_END};
+  HIPCHECK(hipModuleLaunchKernel(Function, 1, 1, 1,
+            N, 1, 1, 0, stream, NULL,
+            reinterpret_cast<void**>(&config)));
+  // Copy buffer from decice to host
+  HIPCHECK(hipMemcpyAsync(B_h, B_d, Nbytes, hipMemcpyDeviceToHost, stream));
+  HIPCHECK(hipDeviceSynchronize());
+  HIPCHECK(hipStreamDestroy(stream));
+
+  // Check the results
+  for (int idx = 0; idx < N; idx++) {
+    if (B_h[idx] != (deviceGlobalInt1H*A_h[idx]
+            + deviceGlobalInt2H
+            + static_cast<int>(deviceGlobalShortH) +
+            + static_cast<int>(deviceGlobalCharH)
+            + static_cast<int>(deviceGlobalFloatH*deviceGlobalFloatH))) {
+        printf("Matrix Addition Failed\n");
+        // exit the current process with failure
+        return false;
+    }
+  }
+  HIPCHECK(hipModuleUnload(Module));
+  // free memory
+  HIPCHECK(hipFree(B_d));
+  HIPCHECK(hipFree(A_d));
+  free(B_h);
+  free(A_h);
+  printf("pid:%u PASSED\n", getpid());
+  return true;
+}
+
+/**
+ * Spawn 1 Process for each device
+ *
+ */
+void spawnProc(int deviceCount, const std::vector<char>& buffer) {
+  int numDevices = deviceCount;
+  bool TestPassed = true;
+#ifdef __linux__
+  pid_t pid = 0;
+  // spawn a process for each device
+  for (int deviceNo = 0; deviceNo < numDevices; deviceNo++) {
+    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;
+      // set the device id for the current process
+      HIPCHECK(hipSetDevice(deviceNo));
+      TestPassedChild = testhipModuleLoadUnloadFunc(buffer);
+
+      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
+      }
+    }
+  }
+  int cumStatus = 0;
+  // Parent shall wait for child to complete
+  for (int i = 0; i < numDevices; i++) {
+    int pidwait = 0, exitStatus;
+    pidwait = wait(&exitStatus);
+    cumStatus |= WEXITSTATUS(exitStatus);
+  }
+  if (cumStatus) {
+    TestPassed &= false;
+  }
+#else
+  for (int deviceNo = 0; deviceNo < numDevices; deviceNo++) {
+    // set the device id for the current process
+    HIPCHECK(hipSetDevice(deviceNo));
+    TestPassed &= testhipModuleLoadUnloadFunc(buffer);
+  }
+#endif
+  if (TestPassed) {
+    passed();
+  } else {
+    failed("hipMallocChild_Concurrency_MultiGpu Failed!");
+  }
+}
+
+int main(int argc, char* argv[]) {
+  HipTest::parseStandardArguments(argc, argv, true);
+  int numDevices = 0;
+  getDeviceCount(&numDevices);
+  if (1 == numDevices) {
+    printf("Testing on Single GPU machine.\n");
+  }
+  std::ifstream file(CODEOBJ_FILE,
+                    std::ios::binary | std::ios::ate);
+  std::streamsize fsize = file.tellg();
+  file.seekg(0, std::ios::beg);
+
+  std::vector<char> buffer(fsize);
+  if (!file.read(buffer.data(), fsize)) {
+    failed("could not open code object '%s'\n", CODEOBJ_FILE);
+  }
+  file.close();
+  if (p_tests == 0x1) {
+    globTestID = 1;
+    spawnProc(numDevices, buffer);
+  } else if (p_tests == 0x2) {
+    globTestID = 2;
+    spawnProc(numDevices, buffer);
+  } else if (p_tests == 0x3) {
+    globTestID = 3;
+    spawnProc(numDevices, buffer);
+  }
+}
diff --git a/projects/clr/hipamd/tests/src/runtimeApi/module/hipModuleLoadUnloadStress.cpp b/projects/clr/hipamd/tests/src/runtimeApi/module/hipModuleLoadUnloadStress.cpp
new file mode 100644
index 0000000000..bc72c5b3b8
--- /dev/null
+++ b/projects/clr/hipamd/tests/src/runtimeApi/module/hipModuleLoadUnloadStress.cpp
@@ -0,0 +1,119 @@
+/*
+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 WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS 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 IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
+ * TEST: %t --tests 0x1
+ * TEST: %t --tests 0x2
+ * TEST: %t --tests 0x3
+ * HIT_END
+ */
+#include <stdio.h>
+#include <stdlib.h>
+#include <iostream>
+#include <fstream>
+#include <cstddef>
+#include <vector>
+#include "test_common.h"
+
+#define TEST_ITERATIONS 1000
+#define CODEOBJ_FILE "kernel_composite_test.code"
+/**
+ * Run Valgrind tool with these test cases to validate memory leakage.
+ * E.g. valgrind --leak-check=yes ./a.out --tests 0x1
+ */
+
+/**
+ * Internal Function
+ */
+std::vector<char> load_file() {
+  std::ifstream file(CODEOBJ_FILE, std::ios::binary | std::ios::ate);
+  std::streamsize fsize = file.tellg();
+  file.seekg(0, std::ios::beg);
+  std::vector<char> buffer(fsize);
+  if (!file.read(buffer.data(), fsize)) {
+    failed("could not open code object '%s'\n", CODEOBJ_FILE);
+  }
+  file.close();
+  return buffer;
+}
+/**
+ * Validates no memory leakage for hipModuleLoad
+ */
+void testhipModuleLoadUnloadStress() {
+  for (int count = 0; count < TEST_ITERATIONS; count++) {
+    hipModule_t Module;
+    HIPCHECK(hipModuleLoad(&Module, CODEOBJ_FILE));
+    hipFunction_t Function;
+    HIPCHECK(hipModuleGetFunction(&Function, Module, "testWeightedCopy"));
+    HIPCHECK(hipModuleUnload(Module));
+  }
+}
+/**
+ * Validates no memory leakage for hipModuleLoadData
+ */
+void testhipModuleLoadDataUnloadStress() {
+  auto buffer = load_file();
+  for (int count = 0; count < TEST_ITERATIONS; count++) {
+    hipModule_t Module;
+    HIPCHECK(hipModuleLoadData(&Module, &buffer[0]));
+    hipFunction_t Function;
+    HIPCHECK(hipModuleGetFunction(&Function, Module, "testWeightedCopy"));
+    HIPCHECK(hipModuleUnload(Module));
+  }
+}
+/**
+ * Validates no memory leakage for hipModuleLoadDataEx
+ */
+void testhipModuleLoadDataExUnloadStress() {
+  auto buffer = load_file();
+  for (int count = 0; count < TEST_ITERATIONS; count++) {
+    hipModule_t Module;
+    HIPCHECK(hipModuleLoadDataEx(&Module, &buffer[0], 0,
+                                nullptr, nullptr));
+    hipFunction_t Function;
+    HIPCHECK(hipModuleGetFunction(&Function, Module, "testWeightedCopy"));
+    HIPCHECK(hipModuleUnload(Module));
+  }
+}
+
+int main(int argc, char* argv[]) {
+  HipTest::parseStandardArguments(argc, argv, true);
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipInit(0));
+  hipDevice_t device;
+  hipCtx_t context;
+  HIPCHECK(hipDeviceGet(&device, 0));
+  HIPCHECK(hipCtxCreate(&context, 0, device));
+#endif
+  if (p_tests == 0x1) {
+    testhipModuleLoadUnloadStress();
+  } else if (p_tests == 0x2) {
+    testhipModuleLoadDataUnloadStress();
+  } else if (p_tests == 0x3) {
+    testhipModuleLoadDataExUnloadStress();
+  }
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  passed();
+}
diff --git a/projects/clr/hipamd/tests/src/runtimeApi/module/hipModuleNegative.cpp b/projects/clr/hipamd/tests/src/runtimeApi/module/hipModuleNegative.cpp
new file mode 100644
index 0000000000..ede25307bc
--- /dev/null
+++ b/projects/clr/hipamd/tests/src/runtimeApi/module/hipModuleNegative.cpp
@@ -0,0 +1,866 @@
+/*
+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 WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS 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 IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
+ * TEST: %t --tests 0x10 EXCLUDE_HIP_PLATFORM hcc rocclr
+ * TEST: %t --tests 0x11 EXCLUDE_HIP_PLATFORM hcc rocclr
+ * TEST: %t --tests 0x12 EXCLUDE_HIP_PLATFORM hcc rocclr
+ * TEST: %t --tests 0x13
+ * TEST: %t --tests 0x14
+ * TEST: %t --tests 0x15
+ * TEST: %t --tests 0x20 EXCLUDE_HIP_PLATFORM hcc rocclr
+ * TEST: %t --tests 0x21 EXCLUDE_HIP_PLATFORM hcc rocclr
+ * TEST: %t --tests 0x22
+ * TEST: %t --tests 0x30 EXCLUDE_HIP_PLATFORM hcc rocclr
+ * TEST: %t --tests 0x31 EXCLUDE_HIP_PLATFORM hcc rocclr
+ * TEST: %t --tests 0x32
+ * TEST: %t --tests 0x40 EXCLUDE_HIP_PLATFORM hcc rocclr
+ * TEST: %t --tests 0x41
+ * TEST: %t --tests 0x42
+ * TEST: %t --tests 0x43 EXCLUDE_HIP_PLATFORM hcc rocclr
+ * TEST: %t --tests 0x44 EXCLUDE_HIP_PLATFORM nvcc
+ * TEST: %t --tests 0x45
+ * TEST: %t --tests 0x50 EXCLUDE_HIP_PLATFORM hcc rocclr nvcc
+ * TEST: %t --tests 0x51 EXCLUDE_HIP_PLATFORM hcc rocclr nvcc
+ * TEST: %t --tests 0x52 EXCLUDE_HIP_PLATFORM hcc rocclr
+ * TEST: %t --tests 0x53
+ * TEST: %t --tests 0x54
+ * TEST: %t --tests 0x55 EXCLUDE_HIP_PLATFORM nvcc
+ * TEST: %t --tests 0x56
+ * TEST: %t --tests 0x60 EXCLUDE_HIP_PLATFORM nvcc
+ * HIT_END
+ */
+#include <stdio.h>
+#include <stdlib.h>
+#include <ctime>
+#include <iostream>
+#include <fstream>
+#include <cstddef>
+#include <vector>
+#include "test_common.h"
+
+#define FILENAME_NONEXST "sample_nonexst.code"
+#define FILENAME_EMPTY "emptyfile.code"
+#define FILENAME_RAND "rand_file.code"
+#define RANDOMFILE_LEN 2048
+#define CODEOBJ_FILE "vcpy_kernel.code"
+#define KERNEL_NAME "hello_world"
+#define KERNEL_NAME_NONEXST "xyz"
+#define CODEOBJ_GLOBAL "global_kernel.code"
+#define DEVGLOB_VAR_NONEXIST "xyz"
+#define DEVGLOB_VAR "myDeviceGlobal"
+/**
+ * Internal Function
+ */
+std::vector<char> load_file(const char* filename) {
+  std::ifstream file(filename, std::ios::binary | std::ios::ate);
+  std::streamsize fsize = file.tellg();
+  file.seekg(0, std::ios::beg);
+  std::vector<char> buffer(fsize);
+  if (!file.read(buffer.data(), fsize)) {
+    failed("could not open code object '%s'\n", filename);
+  }
+  file.close();
+  return buffer;
+}
+
+/**
+ * Internal Function
+ */
+void createRandomFile(const char* filename) {
+  std::ofstream outfile(filename, std::ios::binary);
+  char buf[RANDOMFILE_LEN];
+  unsigned int seed = 1;
+  for (int i = 0; i < RANDOMFILE_LEN; i++) {
+    buf[i] = rand_r(&seed) % 256;
+  }
+  outfile.write(buf, RANDOMFILE_LEN);
+  outfile.close();
+}
+
+/**
+ * Internal Function
+ */
+#ifdef __HIP_PLATFORM_NVCC__
+void initHipCtx(hipCtx_t *pcontext) {
+  HIPCHECK(hipInit(0));
+  hipDevice_t device;
+  HIPCHECK(hipDeviceGet(&device, 0));
+  HIPCHECK(hipCtxCreate(pcontext, 0, device));
+}
+#endif
+
+/**
+ * Validates negative scenarios for hipModuleLoad
+ * module = nullptr
+ */
+bool testhipModuleLoadNeg10() {
+  bool TestPassed = false;
+  hipError_t ret;
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  if ((ret = hipModuleLoad(nullptr, CODEOBJ_FILE))
+       != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  return TestPassed;
+}
+
+/**
+ * Validates negative scenarios for hipModuleLoad
+ * fname = nullptr
+ */
+bool testhipModuleLoadNeg11() {
+  bool TestPassed = false;
+  hipError_t ret;
+  hipModule_t Module;
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  if ((ret = hipModuleLoad(&Module, nullptr))
+       != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  return TestPassed;
+}
+/**
+ * Validates negative scenarios for hipModuleLoad
+ * fname = empty file
+ */
+bool testhipModuleLoadNeg12() {
+  bool TestPassed = false;
+  hipError_t ret;
+  hipModule_t Module;
+  // Create an empty
+  std::fstream fs;
+  fs.open(FILENAME_EMPTY, std::ios::out);
+  fs.close();
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  if ((ret = hipModuleLoad(&Module, FILENAME_EMPTY))
+       != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  remove(FILENAME_EMPTY);
+  return TestPassed;
+}
+
+/**
+ * Validates negative scenarios for hipModuleLoad
+ * fname = ramdom file
+ */
+bool testhipModuleLoadNeg13() {
+  bool TestPassed = false;
+  hipError_t ret;
+  hipModule_t Module;
+  // Create a binary file with random numbers
+  createRandomFile(FILENAME_RAND);
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  if ((ret = hipModuleLoad(&Module, FILENAME_RAND))
+       != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  remove(FILENAME_RAND);
+  return TestPassed;
+}
+
+/**
+ * Validates negative scenarios for hipModuleLoad
+ * fname = non existent file
+ */
+bool testhipModuleLoadNeg14() {
+  bool TestPassed = false;
+  hipError_t ret;
+  hipModule_t Module;
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  if ((ret = hipModuleLoad(&Module, FILENAME_NONEXST))
+       != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  return TestPassed;
+}
+
+/**
+ * Validates negative scenarios for hipModuleLoad
+ * fname = empty string ""
+ */
+bool testhipModuleLoadNeg15() {
+  bool TestPassed = false;
+  hipError_t ret;
+  hipModule_t Module;
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  if ((ret = hipModuleLoad(&Module, "")) != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  return TestPassed;
+}
+
+/**
+ * Validates negative scenarios for hipModuleLoadData
+ * module = nullptr
+ */
+bool testhipModuleLoadDataNeg20() {
+  bool TestPassed = false;
+  hipError_t ret;
+  auto buffer = load_file(CODEOBJ_FILE);
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  if ((ret = hipModuleLoadData(nullptr, &buffer[0]))
+       != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  return TestPassed;
+}
+
+/**
+ * Validates negative scenarios for hipModuleLoadData
+ * image = nullptr
+ */
+bool testhipModuleLoadDataNeg21() {
+  bool TestPassed = false;
+  hipError_t ret;
+  hipModule_t Module;
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  if ((ret = hipModuleLoadData(&Module, nullptr))
+      != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  return TestPassed;
+}
+
+/**
+ * Validates negative scenarios for hipModuleLoadData
+ * image = ramdom file
+ */
+bool testhipModuleLoadDataNeg22() {
+  bool TestPassed = false;
+  hipError_t ret;
+  hipModule_t Module;
+  // Create a binary file with random numbers
+  createRandomFile(FILENAME_RAND);
+  // Open the code object file and copy it in a buffer
+  auto buffer = load_file(FILENAME_RAND);
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  if ((ret = hipModuleLoadData(&Module, &buffer[0]))
+       != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  remove(FILENAME_RAND);
+  return TestPassed;
+}
+/**
+ * Validates negative scenarios for hipModuleLoadDataEx
+ * module = nullptr
+ */
+bool testhipModuleLoadDataExNeg30() {
+  bool TestPassed = false;
+  hipError_t ret;
+  // Open the code object file and copy it in a buffer
+  auto buffer = load_file(CODEOBJ_FILE);
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  if ((ret = hipModuleLoadDataEx(nullptr, &buffer[0], 0, nullptr, nullptr))
+       != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  return TestPassed;
+}
+
+/**
+ * Validates negative scenarios for hipModuleLoadDataEx
+ * image = nullptr
+ */
+bool testhipModuleLoadDataExNeg31() {
+  bool TestPassed = false;
+  hipError_t ret;
+  hipModule_t Module;
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  if ((ret = hipModuleLoadDataEx(&Module, nullptr, 0, nullptr, nullptr))
+      != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  return TestPassed;
+}
+
+/**
+ * Validates negative scenarios for hipModuleLoadDataEx
+ * image = ramdom file
+ */
+bool testhipModuleLoadDataExNeg32() {
+  bool TestPassed = false;
+  hipError_t ret;
+  hipModule_t Module;
+  // Create a binary file with random numbers
+  createRandomFile(FILENAME_RAND);
+  // Open the code object file and copy it in a buffer
+  auto buffer = load_file(FILENAME_RAND);
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  if ((ret = hipModuleLoadDataEx(&Module, &buffer[0], 0, nullptr, nullptr))
+      != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  remove(FILENAME_RAND);
+  return TestPassed;
+}
+
+/**
+ * Validates negative scenarios for hipModuleGetFunction
+ * Function = nullptr
+ */
+bool testhipModuleGetFunctionNeg40() {
+  bool TestPassed = false;
+  hipError_t ret;
+  hipModule_t Module;
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  HIPCHECK(hipModuleLoad(&Module, CODEOBJ_FILE));
+  if ((ret = hipModuleGetFunction(nullptr, Module, KERNEL_NAME))
+      != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+  HIPCHECK(hipModuleUnload(Module));
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  return TestPassed;
+}
+
+/**
+ * Validates negative scenarios for hipModuleGetFunction
+ * Module is uninitialized
+ */
+bool testhipModuleGetFunctionNeg41() {
+  bool TestPassed = false;
+  hipError_t ret;
+  hipModule_t Module;
+  hipFunction_t Function;
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  if ((ret = hipModuleGetFunction(&Function, Module, KERNEL_NAME))
+      != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  return TestPassed;
+}
+
+/**
+ * Validates negative scenarios for hipModuleGetFunction
+ * kname = non existing function
+ */
+bool testhipModuleGetFunctionNeg42() {
+  bool TestPassed = false;
+  hipError_t ret;
+  hipModule_t Module;
+  hipFunction_t Function;
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  HIPCHECK(hipModuleLoad(&Module, CODEOBJ_FILE));
+  if ((ret = hipModuleGetFunction(&Function, Module, KERNEL_NAME_NONEXST))
+      != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+  HIPCHECK(hipModuleUnload(Module));
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  return TestPassed;
+}
+
+/**
+ * Validates negative scenarios for hipModuleGetFunction
+ * kname = nullptr
+ */
+bool testhipModuleGetFunctionNeg43() {
+  bool TestPassed = false;
+  hipError_t ret;
+  hipModule_t Module;
+  hipFunction_t Function;
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  HIPCHECK(hipModuleLoad(&Module, CODEOBJ_FILE));
+  if ((ret = hipModuleGetFunction(&Function, Module, nullptr))
+      != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+  HIPCHECK(hipModuleUnload(Module));
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  return TestPassed;
+}
+
+/**
+ * Validates negative scenarios for hipModuleGetFunction
+ * Module = Unloaded Module
+ */
+bool testhipModuleGetFunctionNeg44() {
+  bool TestPassed = false;
+  hipError_t ret;
+  hipModule_t Module;
+  hipFunction_t Function;
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  HIPCHECK(hipModuleLoad(&Module, CODEOBJ_FILE));
+  HIPCHECK(hipModuleUnload(Module));
+  if ((ret = hipModuleGetFunction(&Function, Module, KERNEL_NAME))
+      != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  return TestPassed;
+}
+
+/**
+ * Validates negative scenarios for hipModuleGetFunction
+ * kname = Empty String ""
+ */
+bool testhipModuleGetFunctionNeg45() {
+  bool TestPassed = false;
+  hipError_t ret;
+  hipModule_t Module;
+  hipFunction_t Function;
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  HIPCHECK(hipModuleLoad(&Module, CODEOBJ_FILE));
+  if ((ret = hipModuleGetFunction(&Function,
+                       Module, "")) != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+  HIPCHECK(hipModuleUnload(Module));
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  return TestPassed;
+}
+
+/**
+ * Validates negative scenarios for hipModuleGetGlobal
+ * dptr = nullptr
+ */
+bool testhipModuleGetGlobalNeg50() {
+  bool TestPassed = false;
+  hipError_t ret;
+  hipModule_t Module;
+  size_t deviceGlobalSize;
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  HIPCHECK(hipModuleLoad(&Module, CODEOBJ_GLOBAL));
+  if ((ret = hipModuleGetGlobal(nullptr,
+    &deviceGlobalSize, Module, DEVGLOB_VAR)) != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+  HIPCHECK(hipModuleUnload(Module));
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  return TestPassed;
+}
+
+/**
+ * Validates negative scenarios for hipModuleGetGlobal
+ * bytes = nullptr
+ */
+bool testhipModuleGetGlobalNeg51() {
+  bool TestPassed = false;
+  hipError_t ret;
+  hipModule_t Module;
+  hipDeviceptr_t deviceGlobal;
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  HIPCHECK(hipModuleLoad(&Module, CODEOBJ_GLOBAL));
+  if ((ret = hipModuleGetGlobal(&deviceGlobal, nullptr,
+      Module, DEVGLOB_VAR)) != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+  HIPCHECK(hipModuleUnload(Module));
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  return TestPassed;
+}
+
+/**
+ * Validates negative scenarios for hipModuleGetGlobal
+ * name = nullptr
+ */
+bool testhipModuleGetGlobalNeg52() {
+  bool TestPassed = false;
+  hipError_t ret;
+  hipModule_t Module;
+  hipDeviceptr_t deviceGlobal;
+  size_t deviceGlobalSize;
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  HIPCHECK(hipModuleLoad(&Module, CODEOBJ_GLOBAL));
+  if ((ret = hipModuleGetGlobal(&deviceGlobal,
+    &deviceGlobalSize, Module, nullptr)) != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+  HIPCHECK(hipModuleUnload(Module));
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  return TestPassed;
+}
+
+/**
+ * Validates negative scenarios for hipModuleGetGlobal
+ * name = wrong name
+ */
+bool testhipModuleGetGlobalNeg53() {
+  bool TestPassed = false;
+  hipError_t ret;
+  hipModule_t Module;
+  hipDeviceptr_t deviceGlobal;
+  size_t deviceGlobalSize;
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  HIPCHECK(hipModuleLoad(&Module, CODEOBJ_GLOBAL));
+  if ((ret = hipModuleGetGlobal(&deviceGlobal, &deviceGlobalSize,
+    Module, DEVGLOB_VAR_NONEXIST)) != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+  HIPCHECK(hipModuleUnload(Module));
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  return TestPassed;
+}
+
+/**
+ * Validates negative scenarios for hipModuleGetGlobal
+ * name = Empty String ""
+ */
+bool testhipModuleGetGlobalNeg54() {
+  bool TestPassed = false;
+  hipError_t ret;
+  hipModule_t Module;
+  hipDeviceptr_t deviceGlobal;
+  size_t deviceGlobalSize;
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  HIPCHECK(hipModuleLoad(&Module, CODEOBJ_GLOBAL));
+  if ((ret = hipModuleGetGlobal(&deviceGlobal,
+    &deviceGlobalSize, Module, "")) != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+  HIPCHECK(hipModuleUnload(Module));
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  return TestPassed;
+}
+
+/**
+ * Validates negative scenarios for hipModuleGetGlobal
+ * Module = Unloaded Module
+ */
+bool testhipModuleGetGlobalNeg55() {
+  bool TestPassed = false;
+  hipError_t ret;
+  hipModule_t Module;
+  hipDeviceptr_t deviceGlobal;
+  size_t deviceGlobalSize;
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  HIPCHECK(hipModuleLoad(&Module, CODEOBJ_GLOBAL));
+  HIPCHECK(hipModuleUnload(Module));
+  if ((ret = hipModuleGetGlobal(&deviceGlobal,
+    &deviceGlobalSize, Module, DEVGLOB_VAR)) != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  return TestPassed;
+}
+
+/**
+ * Validates negative scenarios for hipModuleGetGlobal
+ * Module = Uninitialized Module
+ */
+bool testhipModuleGetGlobalNeg56() {
+  bool TestPassed = false;
+  hipError_t ret;
+  hipModule_t Module;
+  hipDeviceptr_t deviceGlobal;
+  size_t deviceGlobalSize;
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  if ((ret = hipModuleGetGlobal(&deviceGlobal,
+    &deviceGlobalSize, Module, DEVGLOB_VAR)) != hipSuccess) {
+    TestPassed = true;
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  }
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  return TestPassed;
+}
+
+/**
+ * Validates negative scenarios for hipModuleUnload
+ * 1. Unload an uninitialized module
+ * 2. Unload an unloaded module
+ */
+bool testhipModuleLoadNeg60() {
+  bool TestPassed = true;
+  hipError_t ret;
+  hipModule_t Module;
+#ifdef __HIP_PLATFORM_NVCC__
+  hipCtx_t context;
+  initHipCtx(&context);
+#endif
+  // test case 1
+  if ((ret = hipModuleUnload(Module)) != hipSuccess) {
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  } else {
+    TestPassed &= false;
+  }
+  // test case 2
+  HIPCHECK(hipModuleLoad(&Module, CODEOBJ_FILE));
+  HIPCHECK(hipModuleUnload(Module));
+  if ((ret = hipModuleUnload(Module)) != hipSuccess) {
+    printf("Test Passed: Error Code Returned: '%s'(%d)\n",
+           hipGetErrorString(ret), ret);
+  } else {
+    TestPassed &= false;
+  }
+#ifdef __HIP_PLATFORM_NVCC__
+  HIPCHECK(hipCtxDestroy(context));
+#endif
+  return TestPassed;
+}
+
+int main(int argc, char* argv[]) {
+  HipTest::parseStandardArguments(argc, argv, true);
+  bool TestPassed = true;
+  if (p_tests == 0x10) {
+    TestPassed = testhipModuleLoadNeg10();
+  } else if (p_tests == 0x11) {
+    TestPassed = testhipModuleLoadNeg11();
+  } else if (p_tests == 0x12) {
+    TestPassed = testhipModuleLoadNeg12();
+  } else if (p_tests == 0x13) {
+    TestPassed = testhipModuleLoadNeg13();
+  } else if (p_tests == 0x14) {
+    TestPassed = testhipModuleLoadNeg14();
+  } else if (p_tests == 0x15) {
+    TestPassed = testhipModuleLoadNeg15();
+  } else if (p_tests == 0x20) {
+    TestPassed = testhipModuleLoadDataNeg20();
+  } else if (p_tests == 0x21) {
+    TestPassed = testhipModuleLoadDataNeg21();
+  } else if (p_tests == 0x22) {
+    TestPassed = testhipModuleLoadDataNeg22();
+  } else if (p_tests == 0x30) {
+    TestPassed = testhipModuleLoadDataExNeg30();
+  } else if (p_tests == 0x31) {
+    TestPassed = testhipModuleLoadDataExNeg31();
+  } else if (p_tests == 0x32) {
+    TestPassed = testhipModuleLoadDataExNeg32();
+  } else if (p_tests == 0x40) {
+    TestPassed = testhipModuleGetFunctionNeg40();
+  } else if (p_tests == 0x41) {
+    TestPassed = testhipModuleGetFunctionNeg41();
+  } else if (p_tests == 0x42) {
+    TestPassed = testhipModuleGetFunctionNeg42();
+  } else if (p_tests == 0x43) {
+    TestPassed = testhipModuleGetFunctionNeg43();
+  } else if (p_tests == 0x44) {
+    TestPassed = testhipModuleGetFunctionNeg44();
+  } else if (p_tests == 0x45) {
+    TestPassed = testhipModuleGetFunctionNeg45();
+  } else if (p_tests == 0x50) {
+    TestPassed = testhipModuleGetGlobalNeg50();
+  } else if (p_tests == 0x51) {
+    TestPassed = testhipModuleGetGlobalNeg51();
+  } else if (p_tests == 0x52) {
+    TestPassed = testhipModuleGetGlobalNeg52();
+  } else if (p_tests == 0x53) {
+    TestPassed = testhipModuleGetGlobalNeg53();
+  } else if (p_tests == 0x54) {
+    TestPassed = testhipModuleGetGlobalNeg54();
+  } else if (p_tests == 0x55) {
+    TestPassed = testhipModuleGetGlobalNeg55();
+  } else if (p_tests == 0x56) {
+    TestPassed = testhipModuleGetGlobalNeg56();
+  } else if (p_tests == 0x60) {
+    TestPassed = testhipModuleLoadNeg60();
+  } else {
+    printf("Invalid Test Case \n");
+    exit(1);
+  }
+  if (TestPassed) {
+    passed();
+  } else {
+    failed("Test Case %x Failed!", p_tests);
+  }
+}
diff --git a/projects/clr/hipamd/tests/src/runtimeApi/module/kernel_composite_test.cpp b/projects/clr/hipamd/tests/src/runtimeApi/module/kernel_composite_test.cpp
new file mode 100644
index 0000000000..e6cdf35564
--- /dev/null
+++ b/projects/clr/hipamd/tests/src/runtimeApi/module/kernel_composite_test.cpp
@@ -0,0 +1,41 @@
+/*
+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 WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS 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 IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "hip/hip_runtime.h"
+#define GLOBAL_BUF_SIZE 2048
+
+__device__ float deviceGlobalFloat;
+__device__ int   deviceGlobalInt1;
+__device__ int   deviceGlobalInt2;
+__device__ short deviceGlobalShort;
+__device__ char  deviceGlobalChar;
+
+__device__ int getSquareOfGlobalFloat() {
+  return static_cast<int>(deviceGlobalFloat*deviceGlobalFloat);
+}
+
+extern "C" __global__ void testWeightedCopy(int* a, int* b) {
+  int tx = hipThreadIdx_x;
+  b[tx] = deviceGlobalInt1*a[tx] + deviceGlobalInt2 +
+  static_cast<int>(deviceGlobalShort) + static_cast<int>(deviceGlobalChar)
+  + getSquareOfGlobalFloat();
+}