diff --git a/samples/2_Cookbook/11_texture_driver/Makefile b/samples/2_Cookbook/11_texture_driver/Makefile
new file mode 100644
index 0000000000..b68c5c31c7
--- /dev/null
+++ b/samples/2_Cookbook/11_texture_driver/Makefile
@@ -0,0 +1,17 @@
+HIP_PATH?= $(wildcard /opt/rocm/hip)
+ifeq (,$(HIP_PATH))
+	HIP_PATH=../../..
+endif
+HIPCC=$(HIP_PATH)/bin/hipcc
+HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --compiler)
+
+all: tex2dKernel.code texture2dDrv.out
+
+texture2dDrv.out: texture2dDrv.cpp
+	$(HIPCC) $(HIPCC_FLAGS) $< -o $@
+	
+tex2dKernel.code: tex2dKernel.cpp
+	$(HIPCC)  --genco  $(GENCO_FLAGS) $^ -o $@
+
+clean:
+	rm -f *.code *.out
diff --git a/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp b/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp
new file mode 100644
index 0000000000..17ed911808
--- /dev/null
+++ b/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp
@@ -0,0 +1,33 @@
+/*
+Copyright (c) 2015 - 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"
+extern texture<float, 2, hipReadModeElementType> tex;
+
+__global__ void tex2dKernel(hipLaunchParm lp, float* outputData,
+                             int width,
+                             int height)
+{
+    int x = hipBlockIdx_x*hipBlockDim_x + hipThreadIdx_x;
+    int y = hipBlockIdx_y*hipBlockDim_y + hipThreadIdx_y;
+    outputData[y*width + x] = tex2D(tex, x, y);
+}
diff --git a/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp b/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp
new file mode 100644
index 0000000000..87e23d6e94
--- /dev/null
+++ b/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp
@@ -0,0 +1,153 @@
+/*
+Copyright (c) 2015 - 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"
+#include "hip/hip_runtime_api.h"
+#include <iostream>
+#include <fstream>
+#include <vector>
+#include <hip/hip_hcc.h>
+
+#define fileName "tex2dKernel.code"
+
+texture<float, 2, hipReadModeElementType> tex;
+bool testResult = true;
+
+#define HIP_CHECK(cmd) \
+{\
+    hipError_t status = cmd;\
+    if(status != hipSuccess) {std::cout<<"error: #"<<status<<" ("<< hipGetErrorString(status) << ") at line:"<<__LINE__<<":  "<<#cmd<<std::endl;abort();}\
+}
+
+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;
+        }
+    }
+    hipModule_t Module;
+    HIP_CHECK(hipModuleLoad(&Module, fileName));
+
+    hipArray* array;
+	HIP_ARRAY_DESCRIPTOR desc;
+	desc.format = HIP_AD_FORMAT_FLOAT;
+	desc.numChannels = 1;
+	desc.width = width;
+	desc.height = height;
+    hipArrayCreate(&array, &desc);
+
+    hip_Memcpy2D copyParam;
+	memset(&copyParam, 0, sizeof(copyParam));
+	copyParam.dstMemoryType = hipMemoryTypeArray;
+	copyParam.dstArray = array;
+	copyParam.srcMemoryType = hipMemoryTypeHost;
+	copyParam.srcHost = hData;
+	copyParam.srcPitch = width * sizeof(float);
+	copyParam.widthInBytes = copyParam.srcPitch;
+	copyParam.height = height;
+    hipMemcpy_2D(&copyParam);
+
+    hipTexRefSetAddressMode(&tex, 0, hipAddressModeWrap);
+	hipTexRefSetAddressMode(&tex, 1, hipAddressModeWrap);
+	hipTexRefSetFilterMode(&tex, hipFilterModePoint);
+    hipTexRefSetFlags(&tex, 0);
+	hipTexRefSetFormat(&tex, HIP_AD_FORMAT_FLOAT, 1);
+    hipTexRefSetArray(&tex, array, HIP_TRSA_OVERRIDE_FORMAT);
+
+    float* dData = NULL;
+    hipMalloc((void **) &dData, size);
+
+#ifdef __HIP_PLATFORM_HCC__
+
+	struct {
+		uint32_t _hidden[6];  // genco path + wrapper-gen pass used hidden arguments.
+		void * _Ad;
+		unsigned int _Bd;
+		unsigned int _Cd;
+	} args;
+    args._Ad = dData;
+	args._Bd = width;
+	args._Cd = height;
+
+#endif
+
+#ifdef __HIP_PLATFORM_NVCC__
+	struct {
+		uint32_t _hidden[1];
+		void * _Ad;
+		unsigned int _Bd;
+		unsigned int _Cd;
+	} args;
+
+	args._hidden[0] = 0;
+	args._Ad = dData;
+    args._Bd = width;
+	args._Cd = height;
+#endif
+
+
+	size_t sizeTemp = sizeof(args);
+
+	void *config[] = {
+	  HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
+	  HIP_LAUNCH_PARAM_BUFFER_SIZE, &sizeTemp,
+	  HIP_LAUNCH_PARAM_END
+	};
+
+	hipFunction_t Function;
+	HIP_CHECK(hipModuleGetFunction(&Function, Module, "tex2dKernel"));
+
+    int temp1= width/16;
+    int temp2 = height/16;
+    HIP_CHECK(hipModuleLaunchKernel(Function, 16, 16, 1, temp1, temp2, 1, 0, 0, NULL, (void**)&config));
+    hipDeviceSynchronize();
+
+    float *hOutputData = (float *) malloc(size);
+    memset(hOutputData, 0,  size);
+    hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost);
+
+    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;
+            }
+        }
+    }
+    hipFree(dData);
+    hipFreeArray(array);
+}
+
+int main(int argc, char **argv){
+    hipInit(0);
+    runTest(argc, argv);
+    printf("%s ...\n", testResult ? "PASSED" : "FAILED");
+    exit(testResult ? EXIT_SUCCESS : EXIT_FAILURE);
+    return 0;
+}