Merge 'master' into 'amd-master'

Change-Id: I31b31fc6614c32b88df3ddb0796af42d19deb65f

include/hip/hcc_detail/device_functions.h

@@ -27,11 +27,12 @@ THE SOFTWARE.
 #include "math_fwd.h"
 
 #include <hip/hip_runtime_api.h>
+#include <stddef.h>
+
+
 #include <hip/hip_vector_types.h>
 #include <hip/hcc_detail/device_library_decls.h>
 #include <hip/hcc_detail/llvm_intrinsics.h>
-#include <stddef.h>
-
 /*
 Integer Intrinsics
 */

include/hip/hcc_detail/hip_runtime.h

@@ -131,7 +131,7 @@ extern int HIP_TRACE_API;
 
 
 // Feature tests:
-#if defined(__HCC_ACCELERATOR__) && (__HCC_ACCELERATOR__ != 0)
+#if (defined(__HCC_ACCELERATOR__) && (__HCC_ACCELERATOR__ != 0)) || __HIP_DEVICE_COMPILE__
 // Device compile and not host compile:
 
 // 32-bit Atomics:

include/hip/hcc_detail/math_functions.h

@@ -22,14 +22,16 @@ THE SOFTWARE.
 
 #pragma once
 
-#include "math_fwd.h"
-
-#include <hip/hcc_detail/host_defines.h>
-
 #include <assert.h>
 #include <limits.h>
 #include <limits>
 #include <stdint.h>
+#include <algorithm>
+
+#include <hip/hcc_detail/host_defines.h>
+
+#include "hip_fp16_math_fwd.h"
+#include "math_fwd.h"
 
 // HCC's own math functions should be included first, otherwise there will
 // be conflicts when hip/math_functions.h is included before hip/hip_runtime.h.

tests/src/texture/hipBindTexRef1DFetch.cpp

@@ -56,7 +56,7 @@ int runTest() {
     int testResult = 1;
     float *texBuf;
     float val[N], output[N];
-    size_t size = 0;
+    size_t offset = 0;
     float *devBuf;
     for (int i = 0; i < N; i++) {
         val[i] = (float)i;
@@ -74,7 +74,8 @@ int runTest() {
     tex.filterMode = hipFilterModePoint;
     tex.normalized = 0;
 
-    HIPCHECK(hipBindTexture(&size, tex, (void *)texBuf, chanDesc, N * sizeof(float)));
+    HIPCHECK(hipBindTexture(&offset, tex, (void *)texBuf, chanDesc, N * sizeof(float)));
+    HIPCHECK(hipGetTextureAlignmentOffset(&offset,&tex));
 
     dim3 dimBlock(64, 1, 1);
     dim3 dimGrid(N / dimBlock.x, 1, 1);

tests/src/texture/hipGetAlgntoffset2D.cpp

@@ -1,78 +0,0 @@
-/*
-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.
-*/
-
-/* HIT_START
- * BUILD: %t %s ../test_common.cpp
- * RUN: %t
- * HIT_END
- */
-#include <hip/hip_runtime.h>
-#include "test_common.h"
-
-using namespace std;
-#define R 8 //rows, height
-#define C 8 //columns, width
-
-texture<int, hipTextureType2D,hipReadModeElementType> tex;
-
-bool runTest(void);
-
-int main(int argc, char** argv) {
-    bool testResult=runTest();
-
-    if (testResult) {
-        passed();
-    } else {
-        exit(EXIT_FAILURE);
-    }
-}
-
-bool runTest()
-{
-int val[R][C],i,j;
-size_t offset;
-
-for(i=0;i<R;i++)
-    for(j=0;j<C;j++)
-	{
-	 val[i][j]=(i+1)*(j+1);
-	}
-hipChannelFormatDesc chan_desc=hipCreateChannelDesc(32,0,0,0,hipChannelFormatKindSigned);
-hipArray *hipArray;
-HIPCHECK(hipMallocArray(&hipArray, &chan_desc,C,R,0));
-
-HIPCHECK(hipMemcpyToArray(hipArray,0,0, val, R*C*sizeof(int), hipMemcpyHostToDevice));
-
-tex.addressMode[0]=hipAddressModeWrap;
-tex.addressMode[1]=hipAddressModeWrap;
-tex.filterMode=hipFilterModePoint;
-tex.normalized=0;
-
-HIPCHECK(hipBindTextureToArray(&tex, hipArray, &chan_desc));
-HIPCHECK(hipGetTextureAlignmentOffset(&offset,&tex));
-HIPCHECK(hipUnbindTexture(&tex));
-HIPCHECK(hipFreeArray(hipArray));
-if(offset != 0)
-    return false;
-  else
-    return true;
-}