diff --git a/projects/hip/tests/src/surface/hipSurfaceObj2D.cpp b/projects/hip/tests/src/surface/hipSurfaceObj2D.cpp
new file mode 100644
index 0000000000..cd87bc6701
--- /dev/null
+++ b/projects/hip/tests/src/surface/hipSurfaceObj2D.cpp
@@ -0,0 +1,115 @@
+/* HIT_START
+ * BUILD: %t %s ../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+#include <stdlib.h>
+#include <stdio.h>
+#include <string.h>
+
+#include <hip/hip_runtime.h>
+#include "test_common.h"
+
+bool testResult = true;
+
+__global__ void tex2DKernel(hipSurfaceObject_t surfaceObject,
+                            hipSurfaceObject_t outputSurfObj,
+                            int width,
+                            int height)
+{
+    int x = blockIdx.x*blockDim.x + threadIdx.x;
+    int y = blockIdx.y*blockDim.y + threadIdx.y;
+    float data;
+    surf2Dread(&data, surfaceObject, x*4, y, hipBoundaryModeZero);
+    surf2Dwrite(data, outputSurfObj, x*4, y, hipBoundaryModeZero);
+}
+
+void runTest(int argc, char **argv);
+
+int main(int argc, char **argv)
+{
+    runTest(argc, argv);
+
+    if(testResult) {
+        passed();
+    } else {
+        exit(EXIT_FAILURE);
+    }
+
+}
+
+void runTest(int argc, char **argv)
+{
+    unsigned int width = 256;
+    unsigned int height = 256;
+    unsigned int size = width * height * sizeof(float);
+    float* hData = (float*) malloc(size);
+    memset(hData, 0, size);
+    for (int i = 0; i < height; i++) {
+        for (int j = 0; j < width; j++) {
+            hData[i*width+j] = i*width+j;
+        }
+    }
+    printf("hData: ");
+    for (int i = 0; i < 64; i++) {
+        printf("%f  ", hData[i]);
+    }
+    printf("\n");
+
+    hipChannelFormatDesc channelDesc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat);
+    hipArray *hipArray, *hipOutArray;
+    hipMallocArray(&hipArray, &channelDesc, width, height);
+
+    hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice);
+
+    struct hipResourceDesc resDesc;
+    memset(&resDesc, 0, sizeof(resDesc));
+    resDesc.resType = hipResourceTypeArray;
+    resDesc.res.array.array = hipArray;
+    // Create surface object
+    hipSurfaceObject_t surfaceObject = 0;
+    hipCreateSurfaceObject(&surfaceObject, &resDesc);
+
+    hipMallocArray(&hipOutArray, &channelDesc, width, height);
+    struct hipResourceDesc resOutDesc;
+    memset(&resOutDesc, 0, sizeof(resOutDesc));
+    resOutDesc.resType = hipResourceTypeArray;
+    resOutDesc.res.array.array = hipOutArray;
+    hipSurfaceObject_t outSurfaceObject = 0;
+    hipCreateSurfaceObject(&outSurfaceObject, &resOutDesc);
+
+    float* dData = NULL;
+    hipMalloc((void **) &dData, size);
+
+    dim3 dimBlock(16, 16, 1);
+    dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1);
+
+    hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, surfaceObject,outSurfaceObject, width, height);
+
+    hipDeviceSynchronize();
+
+    float *hOutputData = (float *) malloc(size);
+    memset(hOutputData, 0,  size);
+    hipMemcpyFromArray(hOutputData, hipOutArray, 0, 0, size,hipMemcpyDeviceToHost);
+
+    printf("dData: ");
+    for (int i = 0; i < 64; i++) {
+        printf("%f  ", hOutputData[i]);
+    }
+    printf("\n");
+    for (int i = 0; i < height; i++) {
+        for (int j = 0; j < width; j++) {
+            if (hData[i*width+j] != hOutputData[i*width+j]) {
+                printf("Difference [ %d %d ]:%f ----%f\n",i, j, hData[i*width+j] , hOutputData[i*width+j]);
+                testResult = false;
+                break;
+            }
+        }
+    }
+    hipDestroySurfaceObject(surfaceObject);
+    hipDestroySurfaceObject(outSurfaceObject);
+    hipFree(dData);
+    hipFreeArray(hipArray);
+    hipFreeArray(hipOutArray);
+}
+