Merge 'master' into 'amd-master'

Change-Id: If682eb69168da413136c4e5c110468ac658139b7


[ROCm/clr commit: c7366831e3]

projects/clr/hipamd/include/hip/hcc_detail/math_functions.h

@@ -58,6 +58,9 @@ __device__ float expf(float x);
 __device__ float expm1f(float x);
 __device__ int abs(int x);
 __device__ long long abs(long long x);
+__device__ double abs(double x);
+__device__ long int abs(long int x);
+__device__ long long int labs(long long int x);
 __device__ float fabsf(float x);
 __device__ float fdimf(float x, float y);
 __device__ float fdividef(float x, float y);

projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h

@@ -155,6 +155,8 @@ typedef struct cudaArray* hipArray_const_t;
 
 typedef cudaTextureObject_t hipTextureObject_t;
 typedef cudaSurfaceObject_t hipSurfaceObject_t;
+#define hipTextureType1D cudaTextureType1D
+#define hipTextureType1DLayered cudaTextureType1DLayered
 #define hipTextureType2D cudaTextureType2D
 #define hipTextureType3D cudaTextureType3D
 #define hipDeviceMapHost cudaDeviceMapHost
@@ -1161,8 +1163,8 @@ inline static hipError_t hipBindTexture(size_t* offset, const struct texture<T,
 }
 
 template <class T, int dim, enum cudaTextureReadMode readMode>
-inline static hipError_t hipBindTexture(size_t* offset, struct texture<T, dim, readMode>* tex,
-                                        const void* devPtr, const struct hipChannelFormatDesc* desc,
+inline static hipError_t hipBindTexture(size_t* offset, struct texture<T, dim, readMode>& tex,
+                                        const void* devPtr, const struct hipChannelFormatDesc& desc,
                                         size_t size = UINT_MAX) {
     return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, desc, size));
 }
@@ -1172,6 +1174,11 @@ inline static hipError_t hipUnbindTexture(struct texture<T, dim, readMode>* tex)
     return hipCUDAErrorTohipError(cudaUnbindTexture(tex));
 }
 
+inline static hipError_t hipBindTexture(size_t* offset, textureReference* tex, const void* devPtr,
+                                        const hipChannelFormatDesc* desc, size_t size = UINT_MAX){
+    return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, desc, size));
+}
+
 template <class T, int dim, enum hipTextureReadMode readMode>
 inline static hipError_t hipBindTextureToArray(struct texture<T, dim, readMode>& tex,
                                                hipArray_const_t array,

projects/clr/hipamd/src/hip_memory.cpp

@@ -1495,10 +1495,16 @@ __global__ void hip_copy2d_n(T* dst, const T* src, size_t width, size_t height,
 
     size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
     size_t idy = blockIdx.y * blockDim.y + threadIdx.y;
-    if((idx < width) && (idy < height)){
-        T *dstPtr = (T *)((uint8_t*) dst + idy * destPitch);
-        T *srcPtr = (T *)((uint8_t*) src + idy * srcPitch);
+    size_t floorWidth = (width/sizeof(T));
+    T *dstPtr = (T *)((uint8_t*) dst + idy * destPitch);
+    T *srcPtr = (T *)((uint8_t*) src + idy * srcPitch);
+    if((idx < floorWidth) && (idy < height)){
         dstPtr[idx] = srcPtr[idx];
+    } else if((idx < width) && (idy < height)){
+        size_t bytesToCopy = width - (floorWidth * sizeof(T));
+        dstPtr += floorWidth;
+        srcPtr += floorWidth;
+        __builtin_memcpy(reinterpret_cast<uint8_t*>(dstPtr), reinterpret_cast<const uint8_t*>(srcPtr),bytesToCopy);
     }
 }
 }  // namespace
@@ -1515,10 +1521,11 @@ void ihipMemsetKernel(hipStream_t stream, T* ptr, T val, size_t sizeBytes) {
 
 template <typename T>
 void ihipMemcpy2dKernel(hipStream_t stream, T* dst, const T* src, size_t width, size_t height, size_t destPitch, size_t srcPitch) {
-    size_t threadsPerBlock = 16;
-    uint32_t grid_dim_x = clamp_integer<size_t>( (width+(threadsPerBlock-1)) / threadsPerBlock, 1, UINT32_MAX);
-    uint32_t grid_dim_y = clamp_integer<size_t>( (height+(threadsPerBlock-1)) / threadsPerBlock, 1, UINT32_MAX);
-    hipLaunchKernelGGL(hip_copy2d_n, dim3(grid_dim_x,grid_dim_y), dim3(threadsPerBlock,threadsPerBlock), 0u, stream, dst, src,
+    size_t threadsPerBlock_x = 64;
+    size_t threadsPerBlock_y = 4;
+    uint32_t grid_dim_x = clamp_integer<size_t>( (width+(threadsPerBlock_x*sizeof(T)-1)) / (threadsPerBlock_x*sizeof(T)), 1, UINT32_MAX);
+    uint32_t grid_dim_y = clamp_integer<size_t>( (height+(threadsPerBlock_y-1)) / threadsPerBlock_y, 1, UINT32_MAX);
+    hipLaunchKernelGGL(hip_copy2d_n, dim3(grid_dim_x,grid_dim_y), dim3(threadsPerBlock_x,threadsPerBlock_y), 0u, stream, dst, src,
                        width, height, destPitch, srcPitch);
 }
 
@@ -1622,7 +1629,7 @@ hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch
                     stream->locked_copySync((unsigned char*)dst + i * dpitch,
                                     (unsigned char*)src + i * spitch, width, kind);
             } else {
-                ihipMemcpy2dKernel<uint8_t> (stream, static_cast<uint8_t*> (dst), static_cast<const uint8_t*> (src), width, height, dpitch, spitch);
+                ihipMemcpy2dKernel<uint32_t> (stream, static_cast<uint32_t*> (dst), static_cast<const uint32_t*> (src), width, height, dpitch, spitch);
                 stream->locked_wait();
             }
         } catch (ihipException& ex) {
@@ -1661,7 +1668,7 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp
                     e = hip_internal::memcpyAsync((unsigned char*)dst + i * dpitch,
                                           (unsigned char*)src + i * spitch, width, kind, stream);
             } else{
-                ihipMemcpy2dKernel<uint8_t> (stream, static_cast<uint8_t*> (dst), static_cast<const uint8_t*> (src), width, height, dpitch, spitch);
+                ihipMemcpy2dKernel<uint32_t> (stream, static_cast<uint32_t*> (dst), static_cast<const uint32_t*> (src), width, height, dpitch, spitch);
             }
         } catch (ihipException& ex) {
             e = ex._code;

projects/clr/hipamd/src/math_functions.cpp

@@ -59,6 +59,15 @@ __device__ int abs(int x) {
 __device__ long long abs(long long x) {
     return x >= 0 ? x : -x;  
 }
+__device__ double abs(double x) {
+    return x >= 0 ? x : -x;   
+}
+__device__ long int abs(long int x) {
+    return x >= 0 ? x : -x;   
+}
+__device__ long long int labs(long long int x) {
+    return x >= 0 ? x : -x;   
+}
 __device__ float fabsf(float x) { return hc::precise_math::fabsf(x); }
 __device__ float fdimf(float x, float y) { return hc::precise_math::fdimf(x, y); }
 __device__ float fdividef(float x, float y) { return x / y; }

projects/clr/hipamd/tests/src/texture/hipBindTexObj1D.cpp

@@ -0,0 +1,100 @@
+/*
+Copyright (c) 2015-2017 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 EXCLUDE_HIP_PLATFORM nvcc
+ * RUN: %t
+ * HIT_END
+ */
+
+#include "hip/hip_runtime.h"
+#include "hip/hip_runtime_api.h"
+#include "test_common.h"
+#include <iostream>
+#include <string.h>
+
+#define N 512
+using namespace std;
+
+texture<int, hipTextureType1D, hipReadModeElementType> tex;
+
+bool testResult = true;
+
+__global__ void kernel(int *out) {
+  int x = blockIdx.x * blockDim.x + threadIdx.x;
+  out[x] = tex1Dfetch(tex, x);
+}
+
+void runTest(void);
+
+int main(int argc, char **argv) {
+  runTest();
+
+  if (testResult) {
+    passed();
+  } else {
+    exit(EXIT_FAILURE);
+  }
+}
+
+void runTest() {
+  string out;
+  int *tex_buf;
+  int val[N], i, output[N];
+  size_t size = 0;
+
+  for (i = 0; i < N; i++) {
+    val[i] = i;
+    output[i] = 0;
+  }
+  hipChannelFormatDesc chan_desc =
+      hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindUnsigned);
+
+  hipMalloc(&tex_buf, N * sizeof(int));
+
+  hipMemcpy(tex_buf, val, N * sizeof(int), hipMemcpyHostToDevice);
+
+  tex.addressMode[0] = hipAddressModeWrap;
+  tex.filterMode = hipFilterModeLinear;
+  tex.normalized = true;
+
+  hipBindTexture(&size, &tex, (void *)tex_buf, &chan_desc, N * sizeof(int));
+
+  dim3 dimBlock(64, 1, 1);
+  dim3 dimGrid(N / dimBlock.x, 1, 1);
+
+  hipLaunchKernelGGL(kernel, dim3(dimGrid), dim3(dimBlock), 0, 0, output);
+
+  hipDeviceSynchronize();
+
+  hipMemcpy(output, tex_buf, N * sizeof(int), hipMemcpyDeviceToHost);
+
+  for (i = 0; i < N; i++) {
+    if (output[i] != val[i]) {
+      testResult = false;
+      return;
+    }
+  }
+  hipUnbindTexture(&tex);
+  hipFree(tex_buf);
+}

projects/clr/hipamd/tests/src/texture/hipTextureObj2D.cpp

@@ -1,108 +1,108 @@
-/* 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(float* outputData, hipTextureObject_t textureObject, int width,
-                            int height) {
-    int x = blockIdx.x * blockDim.x + threadIdx.x;
-    int y = blockIdx.y * blockDim.y + threadIdx.y;
-    outputData[y * width + x] = tex2D<float>(textureObject, x, y);
-}
-
-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;
-    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;
-
-    // Specify texture object parameters
-    struct hipTextureDesc texDesc;
-    memset(&texDesc, 0, sizeof(texDesc));
-    texDesc.addressMode[0] = hipAddressModeWrap;
-    texDesc.addressMode[1] = hipAddressModeWrap;
-    texDesc.filterMode = hipFilterModePoint;
-    texDesc.readMode = hipReadModeElementType;
-    texDesc.normalizedCoords = 0;
-
-    // Create texture object
-    hipTextureObject_t textureObject = 0;
-    hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL);
-
-    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, dData, textureObject,
-                       width, height);
-
-    hipDeviceSynchronize();
-
-    float* hOutputData = (float*)malloc(size);
-    memset(hOutputData, 0, size);
-    hipMemcpy(hOutputData, dData, 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;
-            }
-        }
-    }
-    hipDestroyTextureObject(textureObject);
-    hipFree(dData);
-    hipFreeArray(hipArray);
-}
+/* 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"
+
+__global__ void tex2DKernel(float* outputData, hipTextureObject_t textureObject, int width,
+                            int height) {
+    int x = blockIdx.x * blockDim.x + threadIdx.x;
+    int y = blockIdx.y * blockDim.y + threadIdx.y;
+    outputData[y * width + x] = tex2D<float>(textureObject, x, y);
+}
+
+int runTest(int argc, char** argv);
+
+int main(int argc, char** argv) {
+    int testResult = runTest(argc, argv);
+
+    if (testResult) {
+        passed();
+    } else {
+        exit(EXIT_FAILURE);
+    }
+}
+
+int runTest(int argc, char** argv) {
+    int testResult = 1;
+    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;
+    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;
+
+    // Specify texture object parameters
+    struct hipTextureDesc texDesc;
+    memset(&texDesc, 0, sizeof(texDesc));
+    texDesc.addressMode[0] = hipAddressModeWrap;
+    texDesc.addressMode[1] = hipAddressModeWrap;
+    texDesc.filterMode = hipFilterModePoint;
+    texDesc.readMode = hipReadModeElementType;
+    texDesc.normalizedCoords = 0;
+
+    // Create texture object
+    hipTextureObject_t textureObject = 0;
+    hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL);
+
+    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, dData, textureObject,
+                       width, height);
+
+    hipDeviceSynchronize();
+
+    float* hOutputData = (float*)malloc(size);
+    memset(hOutputData, 0, size);
+    hipMemcpy(hOutputData, dData, 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 = 0;
+                break;
+            }
+        }
+    }
+    hipDestroyTextureObject(textureObject);
+    hipFree(dData);
+    hipFreeArray(hipArray);
+    return testResult;
+}

projects/clr/hipamd/tests/src/texture/hipTextureRef2D.cpp

@@ -1,93 +1,93 @@
-/* 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"
-texture<float, 2, hipReadModeElementType> tex;
-
-bool testResult = true;
-
-__global__ void tex2DKernel(float* outputData,
-                            int width, int height) {
-    int x = blockIdx.x * blockDim.x + threadIdx.x;
-    int y = blockIdx.y * blockDim.y + threadIdx.y;
-    outputData[y * width + x] = tex2D(tex, x, y);
-}
-
-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;
-    hipMallocArray(&hipArray, &channelDesc, width, height);
-
-    hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice);
-
-    tex.addressMode[0] = hipAddressModeWrap;
-    tex.addressMode[1] = hipAddressModeWrap;
-    tex.filterMode = hipFilterModePoint;
-    tex.normalized = 0;
-
-    hipBindTextureToArray(tex, hipArray, channelDesc);
-
-    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, dData, width, height);
-    hipDeviceSynchronize();
-
-    float* hOutputData = (float*)malloc(size);
-    memset(hOutputData, 0, size);
-    hipMemcpy(hOutputData, dData, 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;
-            }
-        }
-    }
-    hipFree(dData);
-    hipFreeArray(hipArray);
-}
+/* 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"
+texture<float, 2, hipReadModeElementType> tex;
+
+__global__ void tex2DKernel(float* outputData,
+                            int width, int height) {
+    int x = blockIdx.x * blockDim.x + threadIdx.x;
+    int y = blockIdx.y * blockDim.y + threadIdx.y;
+    outputData[y * width + x] = tex2D(tex, x, y);
+}
+
+int runTest(int argc, char** argv);
+
+int main(int argc, char** argv) {
+    int testResult = runTest(argc, argv);
+    if (testResult) {
+        passed();
+    } else {
+        exit(EXIT_FAILURE);
+    }
+}
+
+int runTest(int argc, char** argv) {
+    int testResult = 1;
+    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;
+    hipMallocArray(&hipArray, &channelDesc, width, height);
+
+    hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice);
+
+    tex.addressMode[0] = hipAddressModeWrap;
+    tex.addressMode[1] = hipAddressModeWrap;
+    tex.filterMode = hipFilterModePoint;
+    tex.normalized = 0;
+
+    hipBindTextureToArray(tex, hipArray, channelDesc);
+
+    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, dData, width, height);
+    hipDeviceSynchronize();
+
+    float* hOutputData = (float*)malloc(size);
+    memset(hOutputData, 0, size);
+    hipMemcpy(hOutputData, dData, 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 = 0;
+                break;
+            }
+        }
+    }
+    hipFree(dData);
+    hipFreeArray(hipArray);
+    return testResult;
+}

projects/clr/hipamd/tests/src/texture/tex1Dfetch_linear.cpp

@@ -0,0 +1,112 @@
+/*
+Copyright (c) 2015-2017 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 EXCLUDE_HIP_PLATFORM nvcc
+ * RUN: %t
+ * HIT_END
+ */
+
+#include "hip/hip_runtime.h"
+#include "hip/hip_runtime_api.h"
+#include "test_common.h"
+#include <iostream>
+#include <string.h>
+
+#define N 512
+using namespace std;
+
+bool testResult = true;
+
+__global__ void tex1d_kernel(float *val, hipTextureObject_t obj) {
+  int k = blockIdx.x * blockDim.x + threadIdx.x;
+  val[k] = tex1Dfetch<float>(obj, k);
+}
+
+void runTest(void);
+
+int main(int argc, char **argv) {
+  runTest();
+
+  if (testResult) {
+    passed();
+  } else {
+    exit(EXIT_FAILURE);
+  }
+}
+
+void runTest() {
+
+  // Allocating the required buffer on gpu device
+  float *tex_buf, *tex_buf_check;
+  float val[N], output[N];
+  int i;
+  for (i = 0; i < N; i++)
+    val[i] = (i + 1) * (i + 1);
+  hipMalloc(&tex_buf, N * sizeof(float));
+
+  hipMalloc(&tex_buf_check, N * sizeof(float));
+
+  hipMemcpy(tex_buf, val, N * sizeof(float), hipMemcpyHostToDevice);
+
+  hipMemset(tex_buf_check, 0, N * sizeof(float));
+  hipResourceDesc res_lin;
+
+  memset(&res_lin, 0, sizeof(res_lin));
+
+  res_lin.resType = hipResourceTypeLinear;
+  res_lin.res.linear.devPtr = tex_buf;
+  res_lin.res.linear.desc.f = hipChannelFormatKindFloat;
+  res_lin.res.linear.desc.x = 32;
+  res_lin.res.linear.sizeInBytes = N * sizeof(float);
+
+  hipTextureDesc tex_desc;
+  memset(&tex_desc, 0, sizeof(tex_desc));
+  tex_desc.readMode = hipReadModeElementType;
+
+  // Creating texture object
+
+  hipTextureObject_t tex_obj = 0;
+
+  hipCreateTextureObject(&tex_obj, &res_lin, &tex_desc, NULL);
+
+  dim3 dimBlock(64, 1, 1);
+  dim3 dimGrid(N / dimBlock.x, 1, 1);
+
+  for (i = 0; i < N; i++)
+    output[i] = 0;
+
+  hipLaunchKernelGGL(tex1d_kernel, dim3(dimGrid), dim3(dimBlock), 0, 0,
+                     tex_buf_check, tex_obj);
+  hipDeviceSynchronize();
+
+  hipMemcpy(output, tex_buf_check, N * sizeof(float), hipMemcpyDeviceToHost);
+
+  for (i = 0; i < N; i++)
+    if (output[i] != val[i]) {
+      testResult = false;
+    }
+
+  hipDestroyTextureObject(tex_obj);
+  hipFree(tex_buf);
+  hipFree(tex_buf_check);
+}