Merge 'master' into 'amd-master'
Change-Id: Icb10bd2bba7544534e71e3c0e7e6f8ba46a7ee3c
[ROCm/hip commit: 73fe08d8f0]
This commit is contained in:
Jenkins
فروخته شده
+2
-2
@@ -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
|
||||
|
||||
@@ -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";
|
||||
|
||||
تفاوت فایلی نمایش داده نمی شود زیرا این فایل بسیار بزرگ است
Diff را بارگزاری کن
@@ -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@" )
|
||||
|
||||
|
||||
@@ -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@" )
|
||||
|
||||
|
||||
@@ -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!");
|
||||
}
|
||||
}
|
||||
مرجع در شماره جدید
Block a user