diff --git a/projects/hip/include/hip/hcc_detail/host_defines.h b/projects/hip/include/hip/hcc_detail/host_defines.h
index 689a4530da..11bd577f08 100644
--- a/projects/hip/include/hip/hcc_detail/host_defines.h
+++ b/projects/hip/include/hip/hcc_detail/host_defines.h
@@ -72,6 +72,7 @@ THE SOFTWARE.
 
 #define __noinline__ __attribute__((noinline))
 #define __forceinline__ inline __attribute__((always_inline))
+#define __hip_pinned_shadow__ __attribute__((hip_pinned_shadow))
 
 #else
 
diff --git a/projects/hip/samples/0_Intro/module_api_global/runKernel.cpp b/projects/hip/samples/0_Intro/module_api_global/runKernel.cpp
index 1e10dba5c6..7031b9f898 100644
--- a/projects/hip/samples/0_Intro/module_api_global/runKernel.cpp
+++ b/projects/hip/samples/0_Intro/module_api_global/runKernel.cpp
@@ -31,7 +31,6 @@ THE SOFTWARE.
 #define SIZE LEN * sizeof(float)
 
 #define fileName "vcpy_kernel.code"
-float myDeviceGlobalArray[16];
 #define HIP_CHECK(cmd)                                                                             \
     {                                                                                              \
         hipError_t status = cmd;                                                                   \
@@ -71,14 +70,17 @@ int main() {
     float* deviceGlobal;
     size_t deviceGlobalSize;
     HIP_CHECK(hipModuleGetGlobal((void**)&deviceGlobal, &deviceGlobalSize, Module, "myDeviceGlobal"));
-    *deviceGlobal = 42.0;
+    HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(deviceGlobal), &myDeviceGlobal_h, deviceGlobalSize));
 
 #define ARRAY_SIZE 16
 
     float myDeviceGlobalArray_h[ARRAY_SIZE];
+    float *myDeviceGlobalArray;
+    size_t myDeviceGlobalArraySize;
+    HIP_CHECK(hipModuleGetGlobal((void**)&myDeviceGlobalArray, &myDeviceGlobalArraySize, Module, "myDeviceGlobalArray"));
     for (int i = 0; i < ARRAY_SIZE; i++) {
         myDeviceGlobalArray_h[i] = i * 1000.0f;
-        myDeviceGlobalArray[i] = i * 1000.0f;
+        HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(myDeviceGlobalArray), &myDeviceGlobalArray_h, myDeviceGlobalArraySize));
     }
 
     struct {
diff --git a/projects/hip/samples/0_Intro/module_api_global/vcpy_kernel.cpp b/projects/hip/samples/0_Intro/module_api_global/vcpy_kernel.cpp
index 662dcd0cc3..9e8d7c9e29 100644
--- a/projects/hip/samples/0_Intro/module_api_global/vcpy_kernel.cpp
+++ b/projects/hip/samples/0_Intro/module_api_global/vcpy_kernel.cpp
@@ -25,8 +25,7 @@ THE SOFTWARE.
 #define ARRAY_SIZE (16)
 
 __device__ float myDeviceGlobal;
-extern float myDeviceGlobalArray[16];
-;
+__device__ float myDeviceGlobalArray[16];
 
 extern "C" __global__ void hello_world(const float* a, float* b) {
     int tx = hipThreadIdx_x;
diff --git a/projects/hip/samples/2_Cookbook/9_unroll/Readme.md b/projects/hip/samples/2_Cookbook/9_unroll/Readme.md
index 9b464f1472..0febeaa03c 100644
--- a/projects/hip/samples/2_Cookbook/9_unroll/Readme.md
+++ b/projects/hip/samples/2_Cookbook/9_unroll/Readme.md
@@ -16,17 +16,15 @@ Programmers familiar with CUDA, OpenCL will be able to quickly learn and start c
 
 ## Simple Matrix Transpose
 
-For this tutorial we will be using MatrixTranspose with shfl operation i.e., our 4_shfl tutorial since it is the only examples where we used loops inside the kernel.
+For this tutorial we will be using an example which sums up the row of a 2D matrix and writes it in a 1D array. 
 
-In this tutorial, we'll use `#pragma unroll`. In the same sourcecode, we used for MatrixTranspose. We'll add it just before the for loop as following:
+In this tutorial, we'll use `#pragma unroll`. In the same sourcecode, we used for gpuMatrixRowSum. We'll add it just before the for loop as following:
 
 ```
 #pragma unroll
-	for(int i=0;i<width;i++)
-	{
-		for(int j=0;j<width;j++)
-			out[i*width + j] = __shfl(val,j*width + i);
-	}
+for (int i = 0; i < width; i++) {
+    output[index] += input[index * width + i]
+}
 ```
 
 Specifying the optional parameter, #pragma unroll value, directs the unroller to unroll the loop value times. Be careful while using it.
diff --git a/projects/hip/samples/2_Cookbook/9_unroll/unroll.cpp b/projects/hip/samples/2_Cookbook/9_unroll/unroll.cpp
index 73b1c80f97..eeb688e296 100644
--- a/projects/hip/samples/2_Cookbook/9_unroll/unroll.cpp
+++ b/projects/hip/samples/2_Cookbook/9_unroll/unroll.cpp
@@ -25,100 +25,98 @@ THE SOFTWARE.
 // hip header file
 #include "hip/hip_runtime.h"
 
+#define LENGTH 4
 
-#define WIDTH 4
+#define SIZE (LENGTH * LENGTH)
 
-#define NUM (WIDTH * WIDTH)
+#define THREADS_PER_BLOCK 1
+#define BLOCKS_PER_GRID LENGTH
 
-#define THREADS_PER_BLOCK_X 4
-#define THREADS_PER_BLOCK_Y 4
-#define THREADS_PER_BLOCK_Z 1
-
-// Device (Kernel) function, it must be void
-__global__ void matrixTranspose(float* out, float* in, const int width) {
-    int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
-    float val = in[x];
-
-#pragma unroll
+// CPU function - basically scan each row and save the output in array
+void matrixRowSum(int* input, int* output, int width) {
     for (int i = 0; i < width; i++) {
-        for (int j = 0; j < width; j++) out[i * width + j] = __shfl(val, j * width + i);
-    }
-}
-
-// CPU implementation of matrix transpose
-void matrixTransposeCPUReference(float* output, float* input, const unsigned int width) {
-    for (unsigned int j = 0; j < width; j++) {
-        for (unsigned int i = 0; i < width; i++) {
-            output[i * width + j] = input[j * width + i];
+        for (int j = 0; j < width; j++) {
+            output[i] += input[i * width + j];
         }
     }
 }
 
-int main() {
-    float* Matrix;
-    float* TransposeMatrix;
-    float* cpuTransposeMatrix;
+// Device (kernel) function
+__global__ void gpuMatrixRowSum(int* input, int* output, int width) {
+    int index = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
+#pragma unroll
+    for (int i = 0; i < width; i++) {
+        output[index] += input[index * width + i];
+    }
+}
 
-    float* gpuMatrix;
-    float* gpuTransposeMatrix;
+int main() {
+    int* Matrix;
+    int* sumMatrix;
+    int* cpuSumMatrix;
+
+    int* gpuMatrix;
+    int* gpuSumMatrix;
 
     hipDeviceProp_t devProp;
     hipGetDeviceProperties(&devProp, 0);
 
     std::cout << "Device name " << devProp.name << std::endl;
 
-    int i;
-    int errors;
+    Matrix = (int*)malloc(sizeof(int) * SIZE);
+    sumMatrix = (int*)malloc(sizeof(int) * LENGTH);
+    cpuSumMatrix = (int*)malloc(sizeof(int) * LENGTH);
 
-    Matrix = (float*)malloc(NUM * sizeof(float));
-    TransposeMatrix = (float*)malloc(NUM * sizeof(float));
-    cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));
-
-    // initialize the input data
-    for (i = 0; i < NUM; i++) {
-        Matrix[i] = (float)i * 10.0f;
+    for (int i = 0; i < SIZE; i++) {
+        Matrix[i] = i * 2;
     }
 
-    // allocate the memory on the device side
-    hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
-    hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
+    for (int i = 0; i < LENGTH; i++) {
+        cpuSumMatrix[i] = 0;
+    }
 
-    // Memory transfer from host to device
-    hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice);
+    // Allocated Device Memory
+    hipMalloc((void**)&gpuMatrix, SIZE * sizeof(int));
+    hipMalloc((void**)&gpuSumMatrix, LENGTH * sizeof(int));
 
-    // Lauching kernel from host
-    hipLaunchKernelGGL(matrixTranspose, dim3(1), dim3(THREADS_PER_BLOCK_X * THREADS_PER_BLOCK_Y), 0, 0,
-                    gpuTransposeMatrix, gpuMatrix, WIDTH);
+    // Memory Copy to Device
+    hipMemcpy(gpuMatrix, Matrix, SIZE * sizeof(int), hipMemcpyHostToDevice);
+    hipMemcpy(gpuSumMatrix, cpuSumMatrix, LENGTH * sizeof(float), hipMemcpyHostToDevice);
 
-    // Memory transfer from device to host
-    hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost);
+    // Launch device kernels
+    hipLaunchKernelGGL(gpuMatrixRowSum, dim3(BLOCKS_PER_GRID), dim3(THREADS_PER_BLOCK), 0, 0,
+                       gpuMatrix, gpuSumMatrix, LENGTH);
+
+    // Memory copy back to device
+    hipMemcpy(sumMatrix, gpuSumMatrix, LENGTH * sizeof(int), hipMemcpyDeviceToHost);
+
+    // Cpu implementation
+    matrixRowSum(Matrix, cpuSumMatrix, LENGTH);
 
-    // CPU MatrixTranspose computation
-    matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH);
 
     // verify the results
-    errors = 0;
-    double eps = 1.0E-6;
-    for (i = 0; i < NUM; i++) {
-        if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps) {
-            printf("%d cpu: %f gpu  %f\n", i, cpuTransposeMatrix[i], TransposeMatrix[i]);
+    int errors = 0;
+    for (int i = 0; i < LENGTH; i++) {
+        if (sumMatrix[i] != cpuSumMatrix[i]) {
+            printf("%d - cpu: %d gpu: %d\n", i, sumMatrix[i], cpuSumMatrix[i]);
             errors++;
         }
     }
-    if (errors != 0) {
-        printf("FAILED: %d errors\n", errors);
+
+    if (errors == 0) {
+        printf("PASSED\n");
     } else {
-        printf("PASSED!\n");
+        printf("FAILED with %d errors\n", errors);
     }
 
-    // free the resources on device side
+    // GPU Free
     hipFree(gpuMatrix);
-    hipFree(gpuTransposeMatrix);
+    hipFree(gpuSumMatrix);
 
-    // free the resources on host side
+    // CPU Free
     free(Matrix);
-    free(TransposeMatrix);
-    free(cpuTransposeMatrix);
+    free(sumMatrix);
+    free(cpuSumMatrix);
 
     return errors;
 }
diff --git a/projects/hip/src/hip_event.cpp b/projects/hip/src/hip_event.cpp
index ea014ab292..7438c7d3d0 100644
--- a/projects/hip/src/hip_event.cpp
+++ b/projects/hip/src/hip_event.cpp
@@ -85,7 +85,7 @@ hipError_t ihipEventCreate(hipEvent_t* event, unsigned flags) {
         (flags & ~supportedFlags) ||             // can't set any unsupported flags.
         (flags & releaseFlags) == releaseFlags;  // can't set both release flags
 
-    if (!illegalFlags) {
+    if (event && !illegalFlags) {
         *event = new ihipEvent_t(flags);
     } else {
         e = hipErrorInvalidValue;
@@ -109,46 +109,51 @@ hipError_t hipEventCreate(hipEvent_t* event) {
 
 hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) {
     HIP_INIT_SPECIAL_API(hipEventRecord, TRACE_SYNC, event, stream);
+ 
+    hipError_t status;
+    if (event){
+        auto ecd = event->locked_copyCrit();
+        if( ecd._state != hipEventStatusUnitialized) {
+            stream = ihipSyncAndResolveStream(stream);
 
-    auto ecd = event->locked_copyCrit();
+            if (HIP_SYNC_NULL_STREAM && stream->isDefaultStream()) {
+                // TODO-HIP_SYNC_NULL_STREAM : can remove this code when HIP_SYNC_NULL_STREAM = 0
+                //
+                // If default stream , then wait on all queues.
+                ihipCtx_t* ctx = ihipGetTlsDefaultCtx();
+                ctx->locked_syncDefaultStream(true, true);
 
-    if (event && ecd._state != hipEventStatusUnitialized) {
-        stream = ihipSyncAndResolveStream(stream);
+                { 
+                    LockedAccessor_EventCrit_t eCrit(event->criticalData());
+                    eCrit->_eventData.marker(hc::completion_future());  // reset event
+                    eCrit->_eventData._stream = stream;
+                    eCrit->_eventData._timestamp = hc::get_system_ticks();
+                    eCrit->_eventData._state = hipEventStatusComplete;
+                }
+                status = hipSuccess;
+            } else {
+                // Record the event in the stream:
+                // Keep a copy outside the critical section so we lock stream first, then event - to
+                // avoid deadlock
+                hc::completion_future cf = stream->locked_recordEvent(event);
 
-        if (HIP_SYNC_NULL_STREAM && stream->isDefaultStream()) {
-            // TODO-HIP_SYNC_NULL_STREAM : can remove this code when HIP_SYNC_NULL_STREAM = 0
-            //
-            // If default stream , then wait on all queues.
-            ihipCtx_t* ctx = ihipGetTlsDefaultCtx();
-            ctx->locked_syncDefaultStream(true, true);
+                {
+                    LockedAccessor_EventCrit_t eCrit(event->criticalData());
+                    eCrit->_eventData.marker(cf);
+                    eCrit->_eventData._stream = stream;
+                    eCrit->_eventData._timestamp = 0;
+                    eCrit->_eventData._state = hipEventStatusRecording;
+                }
 
-            {
-                LockedAccessor_EventCrit_t eCrit(event->criticalData());
-                eCrit->_eventData.marker(hc::completion_future());  // reset event
-                eCrit->_eventData._stream = stream;
-                eCrit->_eventData._timestamp = hc::get_system_ticks();
-                eCrit->_eventData._state = hipEventStatusComplete;
+                status = hipSuccess;
             }
-            return ihipLogStatus(hipSuccess);
         } else {
-            // Record the event in the stream:
-            // Keep a copy outside the critical section so we lock stream first, then event - to
-            // avoid deadlock
-            hc::completion_future cf = stream->locked_recordEvent(event);
-
-            {
-                LockedAccessor_EventCrit_t eCrit(event->criticalData());
-                eCrit->_eventData.marker(cf);
-                eCrit->_eventData._stream = stream;
-                eCrit->_eventData._timestamp = 0;
-                eCrit->_eventData._state = hipEventStatusRecording;
-            }
-
-            return ihipLogStatus(hipSuccess);
+            status = hipErrorInvalidResourceHandle;
         }
     } else {
-        return ihipLogStatus(hipErrorInvalidResourceHandle);
+            status = hipErrorInvalidResourceHandle;
     }
+    return ihipLogStatus(status);
 }
 
 
@@ -202,11 +207,13 @@ hipError_t hipEventElapsedTime(float* ms, hipEvent_t start, hipEvent_t stop) {
 
     hipError_t status = hipSuccess;
 
-    *ms = 0.0f;
-
-    if ((start == nullptr) || (stop == nullptr)) {
+    if (ms == nullptr) {
+        status = hipErrorInvalidValue;
+    }
+    else if ((start == nullptr) || (stop == nullptr)) {
         status = hipErrorInvalidResourceHandle;
     } else {
+        *ms = 0.0f;
         auto startEcd = start->locked_copyCrit();
         auto stopEcd = stop->locked_copyCrit();
 
@@ -256,18 +263,26 @@ hipError_t hipEventElapsedTime(float* ms, hipEvent_t start, hipEvent_t stop) {
 
 hipError_t hipEventQuery(hipEvent_t event) {
     HIP_INIT_SPECIAL_API(hipEventQuery, TRACE_QUERY, event);
+ 
+    hipError_t status = hipSuccess;
 
-    if (!(event->_flags & hipEventReleaseToSystem)) {
-        tprintf(DB_WARN,
-                "hipEventQuery on event without system-scope fence ; consider creating with "
-                "hipEventReleaseToSystem\n");
-    }
-
-    auto ecd = event->locked_copyCrit();
-
-    if ((ecd._state == hipEventStatusRecording) && !ecd._stream->locked_eventIsReady(event)) {
-        return ihipLogStatus(hipErrorNotReady);
+    if ( NULL == event)
+    {
+        status = hipErrorInvalidResourceHandle;
     } else {
-        return ihipLogStatus(hipSuccess);
+        if (!(event->_flags & hipEventReleaseToSystem)) {
+            tprintf(DB_WARN,
+                    "hipEventQuery on event without system-scope fence ; consider creating with "
+                    "hipEventReleaseToSystem\n");
+        }
+
+        auto ecd = event->locked_copyCrit();
+
+        if ((ecd._state == hipEventStatusRecording) && !ecd._stream->locked_eventIsReady(event)) {
+            status = hipErrorNotReady;
+        } else {
+            status = hipSuccess;
+        }
     }
+    return ihipLogStatus(status);
 }
diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp
index 36edcdb338..b85ba61584 100644
--- a/projects/hip/src/hip_memory.cpp
+++ b/projects/hip/src/hip_memory.cpp
@@ -243,13 +243,12 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) {
 
     auto ctx = ihipGetTlsDefaultCtx();
     // return NULL pointer when malloc size is 0
-    if (sizeBytes == 0) {
+    if ( nullptr == ctx || nullptr == ptr)  {
+        hip_status = hipErrorInvalidValue;
+    }
+    else if (sizeBytes == 0) {
         *ptr = NULL;
         hip_status = hipSuccess;
-
-    } else if ((ctx == nullptr) || (ptr == nullptr)) {
-        hip_status = hipErrorInvalidValue;
-
     } else {
         auto device = ctx->getWriteableDevice();
         *ptr = hip_internal::allocAndSharePtr("device_mem", sizeBytes, ctx, false /*shareWithAll*/,
@@ -309,12 +308,12 @@ hipError_t ihipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
     }
 
     auto ctx = ihipGetTlsDefaultCtx();
-
-    if (sizeBytes == 0) {
+    if ((ctx == nullptr) || (ptr == nullptr)) {
+        hip_status = hipErrorInvalidValue;
+    }
+    else if (sizeBytes == 0) {
         hip_status = hipSuccess;
         // TODO - should size of 0 return err or be siliently ignored?
-    } else if ((ctx == nullptr) || (ptr == nullptr)) {
-        hip_status = hipErrorInvalidValue;
     } else {
         unsigned trueFlags = flags;
         if (flags == hipHostMallocDefault) {
@@ -400,7 +399,7 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) {
 // width in bytes
 hipError_t ihipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height, size_t depth) {
     hipError_t hip_status = hipSuccess;
-    if(ptr==NULL)
+    if(ptr==NULL || pitch == NULL)
      {
 	hip_status=hipErrorInvalidValue;
        	return hip_status;
@@ -916,11 +915,8 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) {
     am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, hostPtr);
     if (status == AM_SUCCESS) {
         *flagsPtr = amPointerInfo._appAllocationFlags;
-        if (*flagsPtr == 0) {
-            hip_status = hipErrorInvalidValue;
-        } else {
-            hip_status = hipSuccess;
-        }
+        //0 is valid flag hipHostMallocDefault, and during hipHostMalloc if unsupported flags are passed as parameter it throws error
+        hip_status = hipSuccess;
         tprintf(DB_MEM, " %s: host ptr=%p\n", __func__, hostPtr);
     } else {
         hip_status = hipErrorInvalidValue;
diff --git a/projects/hip/src/hip_peer.cpp b/projects/hip/src/hip_peer.cpp
index 32beda2caf..f2fca94404 100644
--- a/projects/hip/src/hip_peer.cpp
+++ b/projects/hip/src/hip_peer.cpp
@@ -41,8 +41,10 @@ THE SOFTWARE.
 hipError_t ihipDeviceCanAccessPeer(int* canAccessPeer, hipCtx_t thisCtx, hipCtx_t peerCtx) {
     hipError_t err = hipSuccess;
 
-
-    if ((thisCtx != NULL) && (peerCtx != NULL)) {
+    if(canAccessPeer == NULL) {
+       err = hipErrorInvalidValue;
+    }
+    else if ((thisCtx != NULL) && (peerCtx != NULL)) {
         if (thisCtx == peerCtx) {
             *canAccessPeer = 0;
             tprintf(DB_MEM, "Can't be peer to self. (this=%s, peer=%s)\n",
diff --git a/projects/hip/src/hip_stream.cpp b/projects/hip/src/hip_stream.cpp
index c328e34b79..b748749538 100644
--- a/projects/hip/src/hip_stream.cpp
+++ b/projects/hip/src/hip_stream.cpp
@@ -56,6 +56,8 @@ hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags, int priorit
     if (ctx) {
         if (HIP_FORCE_NULL_STREAM) {
             *stream = 0;
+        } else if( NULL == stream ){
+            e = hipErrorInvalidValue;
         } else {
             hc::accelerator acc = ctx->getWriteableDevice()->_acc;
 
@@ -65,7 +67,7 @@ hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags, int priorit
             // CUDA stream behavior is that all kernels submitted will automatically
             // wait for prev to complete, this behaviour will be mainatined by 
             // hipModuleLaunchKernel. execute_any_order will help 
-            // hipExtModuleLaunchKernel , which uses a special flag
+	    // hipExtModuleLaunchKernel , which uses a special flag
 
             {
                 // Obtain mutex access to the device critical data, release by destructor
@@ -80,9 +82,9 @@ hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags, int priorit
                 ctxCrit->addStream(istream);
                 *stream = istream;
             }
+            tprintf(DB_SYNC, "hipStreamCreate, %s\n", ToString(*stream).c_str());
         }
 
-        tprintf(DB_SYNC, "hipStreamCreate, %s\n", ToString(*stream).c_str());
     } else {
         e = hipErrorInvalidDevice;
     }
@@ -94,8 +96,10 @@ hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags, int priorit
 //---
 hipError_t hipStreamCreateWithFlags(hipStream_t* stream, unsigned int flags) {
     HIP_INIT_API(hipStreamCreateWithFlags, stream, flags);
-
-    return ihipLogStatus(ihipStreamCreate(stream, flags, priority_normal));
+    if(flags == hipStreamDefault || flags == hipStreamNonBlocking)
+        return ihipLogStatus(ihipStreamCreate(stream, flags, priority_normal));
+    else
+        return ihipLogStatus(hipErrorInvalidValue);
 }
 
 //---
@@ -128,25 +132,25 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
 
     hipError_t e = hipSuccess;
 
-    auto ecd = event->locked_copyCrit();
-
     if (event == nullptr) {
         e = hipErrorInvalidResourceHandle;
 
-    } else if ((ecd._state != hipEventStatusUnitialized) && (ecd._state != hipEventStatusCreated)) {
-        if (HIP_SYNC_STREAM_WAIT || (HIP_SYNC_NULL_STREAM && (stream == 0))) {
-            // conservative wait on host for the specified event to complete:
-            // return _stream->locked_eventWaitComplete(this, waitMode);
-            //
-            ecd._stream->locked_eventWaitComplete(
+    } else {
+        auto ecd = event->locked_copyCrit(); 
+        if ((ecd._state != hipEventStatusUnitialized) && (ecd._state != hipEventStatusCreated)) {
+            if (HIP_SYNC_STREAM_WAIT || (HIP_SYNC_NULL_STREAM && (stream == 0))) {
+                // conservative wait on host for the specified event to complete:
+                // return _stream->locked_eventWaitComplete(this, waitMode);
+                //
+                ecd._stream->locked_eventWaitComplete(
                 ecd.marker(), (event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked
                                                                      : hc::hcWaitModeActive);
-        } else {
-            stream = ihipSyncAndResolveStream(stream);
-            // This will use create_blocking_marker to wait on the specified queue.
-            stream->locked_streamWaitEvent(ecd);
+            } else {
+                stream = ihipSyncAndResolveStream(stream);
+                // This will use create_blocking_marker to wait on the specified queue.
+                stream->locked_streamWaitEvent(ecd);
+            }
         }
-
     }  // else event not recorded, return immediately and don't create marker.
 
     return ihipLogStatus(e);
diff --git a/projects/hip/tests/src/texture/hipBindTexRef1DFetch.cpp b/projects/hip/tests/src/texture/hipBindTexRef1DFetch.cpp
index 546e9775e8..52a0d99ac1 100644
--- a/projects/hip/tests/src/texture/hipBindTexRef1DFetch.cpp
+++ b/projects/hip/tests/src/texture/hipBindTexRef1DFetch.cpp
@@ -32,6 +32,9 @@ THE SOFTWARE.
 
 #define N 512
 
+#if __HIP__
+__hip_pinned_shadow__
+#endif
 texture<float, 1, hipReadModeElementType> tex;
 
 __global__ void kernel(float *out) {
diff --git a/projects/hip/tests/src/texture/hipTextureRef2D.cpp b/projects/hip/tests/src/texture/hipTextureRef2D.cpp
index 2b19d1b5b3..2fb4097b0a 100644
--- a/projects/hip/tests/src/texture/hipTextureRef2D.cpp
+++ b/projects/hip/tests/src/texture/hipTextureRef2D.cpp
@@ -9,6 +9,9 @@
 
 #include <hip/hip_runtime.h>
 #include "test_common.h"
+#if __HIP__
+__hip_pinned_shadow__
+#endif
 texture<float, 2, hipReadModeElementType> tex;
 
 __global__ void tex2DKernel(float* outputData,