AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-299127 - Merge 'develop' into 'amd-staging'

Browse Source 
Change-Id: Idb06f8c83402eaa74de48a30b82d426e58031edf
This commit is contained in:
Jenkins
2021-12-03 07:14:45 -05:00
parent d22c527d2a 610565fe68
commit f50aeef9ec
15 changed files with 180 additions and 11 deletions
Download Patch File Download Diff File Expand all files Collapse all files
docs/markdown/hip_kernel_language.md
+5 -1
View File
@@ -494,7 +494,11 @@ Following is the list of supported floating-point intrinsics. Note that intrinsi
| double __dsqrt_rn ( double x ) <br><sub>Compute `√x` in round-to-nearest-even mode.</sub> |
## Texture Functions
The supported Texture functions are listed in header files "texture_functions.h"(https://github.com/ROCm-Developer-Tools/HIP/blob/main/include/hip/hcc_detail/texture_functions.h) and"texture_indirect_functions.h" (https://github.com/ROCm-Developer-Tools/HIP/blob/main/include/hip/hcc_detail/texture_indirect_functions.h).
The supported Texture functions are listed in header files "texture_fetch_functions.h"(https://github.com/ROCm-Developer-Tools/HIP/blob/main/include/hip/hcc_detail/texture_fetch_functions.h) and"texture_indirect_functions.h" (https://github.com/ROCm-Developer-Tools/HIP/blob/main/include/hip/hcc_detail/texture_indirect_functions.h).
Texture functions are not supported on some devices.
Macro __HIP_NO_IMAGE_SUPPORT == 1 can be used to check whether texture functions are not supported in device code.
Attribute hipDeviceAttributeImageSupport can be queried to check whether texture functions are supported in host runtime code.
## Surface Functions
Surface functions are not supported.
include/hip/driver_types.h
+35 -1
View File
@@ -436,7 +436,41 @@ typedef enum hipFunction_attribute {
HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES,
HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT,
HIP_FUNC_ATTRIBUTE_MAX
}hipFunction_attribute;
} hipFunction_attribute;
typedef enum hipPointer_attribute {
HIP_POINTER_ATTRIBUTE_CONTEXT = 1, ///< The context on which a pointer was allocated
///< @warning - not supported in HIP
HIP_POINTER_ATTRIBUTE_MEMORY_TYPE, ///< memory type describing location of a pointer
HIP_POINTER_ATTRIBUTE_DEVICE_POINTER,///< address at which the pointer is allocated on device
HIP_POINTER_ATTRIBUTE_HOST_POINTER, ///< address at which the pointer is allocated on host
HIP_POINTER_ATTRIBUTE_P2P_TOKENS, ///< A pair of tokens for use with linux kernel interface
///< @warning - not supported in HIP
HIP_POINTER_ATTRIBUTE_SYNC_MEMOPS, ///< Synchronize every synchronous memory operation
///< initiated on this region
HIP_POINTER_ATTRIBUTE_BUFFER_ID, ///< Unique ID for an allocated memory region
HIP_POINTER_ATTRIBUTE_IS_MANAGED, ///< Indicates if the pointer points to managed memory
HIP_POINTER_ATTRIBUTE_DEVICE_ORDINAL,///< device ordinal of a device on which a pointer
///< was allocated or registered
HIP_POINTER_ATTRIBUTE_IS_LEGACY_HIP_IPC_CAPABLE, ///< if this pointer maps to an allocation
///< that is suitable for hipIpcGetMemHandle
///< @warning - not supported in HIP
HIP_POINTER_ATTRIBUTE_RANGE_START_ADDR,///< Starting address for this requested pointer
HIP_POINTER_ATTRIBUTE_RANGE_SIZE, ///< Size of the address range for this requested pointer
HIP_POINTER_ATTRIBUTE_MAPPED, ///< tells if this pointer is in a valid address range
///< that is mapped to a backing allocation
HIP_POINTER_ATTRIBUTE_ALLOWED_HANDLE_TYPES,///< Bitmask of allowed hipmemAllocationHandleType
///< for this allocation @warning - not supported in HIP
HIP_POINTER_ATTRIBUTE_IS_GPU_DIRECT_RDMA_CAPABLE, ///< returns if the memory referenced by
///< this pointer can be used with the GPUDirect RDMA API
///< @warning - not supported in HIP
HIP_POINTER_ATTRIBUTE_ACCESS_FLAGS, ///< Returns the access flags the device associated with
///< for the corresponding memory referenced by the ptr
HIP_POINTER_ATTRIBUTE_MEMPOOL_HANDLE ///< Returns the mempool handle for the allocation if
///< it was allocated from a mempool
///< @warning - not supported in HIP
} hipPointer_attribute;
#endif // !defined(__HIPCC_RTC__)
#else
#error("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
include/hip/hip_runtime_api.h
+39 -5
View File
@@ -442,7 +442,8 @@ typedef enum hipDeviceAttribute_t {
hipDeviceAttributeIsLargeBar, ///< Whether it is LargeBar
hipDeviceAttributeAsicRevision, ///< Revision of the GPU in this device
hipDeviceAttributeCanUseStreamWaitValue, ///< '1' if Device supports hipStreamWaitValue32() and
///< hipStreamWaitValue64() , '0' otherwise.
///< hipStreamWaitValue64(), '0' otherwise.
hipDeviceAttributeImageSupport, ///< '1' if Device supports image, '0' otherwise.
hipDeviceAttributeAmdSpecificEnd = 19999,
hipDeviceAttributeVendorSpecificBegin = 20000,
@@ -1976,15 +1977,48 @@ hipError_t hipEventQuery(hipEvent_t event);
/**
* @brief Return attributes for the specified pointer
*
* @param[out] attributes for the specified pointer
* @param[in] pointer to get attributes for
* @param [out] attributes attributes for the specified pointer
* @param [in] ptr pointer to get attributes for
*
* @return #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue
*
* @see hipGetDeviceCount, hipGetDevice, hipSetDevice, hipChooseDevice
* @see hipPointerGetAttribute
*/
hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void* ptr);
/**
* @brief Returns information about the specified pointer.[BETA]
*
* @param [in, out] data returned pointer attribute value
* @param [in] atribute attribute to query for
* @param [in] ptr pointer to get attributes for
*
* @return #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue
*
* @beta This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*
* @see hipPointerGetAttributes
*/
hipError_t hipPointerGetAttribute(void* data, hipPointer_attribute attribute,
hipDeviceptr_t ptr);
/**
* @brief Returns information about the specified pointer.[BETA]
*
* @param [in] numAttributes number of attributes to query for
* @param [in] attributes attributes to query for
* @param [in, out] data a two-dimensional containing pointers to memory locations
* where the result of each attribute query will be written to
* @param [in] ptr pointer to get attributes for
*
* @return #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue
*
* @beta This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*
* @see hipPointerGetAttribute
*/
hipError_t hipDrvPointerGetAttributes(unsigned int numAttributes, hipPointer_attribute* attributes,
void** data, hipDeviceptr_t ptr);
/**
* @brief Imports an external semaphore.
*
tests/src/runtimeApi/memory/hipMemcpy2DFromArray.cpp
+9
View File
@@ -173,6 +173,7 @@ bool Memcpy2DFromArray<T>::hipMemcpy2DFromArray_PeerDeviceContext() {
err = hipMemcpy2DFromArray(A_h, width, A_d,
0, 0, width,
NUM_H, hipMemcpyDeviceToHost);
if (err == hipSuccess) {
TestPassed = ValidateResult(A_h, INITIAL_VAL);
} else {
@@ -281,7 +282,11 @@ int main(int argc, char **argv) {
TestPassed &= Array_obj.hipMemcpy2DFromArray_SizeCheck();
} else if (p_tests == 3) {
if (numDevices > 1) {
#ifndef _WIN64
TestPassed &= Array_obj.hipMemcpy2DFromArray_PeerDeviceContext();
#else
printf("xgmi memory test not supported on windows\n");
#endif
} else {
printf("skipped the testcase as noof devices <2\n");
}
@@ -289,8 +294,12 @@ int main(int argc, char **argv) {
TestPassed &= Array_obj.hipMemcpy2DFromArray_NegativeTests();
} else if (p_tests == 5) {
if (numDevices > 1) {
#ifndef _WIN64
TestPassed &= Array_obj.hipMemcpy2DFromArray_PinnedHostMemory_SameGPU();
TestPassed &= Array_obj.hipMemcpy2DFromArray_PinnedHostMemory_PeerGPU();
#else
printf("xgmi memory test not supported on windows\n");
#endif
} else {
printf("skipped the testcases as noof devices <2\n");
}
tests/src/runtimeApi/memory/hipMemcpyAtoH.cpp
+2
View File
@@ -209,7 +209,9 @@ int main(int argc, char **argv) {
} else if (p_tests == 2) {
TestPassed = AtoH_obj.hipMemcpyAtoH_ByteCountZero();
} else if (p_tests == 3) {
#ifndef _WIN64
TestPassed = AtoH_obj.hipMemcpyAtoH_PeerDeviceContext();
#endif
} else if (p_tests == 4) {
TestPassed = AtoH_obj.hipMemcpyAtoH_NegativeTests();
} else if (p_tests == 5) {
tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp
+7
View File
@@ -592,6 +592,13 @@ bool testTexSingleStreamMultGPU(unsigned int numOfGPUs,
int main(int argc, char** argv) {
HipTest::parseStandardArguments(argc, argv, true);
int imageSupport = 0;
hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport,
p_gpuDevice);
if (!imageSupport) {
printf("Texture is not support on the device\n");
passed();
}
bool TestPassed = true;
if (p_tests == 0x01) {
TestPassed = testTexType<float>(HIP_AD_FORMAT_FLOAT,
tests/src/runtimeApi/module/tex2d_kernel.cpp
+8
View File
@@ -30,36 +30,44 @@ __device__ float deviceGlobalFloat;
extern "C" __global__ void tex2dKernelFloat(float* outputData,
int width, int height) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
if ((x < width) && (y < width)) {
outputData[y * width + x] = tex2D(ftex, x, y);
}
#endif
}
extern "C" __global__ void tex2dKernelInt(int* outputData,
int width, int height) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
if ((x < width) && (y < width)) {
outputData[y * width + x] = tex2D(itex, x, y);
}
#endif
}
extern "C" __global__ void tex2dKernelInt16(short* outputData,
int width, int height) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
if ((x < width) && (y < width)) {
outputData[y * width + x] = tex2D(stex, x, y);
}
#endif
}
extern "C" __global__ void tex2dKernelInt8(char* outputData,
int width, int height) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
if ((x < width) && (y < width)) {
outputData[y * width + x] = tex2D(ctex, x, y);
}
#endif
}
tests/src/runtimeApi/p2p/hipDeviceGetP2PAttribute.cpp
+4
View File
@@ -42,8 +42,12 @@ int main() {
HIPCHECK(hipGetDeviceCount(&count))
if (count >= 2){
#ifdef _WIN64
std::cout<<"XGMI p2p attribute test no supported on windows"<<std::endl;
#else
for (int dev_idx = 0; dev_idx < (count-1); ++dev_idx) {
runTest(dev_idx, 1 + dev_idx);
#endif
}
} else {
std::cout<<"Not enough GPUs to run the single GPU tests"<<std::endl;
tests/src/runtimeApi/p2p/hipP2pLinkTypeAndHopFunc.cpp
+16
View File
@@ -393,13 +393,29 @@ int main(int argc, char* argv[]) {
printf("This test is skipped due to non linux environment.\n");
#endif
} else if (p_tests == 0x3) {
#ifdef __linux__
TestPassed = testhipInvalidLinkType();
#else
printf("This test is skipped due to non linux environment.\n");
#endif
} else if (p_tests == 0x4) {
#ifdef __linux__
TestPassed = testhipInvalidHopcount();
#else
printf("This test is skipped due to non linux environment.\n");
#endif
} else if (p_tests == 0x5) {
#ifdef __linux__
TestPassed = testhipSameDevice(numDevices);
#else
printf("This test is skipped due to non linux environment.\n");
#endif
} else if (p_tests == 0x6) {
#ifdef __linux__
TestPassed = testhipLinkTypeHopcountDeviceOrderRev(numDevices);
#else
printf("This test is skipped due to non linux environment.\n");
#endif
} else if (p_tests == 0x7) {
/*TODO:This test is currently ommited from directed test due to existing issues
in rocm-smi. Once rocm-smi issues are resolved, this test will be enabled. */
tests/src/texture/hipGetChanDesc.cpp
+7
View File
@@ -36,6 +36,13 @@ using namespace std;
bool runTest(void);
int main(int argc, char** argv) {
int imageSupport = 0;
hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport,
p_gpuDevice);
if (!imageSupport) {
printf("Texture is not support on the device\n");
passed();
}
bool testResult=runTest();
if (testResult) {
tests/src/texture/hipNormalizedFloatValueTex.cpp
+9
View File
@@ -60,6 +60,7 @@ texture<unsigned short, hipTextureType1D, hipReadModeNormalizedFloat> texus;
template<typename T>
__global__ void normalizedValTextureTest(unsigned int numElements, float* pDst)
{
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
unsigned int elementID = hipThreadIdx_x;
if(elementID >= numElements)
return;
@@ -72,6 +73,7 @@ __global__ void normalizedValTextureTest(unsigned int numElements, float* pDst)
pDst[elementID] = tex1D(texs, coord);
else if(std::is_same<T, unsigned short>::value)
pDst[elementID] = tex1D(texus, coord);
#endif
}
bool textureVerifyFilterModePoint(float *hOutputData, float *expected, size_t size) {
@@ -171,6 +173,13 @@ bool runTest() {
int main(int argc, char** argv)
{
int imageSupport = 0;
hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport,
p_gpuDevice);
if (!imageSupport) {
printf("Texture is not support on the device\n");
passed();
}
HipTest::parseStandardArguments(argc, argv, true);
int device = 0;
bool status = false;
tests/src/texture/hipTexObjPitch.cpp
+10 -2
View File
@@ -29,11 +29,12 @@ THE SOFTWARE.
// texture object is a kernel argument
template <typename TYPE_t>
__global__ void texture2dCopyKernel( hipTextureObject_t texObj, TYPE_t* dst,TYPE_t* A) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
for(int i =0;i<SIZE_H;i++)
for(int j = 0;j<SIZE_W;j++)
dst[SIZE_W*i+j] = tex2D<TYPE_t>(texObj, j, i);
__syncthreads();
#endif
}
template <typename TYPE_t>
@@ -76,7 +77,7 @@ void texture2Dtest()
texDescr.readMode = hipReadModeElementType;
hipTextureObject_t texObj;
HIPCHECK( hipCreateTextureObject(&texObj, &texRes, &texDescr, NULL));
HIPCHECK(hipCreateTextureObject(&texObj, &texRes, &texDescr, NULL));
HIPCHECK(hipMalloc((void**)&devPtrB, SIZE_W*sizeof(TYPE_t)*SIZE_H)) ;
@@ -95,6 +96,13 @@ void texture2Dtest()
int main()
{
int imageSupport = 0;
hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport,
p_gpuDevice);
if (!imageSupport) {
printf("Texture is not support on the device\n");
passed();
}
texture2Dtest<float>();
texture2Dtest<int>();
texture2Dtest<unsigned char>();
tests/src/texture/hipTextureMipmapObj2D.cpp
+10 -1
View File
@@ -40,10 +40,12 @@ std::vector<unsigned int> mip_vector = {8, 4, 2, 1};
__global__ void tex2DKernel(float* outputData, hipTextureObject_t textureObject, int width,
int height, float level) {
#ifndef __gfx90a__
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
outputData[y * width + x] = tex2DLod<float>(textureObject, x, y, level);
#endif
#endif
}
bool runMipMapTest(unsigned int width, unsigned int height, unsigned int mipmap_level) {
@@ -148,6 +150,7 @@ bool runMipMapTest(unsigned int width, unsigned int height, unsigned int mipmap_
hipDestroyTextureObject(textureObject);
hipFree(dData);
hipFreeArray(hipArray);
free(hData);
return testResult;
}
@@ -169,7 +172,13 @@ bool runTest(int argc, char** argv) {
int main(int argc, char** argv) {
bool testResult = true;
int imageSupport = 0;
hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport,
p_gpuDevice);
if (!imageSupport) {
printf("Texture is not support on the device\n");
passed();
}
#ifdef _WIN32
testResult = runTest(argc, argv);
#else
tests/src/texture/hipTextureObj2D.cpp
+10 -1
View File
@@ -12,14 +12,23 @@
__global__ void tex2DKernel(float* outputData, hipTextureObject_t textureObject, int width,
int height) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
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);
#endif
}
int runTest(int argc, char** argv);
int main(int argc, char** argv) {
int imageSupport = 0;
hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport,
p_gpuDevice);
if (!imageSupport) {
printf("Texture is not support on the device\n");
passed();
}
int testResult = runTest(argc, argv);
if (testResult) {
@@ -70,7 +79,6 @@ int runTest(int argc, char** argv) {
// Create texture object
hipTextureObject_t textureObject = 0;
hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL);
float* dData = NULL;
hipMalloc((void**)&dData, size);
@@ -104,5 +112,6 @@ int runTest(int argc, char** argv) {
hipDestroyTextureObject(textureObject);
hipFree(dData);
hipFreeArray(hipArray);
free(hData);
return testResult;
}
tests/src/texture/simpleTexture3D.cpp
+9
View File
@@ -42,6 +42,7 @@ __global__ void simpleKernel3DArray(T* outputData,
int width,
int height,int depth)
{
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
for (int i = 0; i < depth; i++) {
for (int j = 0; j < height; j++) {
for (int k = 0; k < width; k++) {
@@ -54,6 +55,7 @@ __global__ void simpleKernel3DArray(T* outputData,
}
}
}
#endif
}
////////////////////////////////////////////////////////////////////////////////
@@ -127,6 +129,13 @@ void runTest(int width,int height,int depth,texture<T, hipTextureType3D, hipRead
////////////////////////////////////////////////////////////////////////////////
int main(int argc, char **argv)
{
int imageSupport = 0;
hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport,
p_gpuDevice);
if (!imageSupport) {
printf("Texture is not support on the device\n");
passed();
}
printf("%s starting...\n", sampleName);
for(int i=1;i<25;i++)
{