Merge 'master' into 'amd-master'

Change-Id: Icb10bd2bba7544534e71e3c0e7e6f8ba46a7ee3c

hipamd/Jenkinsfile

@@ -167,7 +167,7 @@ def docker_build_inside_image( def build_image, String inside_args, String platf
     }
 
     // Cap the maximum amount of testing, in case of hangs
-    // Excluding hipMultiThreadDevice-pyramid test from automation; due to its flakiness which requires some investigation
+    // Excluding hipMultiThreadDevice-pyramid & hipMemoryAllocateCoherentDriver tests from automation; due to its flakiness which requires some investigation
     timeout(time: 1, unit: 'HOURS')
     {
       stage("${platform} unit testing")
@@ -177,7 +177,7 @@ def docker_build_inside_image( def build_image, String inside_args, String platf
             cd ${build_dir_rel}
             make install -j\$(nproc)
             make build_tests -i -j\$(nproc)
-            ctest -E pyramid
+            ctest -E "(hipMultiThreadDevice-pyramid|hipMemoryAllocateCoherentDriver)"
           """
         // If unit tests output a junit or xunit file in the future, jenkins can parse that file
         // to display test results on the dashboard

hipamd/bin/hipconfig

@@ -92,7 +92,7 @@ if (not defined $HIP_PLATFORM) {
 $HIP_PATH=$ENV{'HIP_PATH'} // Cwd::realpath (dirname (dirname $0)); # use parent directory of this tool
 
 if ($HIP_PLATFORM eq "hcc") {
-    $CPP_CONFIG= " -D__HIP_PLATFORM_HCC__=   -I$HIP_PATH/include -I$HCC_HOME/include";
+    $CPP_CONFIG= " -D__HIP_PLATFORM_HCC__=   -I$HIP_PATH/include -I$HCC_HOME/include -I$HSA_PATH/include";
 }
 if ($HIP_PLATFORM eq "nvcc") {
     $CPP_CONFIG = " -D__HIP_PLATFORM_NVCC__=  -I$HIP_PATH/include -I$CUDA_PATH/include";

hipamd/hip-config-clang.cmake.in

@@ -1,4 +1,16 @@
 @PACKAGE_INIT@
+if(NOT DEFINED HSA_PATH)
+    if(NOT DEFINED ENV{HSA_PATH})
+        set(HSA_PATH "/opt/rocm/hsa" CACHE PATH "Path to which HSA runtime has been installed")
+    else()
+        set(HSA_PATH $ENV{HSA_PATH} CACHE PATH "Path to which HSA runtime has been installed")
+    endif()
+endif()
+if(IS_ABSOLUTE ${HSA_PATH} AND EXISTS ${HSA_PATH} AND IS_DIRECTORY ${HSA_PATH})
+    message(STATUS "Looking for HSA runtime in: " ${HSA_PATH})
+else()
+    message(FATAL_ERROR "Don't know where to find HSA runtime. Please specify absolute path using -DHSA_PATH")
+endif()
 
 include(CMakeFindDependencyMacro OPTIONAL RESULT_VARIABLE _CMakeFindDependencyMacro_FOUND)
 if (NOT _CMakeFindDependencyMacro_FOUND)
@@ -43,7 +55,8 @@ set(HIP_COMPILER "@HIP_COMPILER@")
 set(HIP_RUNTIME "@HIP_RUNTIME@")
 
 set_and_check( hip_INCLUDE_DIR "@PACKAGE_INCLUDE_INSTALL_DIR@" )
-set_and_check( hip_INCLUDE_DIRS "${hip_INCLUDE_DIR}" )
+set_and_check( HSA_INCLUDE_PATH "${HSA_PATH}/include" )
+set( hip_INCLUDE_DIRS "${hip_INCLUDE_DIR}" "${HSA_INCLUDE_PATH}" )
 set_and_check( hip_LIB_INSTALL_DIR "@PACKAGE_LIB_INSTALL_DIR@" )
 set_and_check( hip_BIN_INSTALL_DIR "@PACKAGE_BIN_INSTALL_DIR@" )
 

hipamd/hip-config-hcc.cmake.in

@@ -1,4 +1,16 @@
 @PACKAGE_INIT@
+if(NOT DEFINED HSA_PATH)
+    if(NOT DEFINED ENV{HSA_PATH})
+        set(HSA_PATH "/opt/rocm/hsa" CACHE PATH "Path to which HSA runtime has been installed")
+    else()
+        set(HSA_PATH $ENV{HSA_PATH} CACHE PATH "Path to which HSA runtime has been installed")
+    endif()
+endif()
+if(IS_ABSOLUTE ${HSA_PATH} AND EXISTS ${HSA_PATH} AND IS_DIRECTORY ${HSA_PATH})
+    message(STATUS "Looking for HSA runtime in: " ${HSA_PATH})
+else()
+    message(FATAL_ERROR "Don't know where to find HSA runtime. Please specify absolute path using -DHSA_PATH")
+endif()
 
 include(CMakeFindDependencyMacro OPTIONAL RESULT_VARIABLE _CMakeFindDependencyMacro_FOUND)
 if (NOT _CMakeFindDependencyMacro_FOUND)
@@ -43,7 +55,8 @@ set(HIP_COMPILER "@HIP_COMPILER@")
 set(HIP_RUNTIME "@HIP_RUNTIME@")
 
 set_and_check( hip_INCLUDE_DIR "@PACKAGE_INCLUDE_INSTALL_DIR@" )
-set_and_check( hip_INCLUDE_DIRS "${hip_INCLUDE_DIR}" )
+set_and_check( HSA_INCLUDE_PATH "${HSA_PATH}/include" )
+set( hip_INCLUDE_DIRS "${hip_INCLUDE_DIR}" "${HSA_INCLUDE_PATH}" )
 set_and_check( hip_LIB_INSTALL_DIR "@PACKAGE_LIB_INSTALL_DIR@" )
 set_and_check( hip_BIN_INSTALL_DIR "@PACKAGE_BIN_INSTALL_DIR@" )
 

hipamd/tests/src/texture/hipNormalizedFloatValueTex.cpp

@@ -0,0 +1,122 @@
+/*
+Copyright (c) 2019-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 EXCLUDE_HIP_PLATFORM nvcc
+ * TEST: %t
+ * HIT_END
+ */
+
+#define CHECK(cmd) \
+{\
+    hipError_t error = cmd;\
+    if(error != hipSuccess) {\
+        fprintf(stderr, "error: '%s' (%d) at %s:%d\n", hipGetErrorString(error), error, __FILE__, __LINE__);\
+        exit(EXIT_FAILURE);\
+    }\
+}
+#include "hip/hip_runtime.h"
+#include "test_common.h"
+
+#define SIZE 10
+texture<float, hipTextureType1D, hipReadModeElementType> textureNormalizedVal_1D;
+
+__global__ void normalizedValTextureTest(unsigned int numElements, float* pDst)
+{
+    unsigned int elementID = hipThreadIdx_x;
+    if(elementID >= numElements)
+	    return;
+    float coord =(float) elementID/(numElements-1);
+    pDst[elementID] = tex1D(textureNormalizedVal_1D, coord);
+}
+
+template<typename T>
+bool textureTest(enum hipArray_Format texFormat)
+{
+    T hData[] = {65, 66, 67, 68, 69, 70, 71, 72,73,74};
+    T *dData = NULL;
+    CHECK(hipMalloc((void **) &dData, sizeof(T)*SIZE));
+    CHECK(hipMemcpyHtoD((hipDeviceptr_t)dData, hData, sizeof(T)*SIZE));
+    
+    textureReference* texRef = &textureNormalizedVal_1D;
+    CHECK(hipTexRefSetAddressMode(texRef, 0, hipAddressModeClamp));
+    CHECK(hipTexRefSetAddressMode(texRef, 1, hipAddressModeClamp));
+    CHECK(hipTexRefSetFilterMode(texRef, hipFilterModePoint));
+    CHECK(hipTexRefSetFlags(texRef, HIP_TRSF_NORMALIZED_COORDINATES)); 
+    CHECK(hipTexRefSetFormat(texRef, texFormat, 1));
+    
+    HIP_ARRAY_DESCRIPTOR desc;
+    desc.Width = SIZE;
+    desc.Height = 1;
+    desc.Format = texFormat;
+    desc.NumChannels = 1;
+    CHECK(hipTexRefSetAddress2D(texRef, &desc, (hipDeviceptr_t)dData, sizeof(T)*SIZE));
+    
+    bool testResult = true;
+    float *dOutputData = NULL;
+    CHECK(hipMalloc((void **) &dOutputData, sizeof(float)*SIZE));
+ 
+    hipLaunchKernelGGL(HIP_KERNEL_NAME(normalizedValTextureTest), dim3(1,1,1), dim3(SIZE,1,1), 0, 0, SIZE, dOutputData);
+
+    float *hOutputData = new float[SIZE];
+    CHECK(hipMemcpyDtoH(hOutputData, (hipDeviceptr_t)dOutputData, (sizeof(float)*SIZE)));
+    
+    for(int i = 0; i < SIZE; i++)
+    {
+    	if((float)hData[i]/texFormatToSize[texFormat] != hOutputData[i])
+        {
+	    printf("mismatch at index:%d for texType:%d output:%f\n",i,texFormat,hOutputData[i]);
+            testResult = false;
+	    break;
+        }
+    }
+    hipFree(dData);
+    hipFree(dOutputData);
+    delete [] hOutputData;
+    return testResult;
+}
+
+int main(int argc, char** argv)
+{
+    int device = 0;
+    bool status = true;
+    CHECK(hipSetDevice(device));
+    hipDeviceProp_t props;
+    CHECK(hipGetDeviceProperties(&props, device));
+    std::cout << "Device :: " << props.name << std::endl;
+    #ifdef __HIP_PLATFORM_HCC__
+    std::cout << "Arch - AMD GPU :: " << props.gcnArch << std::endl;
+    #endif
+    
+    status &= textureTest<char>          (HIP_AD_FORMAT_SIGNED_INT8);
+    status &= textureTest<unsigned char> (HIP_AD_FORMAT_UNSIGNED_INT8);
+    status &= textureTest<short>         (HIP_AD_FORMAT_SIGNED_INT16);
+    status &= textureTest<unsigned short>(HIP_AD_FORMAT_UNSIGNED_INT16);
+    status &= textureTest<float>         (HIP_AD_FORMAT_FLOAT);
+	
+    if(status){
+        passed();
+    }
+    else{
+        failed("checks failed!");
+    }
+}