AI/rocm-systems
2
0
フォーク 0
コード イシュー プルリクエスト Actions パッケージ プロジェクト リリース Wiki アクティビティ

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

ソースを参照 
Change-Id: I37a15920fd97c85b2c090af8da96d178ee264afc
このコミットが含まれているのは:
Jenkins
2022-07-24 19:10:20 -04:00
親 7cfde0047f 1f11cabcbb
コミット 3aa38b08b4
49個のファイルの変更、2669行の追加、225行の削除
Patchファイルをダウンロード Diffファイルをダウンロード すべてのファイルを展開 すべてのファイルを折り畳む
README.md
+3 -3
ファイルの表示
@@ -70,9 +70,9 @@ Programmers familiar with CUDA will also be able to quickly learn and start codi
Compute kernels are launched with the "hipLaunchKernelGGL" macro call. Here is simple example showing a
snippet of HIP API code:
```cpp
hipMalloc(&A_d, Nbytes));
hipMalloc(&C_d, Nbytes));
```
hipMalloc(&A_d, Nbytes);
hipMalloc(&C_d, Nbytes);
hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice);
include/hip/hip_runtime_api.h
+181 -57
ファイルの表示
@@ -452,6 +452,7 @@ typedef enum hipDeviceAttribute_t {
hipDeviceAttributePhysicalMultiProcessorCount, ///< All available physical compute
///< units for the device
hipDeviceAttributeFineGrainSupport, ///< '1' if Device supports fine grain, '0' otherwise
hipDeviceAttributeWallClockRate, ///< Constant frequency of wall clock in kilohertz.
hipDeviceAttributeAmdSpecificEnd = 19999,
hipDeviceAttributeVendorSpecificBegin = 20000,
@@ -1076,6 +1077,12 @@ typedef struct hipGraphNode* hipGraphNode_t;
*/
typedef struct hipGraphExec* hipGraphExec_t;
/**
* An opaque value that represents a user obj
*/
typedef struct hipUserObject* hipUserObject_t;
/**
* @brief hipGraphNodeType
* @enum
@@ -1188,6 +1195,13 @@ typedef enum hipGraphMemAttributeType {
hipGraphMemAttrReservedMemCurrent, ///< Amount of memory, in bytes, currently allocated for graphs.
hipGraphMemAttrReservedMemHigh, ///< High watermark of memory, in bytes, currently allocated for graphs
}hipGraphMemAttributeType;
typedef enum hipUserObjectFlags {
hipUserObjectNoDestructorSync = 0x1, ///< Destructor execution is not synchronized.
} hipUserObjectFlags;
typedef enum hipUserObjectRetainFlags {
hipGraphUserObjectMove = 0x1, ///< Add new reference or retain.
} hipUserObjectRetainFlags;
typedef enum hipGraphInstantiateFlags {
hipGraphInstantiateFlagAutoFreeOnLaunch =
@@ -2091,7 +2105,7 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback
/**
*-------------------------------------------------------------------------------------------------
*-------------------------------------------------------------------------------------------------
* @defgroup Stream Memory Operations
* @defgroup StreamM Stream Memory Operations
* @{
* This section describes Stream Memory Wait and Write functions of HIP runtime API.
*/
@@ -2590,7 +2604,7 @@ hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int flags);
/**
*-------------------------------------------------------------------------------------------------
*-------------------------------------------------------------------------------------------------
* @addtogroup Memory Managed Memory
* @addtogroup MemoryM Managed Memory
* @{
* @ingroup Memory
* This section describes the managed memory management functions of HIP runtime API.
@@ -2698,7 +2712,7 @@ hipError_t hipStreamAttachMemAsync(hipStream_t stream,
/**
*-------------------------------------------------------------------------------------------------
*-------------------------------------------------------------------------------------------------
* @addtogroup Stream Ordered Memory Allocator
* @addtogroup StreamO Ordered Memory Allocator
* @{
* @ingroup Memory
* This section describes Stream Ordered Memory Allocator functions of HIP runtime API.
@@ -3086,6 +3100,7 @@ hipError_t hipMemPoolImportPointer(
void** dev_ptr,
hipMemPool_t mem_pool,
hipMemPoolPtrExportData* export_data);
// Doxygen end of ordered memory allocator
/**
* @}
*/
@@ -4916,20 +4931,6 @@ hipError_t hipBindTextureToMipmappedArray(
hipMipmappedArray_const_t mipmappedArray,
const hipChannelFormatDesc* desc);
/**
* @brief Gets the texture reference related with the symbol.
*
* @param [out] texref texture reference
* @param [in] symbol pointer to the symbol related with the texture for the reference
*
* @returns hipSuccess, hipErrorInvalidValue
*
*/
DEPRECATED(DEPRECATED_MSG)
hipError_t hipGetTextureReference(
const textureReference** texref,
const void* symbol);
/**
* @brief Creates a texture object.
*
@@ -5010,8 +5011,93 @@ hipError_t hipGetTextureObjectTextureDesc(
hipTextureObject_t textureObject);
/**
* @brief Creates a texture object.
*
* @param [out] pTexObject pointer to texture object to create
* @param [in] pResDesc pointer to resource descriptor
* @param [in] pTexDesc pointer to texture descriptor
* @param [in] pResViewDesc pointer to resource view descriptor
*
* @returns hipSuccess, hipErrorInvalidValue
*
*/
hipError_t hipTexObjectCreate(
hipTextureObject_t* pTexObject,
const HIP_RESOURCE_DESC* pResDesc,
const HIP_TEXTURE_DESC* pTexDesc,
const HIP_RESOURCE_VIEW_DESC* pResViewDesc);
/**
* @brief Destroys a texture object.
*
* @param [in] texObject texture object to destroy
*
* @returns hipSuccess, hipErrorInvalidValue
*
*/
hipError_t hipTexObjectDestroy(
hipTextureObject_t texObject);
/**
* @brief Gets resource descriptor of a texture object.
*
* @param [out] pResDesc pointer to resource descriptor
* @param [in] texObject texture object
*
* @returns hipSuccess, hipErrorNotSupported, hipErrorInvalidValue
*
*/
hipError_t hipTexObjectGetResourceDesc(
HIP_RESOURCE_DESC* pResDesc,
hipTextureObject_t texObject);
/**
* @brief Gets resource view descriptor of a texture object.
*
* @param [out] pResViewDesc pointer to resource view descriptor
* @param [in] texObject texture object
*
* @returns hipSuccess, hipErrorNotSupported, hipErrorInvalidValue
*
*/
hipError_t hipTexObjectGetResourceViewDesc(
HIP_RESOURCE_VIEW_DESC* pResViewDesc,
hipTextureObject_t texObject);
/**
* @brief Gets texture descriptor of a texture object.
*
* @param [out] pTexDesc pointer to texture descriptor
* @param [in] texObject texture object
*
* @returns hipSuccess, hipErrorNotSupported, hipErrorInvalidValue
*
*/
hipError_t hipTexObjectGetTextureDesc(
HIP_TEXTURE_DESC* pTexDesc,
hipTextureObject_t texObject);
/**
*
* @addtogroup TextureD Texture Management [Deprecated]
* @{
* @ingroup Texture
* This section describes the deprecated texture management functions of HIP runtime API.
*/
/**
* @brief Gets the texture reference related with the symbol.
*
* @param [out] texref texture reference
* @param [in] symbol pointer to the symbol related with the texture for the reference
*
* @returns hipSuccess, hipErrorInvalidValue
*
*/
DEPRECATED(DEPRECATED_MSG)
hipError_t hipGetTextureReference(
const textureReference** texref,
const void* symbol);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefSetAddressMode(
textureReference* texRef,
@@ -5035,30 +5121,6 @@ hipError_t hipTexRefSetFormat(
textureReference* texRef,
hipArray_Format fmt,
int NumPackedComponents);
hipError_t hipTexObjectCreate(
hipTextureObject_t* pTexObject,
const HIP_RESOURCE_DESC* pResDesc,
const HIP_TEXTURE_DESC* pTexDesc,
const HIP_RESOURCE_VIEW_DESC* pResViewDesc);
hipError_t hipTexObjectDestroy(
hipTextureObject_t texObject);
hipError_t hipTexObjectGetResourceDesc(
HIP_RESOURCE_DESC* pResDesc,
hipTextureObject_t texObject);
hipError_t hipTexObjectGetResourceViewDesc(
HIP_RESOURCE_VIEW_DESC* pResViewDesc,
hipTextureObject_t texObject);
hipError_t hipTexObjectGetTextureDesc(
HIP_TEXTURE_DESC* pTexDesc,
hipTextureObject_t texObject);
/**
*
* @addtogroup TextureD Texture Management [Deprecated]
* @{
* @ingroup Texture
* This section describes the deprecated texture management functions of HIP runtime API.
*/
DEPRECATED(DEPRECATED_MSG)
hipError_t hipBindTexture(
size_t* offset,
@@ -5145,21 +5207,7 @@ DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefSetMaxAnisotropy(
textureReference* texRef,
unsigned int maxAniso);
// doxygen end deprecated texture management
/**
* @}
*/
// The following are not supported.
/**
*
* @addtogroup TextureU Texture Management [Not supported]
* @{
* @ingroup Texture
* This section describes the texture management functions currently unsupported in HIP runtime.
*/
DEPRECATED(DEPRECATED_MSG)
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefSetBorderColor(
textureReference* texRef,
float* pBorderColor);
@@ -5181,6 +5229,21 @@ hipError_t hipTexRefSetMipmappedArray(
textureReference* texRef,
struct hipMipmappedArray* mipmappedArray,
unsigned int Flags);
// doxygen end deprecated texture management
/**
* @}
*/
// The following are not supported.
/**
*
* @addtogroup TextureU Texture Management [Not supported]
* @{
* @ingroup Texture
* This section describes the texture management functions currently unsupported in HIP runtime.
*/
hipError_t hipMipmappedArrayCreate(
hipMipmappedArray_t* pHandle,
HIP_ARRAY3D_DESCRIPTOR* pMipmappedArrayDesc,
@@ -6243,6 +6306,67 @@ hipError_t hipDeviceSetGraphMemAttribute(int device, hipGraphMemAttributeType at
* it is still open to changes and may have outstanding issues.
*/
hipError_t hipDeviceGraphMemTrim(int device);
/**
* @brief Create an instance of userObject to manage lifetime of a resource.
*
* @param [out] object_out - pointer to instace of userobj.
* @param [in] ptr - pointer to pass to destroy function.
* @param [in] destroy - destroy callback to remove resource.
* @param [in] initialRefcount - reference to resource.
* @param [in] flags - flags passed to API.
* @returns #hipSuccess, #hipErrorInvalidValue
* @warning : This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*/
hipError_t hipUserObjectCreate(hipUserObject_t* object_out, void* ptr, hipHostFn_t destroy, unsigned int initialRefcount, unsigned int flags);
/**
* @brief Release number of references to resource.
*
* @param [in] object - pointer to instace of userobj.
* @param [in] count - reference to resource to be retained.
* @returns #hipSuccess, #hipErrorInvalidValue
* @warning : This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*/
hipError_t hipUserObjectRelease(hipUserObject_t object, unsigned int count);
/**
* @brief Retain number of references to resource.
*
* @param [in] object - pointer to instace of userobj.
* @param [in] count - reference to resource to be retained.
* @returns #hipSuccess, #hipErrorInvalidValue
* @warning : This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*/
hipError_t hipUserObjectRetain(hipUserObject_t object, unsigned int count);
/**
* @brief Retain user object for graphs.
*
* @param [in] graph - pointer to graph to retain the user object for.
* @param [in] object - pointer to instace of userobj.
* @param [in] count - reference to resource to be retained.
* @param [in] flags - flags passed to API.
* @returns #hipSuccess, #hipErrorInvalidValue
* @warning : This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*/
hipError_t hipGraphRetainUserObject(hipGraph_t graph, hipUserObject_t object, unsigned int count, unsigned int flags);
/**
* @brief Release user object from graphs.
*
* @param [in] graph - pointer to graph to retain the user object for.
* @param [in] object - pointer to instace of userobj.
* @param [in] count - reference to resource to be retained.
* @returns #hipSuccess, #hipErrorInvalidValue
* @warning : This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*/
hipError_t hipGraphReleaseUserObject(hipGraph_t graph, hipUserObject_t object, unsigned int count);
// doxygen end graph API
/**
* @}
@@ -6349,7 +6473,7 @@ typedef struct hipArrayMapInfo {
/**
*-------------------------------------------------------------------------------------------------
*-------------------------------------------------------------------------------------------------
* @defgroup Virtual Memory Management
* @defgroup Virtual Virtual Memory Management
* @{
* This section describes the virtual memory management functions of HIP runtime API.
*/
tests/catch/hipTestMain/config/config_amd_linux.json
+1 -4
ファイルの表示
@@ -1,10 +1,7 @@
{
"DisabledTests":
[
"# Following test is related to ticket EXSWCPHIPT-41",
"Unit_hipStreamGetPriority_happy",
"Unit_hipStreamPerThread_DeviceReset_1",
"Unit_hipGraphAddKernelNode_Negative"
"Unit_hipMallocManaged_OverSubscription"
]
}
tests/catch/hipTestMain/config/config_amd_windows.json
+5 -7
ファイルの表示
@@ -7,8 +7,6 @@
"Unit_hipTextureMipmapObj2D_Check",
"Unit_hipManagedKeyword_MultiGpu",
"Unit_hipGraphAddHostNode_ClonedGraphwithHostNode",
"Unit_hipGraphAddChildGraphNode_OrgGraphAsChildGraph",
"Unit_hipGraphAddChildGraphNode_SingleChildNode",
"Unit_hipEventIpc",
"Unit_hipMalloc3D_Negative",
"Unit_hipPointerGetAttribute_MappedMem",
@@ -27,13 +25,9 @@
"Unit_hipMemPoolApi_Opportunistic",
"Unit_hipMemPoolApi_Default",
"Unit_hipDeviceGetUuid",
"Unit_hipGraphMemcpyNodeSetParamsFromSymbol_Functional",
"Unit_hipGraphExecEventWaitNodeSetEvent_Negative",
"Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory",
"Unit_hipGraphMemcpyNodeSetParams_Functional",
"Unit_hipMalloc3D_ValidatePitch",
"Unit_hipArrayCreate_happy",
"Unit_hipGraphAddKernelNode_Negative",
"Unit_hipHostRegister_Negative - int",
"Unit_hipHostRegister_Negative - float",
"Unit_hipHostRegister_Negative - double",
@@ -46,6 +40,10 @@
"Unit_hipArrayCreate_happy - float",
"Unit_hipArrayCreate_happy - float2",
"Unit_hipArrayCreate_happy - float4",
"Unit_hipMemVmm_Basic"
"Unit_hipMemVmm_Basic",
"Unit_hipGraphKernelNodeSetParams_Functional",
"Unit_hipGraphExecKernelNodeSetParams_Negative",
"Unit_hipGraphExecKernelNodeSetParams_Functional",
"Unit_hipGraphMemcpyNodeSetParams1D_Functional"
]
}
tests/catch/include/hip_test_checkers.hh
+7
ファイルの表示
@@ -98,6 +98,13 @@ size_t checkVectorADD(T* A_h, T* B_h, T* result_H, size_t N, bool expectMatch =
A_h, B_h, result_H, N, [](T a, T b) { return a + b; }, expectMatch, reportMismatch);
}
template <typename T>
size_t checkVectorSUB(T* A_h, T* B_h, T* result_H, size_t N, bool expectMatch = true,
bool reportMismatch = true) {
return checkVectors<T>(
A_h, B_h, result_H, N, [](T a, T b) { return a - b; }, expectMatch, reportMismatch);
}
template <typename T>
void checkTest(T* expected_H, T* result_H, size_t N, bool expectMatch = true) {
checkVectors<T>(
tests/catch/include/hip_test_common.hh
+62 -6
ファイルの表示
@@ -203,7 +203,7 @@ static inline int RAND_R(unsigned* rand_seed) {
inline bool isImageSupported() {
int imageSupport = 1;
#ifdef __HIP_PLATFORM_AMD__
#if HT_AMD
int device;
HIP_CHECK(hipGetDevice(&device));
HIPCHECK(hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport, device));
@@ -271,13 +271,69 @@ void launchKernel(K kernel, Dim numBlocks, Dim numThreads, std::uint32_t memPerB
std::forward<Args>(packedArgs)...);
#endif
}
//---
struct Pinned {
static const bool isPinned = true;
static const char* str() { return "Pinned"; };
static void* Alloc(size_t sizeBytes) {
void* p;
HIPCHECK(hipHostMalloc((void**)&p, sizeBytes));
return p;
};
};
//---
struct Unpinned {
static const bool isPinned = false;
static const char* str() { return "Unpinned"; };
static void* Alloc(size_t sizeBytes) {
void* p = malloc(sizeBytes);
HIPASSERT(p);
return p;
};
};
struct Memcpy {
static const char* str() { return "Memcpy"; };
};
struct MemcpyAsync {
static const char* str() { return "MemcpyAsync"; };
};
template <typename C>
struct MemTraits;
template <>
struct MemTraits<Memcpy> {
static void Copy(void* dest, const void* src, size_t sizeBytes, hipMemcpyKind kind,
hipStream_t stream) {
(void)stream;
HIPCHECK(hipMemcpy(dest, src, sizeBytes, kind));
}
};
template <>
struct MemTraits<MemcpyAsync> {
static void Copy(void* dest, const void* src, size_t sizeBytes, hipMemcpyKind kind,
hipStream_t stream) {
HIPCHECK(hipMemcpyAsync(dest, src, sizeBytes, kind, stream));
}
};
} // namespace HipTest
// This must be called in the beginning of image test app's main() to indicate whether image
// is supported.
#define checkImageSupport() \
if (!HipTest::isImageSupported()) { \
printf("Texture is not support on the device. Skipped.\n"); \
return; \
}
#define CHECK_IMAGE_SUPPORT \
if (!HipTest::isImageSupported()) \
{ INFO("Texture is not support on the device. Skipped."); return; }
tests/catch/include/hip_test_kernels.hh
+8
ファイルの表示
@@ -34,6 +34,14 @@ template <typename T> __global__ void vectorADD(const T* A_d, const T* B_d, T* C
}
}
template <typename T> __global__ void vectorSUB(const T* A_d, const T* B_d, T* C_d, size_t NELEM) {
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x;
for (size_t i = offset; i < NELEM; i += stride) {
C_d[i] = A_d[i] - B_d[i];
}
}
template <typename T>
__global__ void vectorADDReverse(const T* A_d, const T* B_d, T* C_d, size_t NELEM) {
tests/catch/include/hip_texture_helper.hh
+227
ファイルの表示
@@ -0,0 +1,227 @@
#pragma once
#define HIP_SAMPLING_VERIFY_EPSILON 0.00001
// The internal precision varies by the GPU family and sometimes within the family.
// Thus the following threshold is subject to change.
#define HIP_SAMPLING_VERIFY_RELATIVE_THRESHOLD 0.05 // 5% for filter mode
#define HIP_SAMPLING_VERIFY_ABSOLUTE_THRESHOLD 0.1
template<typename type, hipTextureFilterMode fMode = hipFilterModePoint>
bool hipTextureSamplingVerify(const type outputData, const type expected) {
bool testResult = false;
if (fMode == hipFilterModePoint) {
testResult = outputData == expected;
} else if (fMode == hipFilterModeLinear) {
const type mean = (fabs(outputData) + fabs(expected)) / 2;
const type diff = fabs(outputData - expected);
const type ratio = diff / (mean + HIP_SAMPLING_VERIFY_EPSILON);
if (ratio <= HIP_SAMPLING_VERIFY_RELATIVE_THRESHOLD) {
testResult = true;
} else if (diff <= HIP_SAMPLING_VERIFY_ABSOLUTE_THRESHOLD) {
// Some small outputs have big ratio due to float operation difference of ALU and GPU
testResult = true;
}
}
return testResult;
}
// Simulate CTS static AddressingTable sAddressingTable
template<hipTextureAddressMode addressMode>
void hipTextureGetAddress(int &value, const int maxValue)
{
switch(addressMode)
{
case hipAddressModeClamp:
value = value < 0 ? 0
: (value > maxValue - 1 ? maxValue - 1 : value);
break;
case hipAddressModeBorder:
value = value < -1 ? -1
: (value > maxValue ? maxValue : value);
break;
default:
break;
}
}
// Simulate logics in CTS read_image_pixel_float().
// x, y and z must be returned by hipTextureGetAddress()
template<hipTextureAddressMode addressMode>
float hipTextureGetValue(const float *data, const int x, const int width,
const int y = 0, const int height = 0,const int z = 0, const int depth = 0) {
float result = std::numeric_limits<float>::lowest();
switch (addressMode) {
case hipAddressModeClamp:
if (width > 0) {
if (height == 0 && depth == 0) {
result = data[x]; // 1D
} else if (depth == 0) {
result = data[y * width + x]; // 2D
} else {
result = data[z * width * height + y * width + x]; // 3D
}
}
break;
case hipAddressModeBorder:
if (width > 0) {
if (height == 0 && depth == 0) {
result = (x >= 0 && x < width) ? data[x] : 0; // 1D
} else if (depth == 0) {
result = (x >= 0 && x < width && y >= 0 && y < height) ?
data[y * width + x] : 0; // 2D
} else {
result = (x >= 0 && x < width && y >= 0 && y < height && z >= 0 && z < depth) ?
data[z * width * height + y * width + x] : 0; // 3D
}
}
break;
default:
break;
}
return result;
}
template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode>
float getExpectedValue(const int width, float x, const float *data) {
float result = std::numeric_limits<float>::lowest();
switch (filterMode) {
case hipFilterModePoint: {
int i1 = static_cast<int>(floor(x));
hipTextureGetAddress < addressMode > (i1, width);
result = hipTextureGetValue < addressMode > (data, i1, width);
}
break;
case hipFilterModeLinear: {
x -= 0.5;
int i1 = static_cast<int>(floor(x));
int i2 = i1 + 1;
float a = x - i1;
hipTextureGetAddress < addressMode > (i1, width);
hipTextureGetAddress < addressMode > (i2, width);
float t1 = hipTextureGetValue < addressMode > (data, i1, width);
float t2 = hipTextureGetValue < addressMode > (data, i2, width);
return (1 - a) * t1 + a * t2;
}
break;
}
return result;
}
template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode>
float getExpectedValue(const int width, const int height, float x, float y, const float *data) {
float result = std::numeric_limits<float>::lowest();
switch (filterMode) {
case hipFilterModePoint: {
int i1 = static_cast<int>(floor(x));
int j1 = static_cast<int>(floor(y));
hipTextureGetAddress < addressMode > (i1, width);
hipTextureGetAddress < addressMode > (j1, height);
result = hipTextureGetValue < addressMode > (data, i1, width, j1, height);
}
break;
case hipFilterModeLinear: {
x -= 0.5;
y -= 0.5;
int i1 = static_cast<int>(floor(x));
int j1 = static_cast<int>(floor(y));
int i2 = i1 + 1;
int j2 = j1 + 1;
float a = x - i1;
float b = y - j1;
hipTextureGetAddress < addressMode > (i1, width);
hipTextureGetAddress < addressMode > (i2, width);
hipTextureGetAddress < addressMode > (j1, height);
hipTextureGetAddress < addressMode > (j2, height);
float t11 = hipTextureGetValue < addressMode
> (data, i1, width, j1, height);
float t21 = hipTextureGetValue < addressMode
> (data, i2, width, j1, height);
float t12 = hipTextureGetValue < addressMode
> (data, i1, width, j2, height);
float t22 = hipTextureGetValue < addressMode
> (data, i2, width, j2, height);
result = (1 - a) * (1 - b) * t11 + a * (1 - b) * t21 + (1 - a) * b * t12
+ a * b * t22;
}
break;
}
return result;
}
template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode>
float getExpectedValue(const int width, const int height, const int depth,
float x, float y, float z, const float *data) {
float result = std::numeric_limits<float>::lowest();
switch (filterMode) {
case hipFilterModePoint: {
int i1 = static_cast<int>(floor(x));
int j1 = static_cast<int>(floor(y));
int k1 = static_cast<int>(floor(z));
hipTextureGetAddress < addressMode > (i1, width);
hipTextureGetAddress < addressMode > (j1, height);
hipTextureGetAddress < addressMode > (k1, depth);
result = hipTextureGetValue < addressMode > (data, i1, width, j1, height, k1, depth);
}
break;
case hipFilterModeLinear: {
x -= 0.5;
y -= 0.5;
z -= 0.5;
int i1 = static_cast<int>(floor(x));
int j1 = static_cast<int>(floor(y));
int k1 = static_cast<int>(floor(z));
int i2 = i1 + 1;
int j2 = j1 + 1;
int k2 = k1 + 1;
float a = x - i1;
float b = y - j1;
float c = z - k1;
hipTextureGetAddress < addressMode > (i1, width);
hipTextureGetAddress < addressMode > (i2, width);
hipTextureGetAddress < addressMode > (j1, height);
hipTextureGetAddress < addressMode > (j2, height);
hipTextureGetAddress < addressMode > (k1, depth);
hipTextureGetAddress < addressMode > (k2, depth);
float t111 = hipTextureGetValue < addressMode
> (data, i1, width, j1, height, k1, depth);
float t211 = hipTextureGetValue < addressMode
> (data, i2, width, j1, height, k1, depth);
float t121 = hipTextureGetValue < addressMode
> (data, i1, width, j2, height, k1, depth);
float t112 = hipTextureGetValue < addressMode
> (data, i1, width, j1, height, k2, depth);
float t122 = hipTextureGetValue < addressMode
> (data, i1, width, j2, height, k2, depth);
float t212 = hipTextureGetValue < addressMode
> (data, i2, width, j1, height, k2, depth);
float t221 = hipTextureGetValue < addressMode
> (data, i2, width, j2, height, k1, depth);
float t222 = hipTextureGetValue < addressMode
> (data, i2, width, j2, height, k2, depth);
result =
(1 - a) * (1 - b) * (1 - c) * t111 + a * (1 - b) * (1 - c) * t211 +
(1 - a) * b * (1 - c) * t121 + a * b * (1 - c) * t221 +
(1 - a) * (1 - b) * c * t112 + a * (1 - b) * c * t212 +
(1 - a) * b * c * t122 + a * b * c * t222;
}
break;
}
return result;
}
tests/catch/unit/CMakeLists.txt
+1
ファイルの表示
@@ -30,3 +30,4 @@ add_subdirectory(printf)
add_subdirectory(texture)
add_subdirectory(streamperthread)
add_subdirectory(kernel)
add_subdirectory(multiThread)
tests/catch/unit/device/hipSetGetDevice.cc
+75 -5
ファイルの表示
@@ -23,21 +23,91 @@
*/
#include <hip_test_common.hh>
#include <thread>
TEST_CASE("Unit_hipSetDevice_BasicSetGet") {
int numDevices = 0;
int device;
int validateCount = 0;
int device{};
HIP_CHECK(hipGetDeviceCount(&numDevices));
REQUIRE(numDevices != 0);
for (int i = 0; i < numDevices; i++) {
HIP_CHECK(hipSetDevice(i));
HIP_CHECK(hipGetDevice(&device));
if (device == i) {
validateCount+= 1;
REQUIRE(device == i);
// Check for hipDevice_t as well
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, i));
}
}
TEST_CASE("Unit_hipGetSetDevice_MultiThreaded") {
auto maxThreads = std::thread::hardware_concurrency();
auto deviceCount = HipTest::getDeviceCount();
auto thread = [&]() {
for (int i = 0; i < deviceCount; i++) {
HIP_CHECK_THREAD(hipSetDevice(i));
int get = -1;
HIP_CHECK_THREAD(hipGetDevice(&get));
REQUIRE_THREAD(get == i);
// check hipDeviceGet
hipDevice_t device;
HIP_CHECK_THREAD(hipDeviceGet(&device, i));
// Alloc some memory and set it
unsigned int* ptr{nullptr};
HIP_CHECK_THREAD(hipMalloc(&ptr, sizeof(unsigned int)));
REQUIRE_THREAD(ptr != nullptr);
HIP_CHECK_THREAD(hipMemset(ptr, 0x0A, sizeof(unsigned int)));
int res{0};
HIP_CHECK_THREAD(hipMemcpy(&res, ptr, sizeof(unsigned int), hipMemcpyDeviceToHost));
REQUIRE_THREAD(res == 0x0A0A0A0A);
HIP_CHECK_THREAD(hipFree(ptr));
}
};
std::vector<std::thread> pool;
pool.reserve(maxThreads);
for (unsigned i = 0; i < maxThreads; i++) {
pool.emplace_back(std::thread(thread));
}
REQUIRE(numDevices == validateCount);
for (auto& i : pool) {
i.join();
}
HIP_CHECK_THREAD_FINALIZE();
}
TEST_CASE("Unit_hipSetGetDevice_Negative") {
SECTION("Get Device - nullptr") { HIP_CHECK_ERROR(hipGetDevice(nullptr), hipErrorInvalidValue); }
SECTION("Set Device - -1") { HIP_CHECK_ERROR(hipSetDevice(-1), hipErrorInvalidDevice); }
SECTION("Set Device - NumDevices + 1") {
HIP_CHECK_ERROR(hipSetDevice(HipTest::getDeviceCount()), hipErrorInvalidDevice);
}
}
TEST_CASE("Unit_hipDeviceGet_Negative") {
// TODO enable after EXSWCPHIPT-104 is fixed
#if HT_NVIDIA
HIP_CHECK(hipInit(0));
SECTION("Nullptr as handle") { HIP_CHECK_ERROR(hipDeviceGet(nullptr, 0), hipErrorInvalidValue); }
#endif
SECTION("Out of bound ordial - positive") {
hipDevice_t device{};
auto totalDevices = HipTest::getDeviceCount();
HIP_CHECK_ERROR(hipDeviceGet(&device, totalDevices), hipErrorInvalidDevice);
}
SECTION("Out of bound ordial - negative") {
hipDevice_t device{};
HIP_CHECK_ERROR(hipDeviceGet(&device, -1), hipErrorInvalidDevice);
}
}
tests/catch/unit/event/Unit_hipEvent_Negative.cc
-6
ファイルの表示
@@ -49,12 +49,6 @@ TEST_CASE("Unit_hipEventDestroy_NullCheck") {
TEST_CASE("Unit_hipEventCreate_IncompatibleFlags") {
hipEvent_t event;
#if HT_AMD
HipTest::HIP_SKIP_TEST("EXSWCPHIPT-106");
return;
#endif
HIP_CHECK_ERROR(hipEventCreateWithFlags(&event, hipEventInterprocess), hipErrorInvalidValue);
#if HT_AMD
tests/catch/unit/graph/CMakeLists.txt
+5
ファイルの表示
@@ -68,6 +68,11 @@ set(TEST_SRC
hipGraphAddKernelNode.cc
hipGraphMemcpyNodeGetParams.cc
hipGraphMemcpyNodeSetParams.cc
hipGraphKernelNodeGetParams.cc
hipGraphKernelNodeSetParams.cc
hipGraphExecKernelNodeSetParams.cc
hipGraphLaunch.cc
hipGraphMemcpyNodeSetParams1D.cc
)
hip_add_exe_to_target(NAME GraphsTest
tests/catch/unit/graph/hipGraphExecKernelNodeSetParams.cc
+186
ファイルの表示
@@ -0,0 +1,186 @@
/*
Copyright (c) 2022 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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/**
Testcase Scenarios :
Negative -
1) Pass hGraphExec as nullptr and verify api returns error code.
2) Pass node as nullptr and verify api returns error code.
3) Pass NodeParams as un-initialized structure object and verify api returns error code.
4) Pass pNodeParams as nullptr and verify api returns error code.
5) Pass NodeParams:func datamember as nullptr and verify api returns error code.
Functional -
1) Instantiate a graph with kernel node, obtain executable graph and update
the kernel node params with set and check it is taking effect.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
/**
* Negative Test for API hipGraphExecKernelNodeSetParams
*/
TEST_CASE("Unit_hipGraphExecKernelNodeSetParams_Negative") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(int);
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
hipGraph_t graph;
hipError_t ret;
hipGraphNode_t memcpyNode, kNode{};
hipKernelNodeParams kNodeParams{};
hipStream_t streamForGraph;
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
std::vector<hipGraphNode_t> dependencies;
hipGraphExec_t graphExec;
size_t NElem{N};
HIP_CHECK(hipStreamCreate(&streamForGraph));
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
dependencies.push_back(memcpyNode);
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, B_d, B_h,
Nbytes, hipMemcpyHostToDevice));
dependencies.push_back(memcpyNode);
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
kNodeParams.gridDim = dim3(blocks);
kNodeParams.blockDim = dim3(threadsPerBlock);
kNodeParams.sharedMemBytes = 0;
kNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
kNodeParams.extra = nullptr;
// Instantiate and launch the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
SECTION("Pass hipGraphExec as nullptr") {
ret = hipGraphExecKernelNodeSetParams(nullptr, kNode, &kNodeParams);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass Node as nullptr") {
ret = hipGraphExecKernelNodeSetParams(graphExec, nullptr, &kNodeParams);
REQUIRE(hipErrorInvalidValue == ret);
}
#if HT_AMD
/* NodeParams null check is disabled on Nvedia as
* this call gives SIGSEGV error in CUDA setup */
SECTION("Pass NodeParams as nullptr") {
ret = hipGraphExecKernelNodeSetParams(graphExec, kNode, nullptr);
REQUIRE(hipErrorInvalidValue == ret);
}
#endif
/* For below 2 scenarios -
In AMD setup this API return - hipErrorInvalidValue and
In CUDA setup this API return - hipErrorInvalidDeviceFunction
As per Cuda spec API can only return "cudaSuccess, cudaErrorInvalidValue".
*/
SECTION("Pass NodeParams as un-initialized structure object") {
hipKernelNodeParams kNodeParams1{};
ret = hipGraphExecKernelNodeSetParams(graphExec, kNode, &kNodeParams1);
REQUIRE(hipSuccess != ret);
}
SECTION("Pass NodeParams func datamember as nullptr") {
kNodeParams.func = nullptr;
ret = hipGraphExecKernelNodeSetParams(graphExec, kNode, &kNodeParams);
REQUIRE(hipSuccess != ret);
}
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(streamForGraph));
}
/**
* Functional Test for API Exec Kernel Params
*/
TEST_CASE("Unit_hipGraphExecKernelNodeSetParams_Functional") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(int);
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
hipGraph_t graph;
hipGraphNode_t memcpyNode, kNode;
hipKernelNodeParams kNodeParams{}, kNodeParams1{};
hipStream_t streamForGraph;
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
std::vector<hipGraphNode_t> dependencies;
hipGraphExec_t graphExec;
size_t NElem{N};
HIP_CHECK(hipStreamCreate(&streamForGraph));
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
dependencies.push_back(memcpyNode);
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, B_d, B_h,
Nbytes, hipMemcpyHostToDevice));
dependencies.push_back(memcpyNode);
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
kNodeParams.gridDim = dim3(blocks);
kNodeParams.blockDim = dim3(threadsPerBlock);
kNodeParams.sharedMemBytes = 0;
kNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
kNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, dependencies.data(),
dependencies.size(), &kNodeParams));
memset(&kNodeParams1, 0, sizeof(kNodeParams1));
kNodeParams1.func = reinterpret_cast<void *>(HipTest::vectorSUB<int>);
kNodeParams1.gridDim = dim3(blocks);
kNodeParams1.blockDim = dim3(threadsPerBlock);
kNodeParams1.sharedMemBytes = 0;
kNodeParams1.kernelParams = reinterpret_cast<void**>(kernelArgs);
kNodeParams1.extra = nullptr;
dependencies.clear();
dependencies.push_back(kNode);
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(),
dependencies.size(), C_h, C_d,
Nbytes, hipMemcpyDeviceToHost));
// Instantiate and launch the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
REQUIRE(hipSuccess == hipGraphExecKernelNodeSetParams(graphExec, kNode,
&kNodeParams1));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
// Verify graph execution result
HipTest::checkVectorSUB<int>(A_h, B_h, C_h, N);
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(streamForGraph));
}
tests/catch/unit/graph/hipGraphKernelNodeGetParams.cc
+159
ファイルの表示
@@ -0,0 +1,159 @@
/*
Copyright (c) 2022 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.
*/
/**
Testcase Scenarios :
Negative -
1) Pass node as nullptr and verify api returns error code.
2) Pass pNodeParams as nullptr and verify api returns error code.
Functional -
1) Create a graph, add kernel node to graph with desired kernel node params.
Verify api fetches the node params mentioned while adding kernel node.
2) Set kernel node params with hipGraphKernelNodeSetParams,
now get the params and verify both are same.
*/
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
#define THREADS_PER_BLOCK 512
/* Test verifies hipGraphKernelNodeGetParams API Negative scenarios.
*/
TEST_CASE("Unit_hipGraphKernelNodeGetParams_Negative") {
constexpr int N = 1024;
size_t NElem{N};
int *A_d, *B_d, *C_d;
hipError_t ret;
hipGraph_t graph;
hipGraphNode_t kNode;
hipKernelNodeParams kNodeParams{};
HIP_CHECK(hipMalloc(&A_d, sizeof(int) * N));
HIP_CHECK(hipMalloc(&B_d, sizeof(int) * N));
HIP_CHECK(hipMalloc(&C_d, sizeof(int) * N));
HIP_CHECK(hipGraphCreate(&graph, 0));
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
kNodeParams.gridDim = dim3(N / THREADS_PER_BLOCK, 1, 1);
kNodeParams.blockDim = dim3(THREADS_PER_BLOCK, 1, 1);
kNodeParams.sharedMemBytes = 0;
kNodeParams.kernelParams = reinterpret_cast<void **>(kernelArgs);
kNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams));
SECTION("Pass node as nullptr") {
ret = hipGraphKernelNodeGetParams(nullptr, &kNodeParams);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass kNodeParams as nullptr") {
ret = hipGraphKernelNodeGetParams(kNode, nullptr);
REQUIRE(hipErrorInvalidValue == ret);
}
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(B_d));
HIP_CHECK(hipFree(C_d));
HIP_CHECK(hipGraphDestroy(graph));
}
static bool dim3_compare(dim3 node1, dim3 node2) {
if ((node1.x == node2.x) && (node1.y == node2.y) && (node1.z == node2.z))
return true;
else
return false;
}
static bool kernelParam_compare(void **p1, void ** p2) {
for (int i = 0; i < 4; i++) {
if (*reinterpret_cast<int *>(p1[i]) != *reinterpret_cast<int *>(p2[i]))
return false;
}
return true;
}
static bool node_compare(hipKernelNodeParams *kNode1,
hipKernelNodeParams *kNode2) {
if (!dim3_compare(kNode1->blockDim, kNode2->blockDim))
return false;
if (kNode1->extra != kNode2->extra)
return false;
if (kNode1->func != kNode2->func)
return false;
if (!dim3_compare(kNode1->gridDim, kNode2->gridDim))
return false;
if (!kernelParam_compare(kNode1->kernelParams, kNode2->kernelParams))
return false;
if (kNode1->sharedMemBytes != kNode2->sharedMemBytes)
return false;
return true;
}
/* Test verifies hipGraphKernelNodeGetParams API Functional scenarios.
*/
TEST_CASE("Unit_hipGraphKernelNodeGetParams_Functional") {
constexpr int N = 1024;
size_t NElem{N};
int *A_d, *B_d, *C_d;
hipGraph_t graph;
hipGraphNode_t kNode;
hipKernelNodeParams kNodeParams{};
HIP_CHECK(hipMalloc(&A_d, sizeof(int) * N));
HIP_CHECK(hipMalloc(&B_d, sizeof(int) * N));
HIP_CHECK(hipMalloc(&C_d, sizeof(int) * N));
HIP_CHECK(hipGraphCreate(&graph, 0));
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kNodeParams.func = reinterpret_cast<void*>(HipTest::vectorADD<int>);
kNodeParams.gridDim = dim3(N / THREADS_PER_BLOCK, 1, 1);
kNodeParams.blockDim = dim3(THREADS_PER_BLOCK, 1, 1);
kNodeParams.sharedMemBytes = 0;
kNodeParams.kernelParams = reinterpret_cast<void **>(kernelArgs);
kNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams));
SECTION("Get Kernel Param and verify.") {
hipKernelNodeParams kNodeGetParams;
HIP_CHECK(hipGraphKernelNodeGetParams(kNode, &kNodeGetParams));
REQUIRE(true == node_compare(&kNodeParams, &kNodeGetParams));
}
SECTION("Set kernel node params then Get Kernel Param and verify.") {
hipKernelNodeParams kNodeParams1;
kNodeParams1.func =
reinterpret_cast<void *>(HipTest::vectorADDReverse<int>);
kNodeParams1.gridDim = dim3(N / THREADS_PER_BLOCK, 1, 1);
kNodeParams1.blockDim = dim3(THREADS_PER_BLOCK, 1, 1);
kNodeParams1.sharedMemBytes = 0;
kNodeParams1.kernelParams = reinterpret_cast<void **>(kernelArgs);
kNodeParams1.extra = nullptr;
hipKernelNodeParams kNodeGetParams1;
HIP_CHECK(hipGraphKernelNodeSetParams(kNode, &kNodeParams1));
HIP_CHECK(hipGraphKernelNodeGetParams(kNode, &kNodeGetParams1));
REQUIRE(true == node_compare(&kNodeParams1, &kNodeGetParams1));
}
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(B_d));
HIP_CHECK(hipFree(C_d));
HIP_CHECK(hipGraphDestroy(graph));
}
tests/catch/unit/graph/hipGraphKernelNodeSetParams.cc
+145
ファイルの表示
@@ -0,0 +1,145 @@
/*
Copyright (c) 2022 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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/**
Testcase Scenarios :
Negative -
1) Pass node as nullptr and verify api returns error code.
2) Pass pNodeParams as nullptr and verify api returns error code.
Functional -
1) Add kernel node to graph with certain kernel params, now update the kernel
node params with set and check taking effect after launching graph.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
/* Test verifies hipGraphKernelNodeSetParams API Negative scenarios.
*/
TEST_CASE("Unit_hipGraphKernelNodeSetParams_Negative") {
constexpr int N = 1024;
size_t NElem{N};
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
int *A_d, *B_d, *C_d;
hipError_t ret;
hipGraph_t graph;
hipGraphNode_t kNode;
hipKernelNodeParams kNodeParams{};
HIP_CHECK(hipMalloc(&A_d, sizeof(int) * N));
HIP_CHECK(hipMalloc(&B_d, sizeof(int) * N));
HIP_CHECK(hipMalloc(&C_d, sizeof(int) * N));
HIP_CHECK(hipGraphCreate(&graph, 0));
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
kNodeParams.gridDim = dim3(blocks);
kNodeParams.blockDim = dim3(threadsPerBlock);
kNodeParams.sharedMemBytes = 0;
kNodeParams.kernelParams = reinterpret_cast<void **>(kernelArgs);
kNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams));
SECTION("Pass node as nullptr") {
ret = hipGraphKernelNodeSetParams(nullptr, &kNodeParams);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass kNodeParams as nullptr") {
ret = hipGraphKernelNodeSetParams(kNode, nullptr);
REQUIRE(hipErrorInvalidValue == ret);
}
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(B_d));
HIP_CHECK(hipFree(C_d));
HIP_CHECK(hipGraphDestroy(graph));
}
/**
* Functional Test for API Set Kernel Params
*/
TEST_CASE("Unit_hipGraphKernelNodeSetParams_Functional") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(int);
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
hipGraph_t graph;
hipGraphNode_t memcpyNode, kNode;
hipKernelNodeParams kNodeParams{}, kNodeParams1{};
hipStream_t streamForGraph;
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
std::vector<hipGraphNode_t> dependencies;
hipGraphExec_t graphExec;
size_t NElem{N};
HIP_CHECK(hipStreamCreate(&streamForGraph));
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
dependencies.push_back(memcpyNode);
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, B_d, B_h,
Nbytes, hipMemcpyHostToDevice));
dependencies.push_back(memcpyNode);
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
kNodeParams.gridDim = dim3(blocks);
kNodeParams.blockDim = dim3(threadsPerBlock);
kNodeParams.sharedMemBytes = 0;
kNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
kNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, dependencies.data(),
dependencies.size(), &kNodeParams));
kNodeParams1.func = reinterpret_cast<void *>(HipTest::vectorSUB<int>);
kNodeParams1.gridDim = dim3(blocks);
kNodeParams1.blockDim = dim3(threadsPerBlock);
kNodeParams1.sharedMemBytes = 0;
kNodeParams1.kernelParams = reinterpret_cast<void**>(kernelArgs);
kNodeParams1.extra = nullptr;
HIP_CHECK(hipGraphKernelNodeSetParams(kNode, &kNodeParams1));
dependencies.clear();
dependencies.push_back(kNode);
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(),
dependencies.size(), C_h, C_d,
Nbytes, hipMemcpyDeviceToHost));
// Instantiate and launch the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
// Verify graph execution result
HipTest::checkVectorSUB<int>(A_h, B_h, C_h, N);
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(streamForGraph));
}
tests/catch/unit/graph/hipGraphLaunch.cc
+307
ファイルの表示
@@ -0,0 +1,307 @@
/*
Copyright (c) 2022 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.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
/* Test verifies hipGraphLaunch API
Negative scenarios -
1) Pass graphExec as nullptr and verify api returns error code.
2) Pass pGraphExec as nullptr and stream as hipStreamPerThread and verify api returns error code.
3) Pass pGraphExec as empty object and verify api returns error code.
4) Destroy executable graph and try to launch it. Make sure api should not crash and it should returns error code.
5) Destroy stream and try to launch respective executable graph. Make sure api should not crash and it should returns error code.
6) Destroy actual graph created and try to launch respective executable graph.
Check api should execute properly without crash or error code.
Functional Scenario -
1) Check basic functionality with stream as hipStreamPerThread
2) Test hipGraphLaunch call on multiple devices.
*/
TEST_CASE("Unit_hipGraphLaunch_Negative") {
hipError_t ret;
SECTION("Pass pGraphExec as nullptr") {
hipStream_t stream{};
ret = hipGraphLaunch(nullptr, stream);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass pGraphExec as nullptr and stream as hipStreamPerThread") {
ret = hipGraphLaunch(nullptr, hipStreamPerThread);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass pGraphExec as empty object") {
hipGraphExec_t graphExec{};
hipStream_t stream{};
ret = hipGraphLaunch(graphExec, stream);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Destroy executable graph and try to launch it") {
constexpr size_t Nbytes = 1024;
hipGraph_t graph;
hipGraphExec_t graphExec;
hipStream_t stream;
hipGraphNode_t memsetNode;
char *devData;
HIP_CHECK(hipMalloc(&devData, Nbytes));
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipStreamCreate(&stream));
hipMemsetParams memsetParams{};
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(devData);
memsetParams.value = 0;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = Nbytes;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
&memsetParams));
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, stream));
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipGraphExecDestroy(graphExec));
// Launch again after destroy graph exec object.
ret = hipGraphLaunch(graphExec, stream);
REQUIRE(hipErrorInvalidValue == ret);
HIP_CHECK(hipFree(devData));
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(stream));
}
/* In this case in CUDA setup this api call is giving - unknown error (999)
So enabling this test for both AMD and CUDA by checking with hipSuccess */
SECTION("Destroy stream and try to launch respective executable graph") {
constexpr size_t Nbytes = 1024;
hipGraph_t graph;
hipGraphExec_t graphExec;
hipStream_t stream;
hipGraphNode_t memsetNode;
char *devData;
HIP_CHECK(hipMalloc(&devData, Nbytes));
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipStreamCreate(&stream));
hipMemsetParams memsetParams{};
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(devData);
memsetParams.value = 0;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = Nbytes;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
&memsetParams));
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, stream));
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipStreamDestroy(stream));
// Launch again after destroy stream
ret = hipGraphLaunch(graphExec, stream);
REQUIRE(hipSuccess != ret);
HIP_CHECK(hipFree(devData));
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
}
SECTION("Destroy graph and try to launch respective executable graph") {
constexpr size_t Nbytes = 1024;
hipGraph_t graph;
hipGraphExec_t graphExec;
hipStream_t stream;
hipGraphNode_t memsetNode;
char *devData;
HIP_CHECK(hipMalloc(&devData, Nbytes));
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipStreamCreate(&stream));
hipMemsetParams memsetParams{};
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(devData);
memsetParams.value = 0;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = Nbytes;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
&memsetParams));
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, stream));
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipGraphDestroy(graph));
// Launch again after destroy graph
ret = hipGraphLaunch(graphExec, stream);
REQUIRE(hipSuccess == ret);
HIP_CHECK(hipFree(devData));
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipStreamDestroy(stream));
}
}
TEST_CASE("Unit_hipGraphLaunch_Functional_hipStreamPerThread") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(char);
constexpr size_t val = 0;
constexpr size_t updateVal = 2;
char *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr};
char *A_h{nullptr}, *B_h{nullptr};
HipTest::initArrays<char>(&A_d, &B_d, &C_d,
&A_h, &B_h, nullptr, N, false);
hipGraph_t graph;
hipGraphExec_t graphExec;
hipGraphNode_t memsetNode;
HIP_CHECK(hipGraphCreate(&graph, 0));
hipMemsetParams memsetParams{};
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(C_d);
memsetParams.value = val;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = Nbytes;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
&memsetParams));
std::vector<hipGraphNode_t> dependencies;
dependencies.push_back(memsetNode);
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(A_d);
memsetParams.value = updateVal;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = Nbytes;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, dependencies.data(),
dependencies.size(), &memsetParams));
HIP_CHECK(hipGraphMemsetNodeSetParams(memsetNode, &memsetParams));
dependencies.push_back(memsetNode);
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, hipStreamPerThread));
HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
// Validating the result
for (size_t i = 0; i < Nbytes; i++) {
if (A_h[i] != updateVal) {
WARN("Validation failed at- " << i << " A_h[i] " << A_h[i]);
REQUIRE(false);
}
}
HipTest::freeArrays<char>(A_d, B_d, C_d,
A_h, B_h, nullptr, false);
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
}
static void hipGraphLaunch_test() {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(char);
constexpr size_t val = 0;
constexpr size_t updateVal = 1;
char *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr};
char *A_h{nullptr}, *B_h{nullptr};
HipTest::initArrays<char>(&A_d, &B_d, &C_d,
&A_h, &B_h, nullptr, N, false);
hipGraph_t graph;
hipGraphExec_t graphExec;
hipStream_t streamForGraph;
hipGraphNode_t memsetNode;
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipStreamCreate(&streamForGraph));
hipMemsetParams memsetParams{};
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(C_d);
memsetParams.value = val;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = Nbytes;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
&memsetParams));
std::vector<hipGraphNode_t> dependencies;
dependencies.push_back(memsetNode);
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(A_d);
memsetParams.value = updateVal;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = Nbytes;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, dependencies.data(),
dependencies.size(), &memsetParams));
HIP_CHECK(hipGraphMemsetNodeSetParams(memsetNode, &memsetParams));
dependencies.push_back(memsetNode);
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
// Validating the result
for (size_t i = 0; i < Nbytes; i++) {
if (A_h[i] != updateVal) {
WARN("Validation failed at- " << i << " A_h[i] " << A_h[i]);
REQUIRE(false);
}
}
HipTest::freeArrays<char>(A_d, B_d, C_d,
A_h, B_h, nullptr, false);
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(streamForGraph));
}
TEST_CASE("Unit_hipGraphLaunch_Functional_multidevice_test") {
int numDevices = 0;
HIP_CHECK(hipGetDeviceCount(&numDevices));
if (numDevices > 0) {
for (int i = 0; i < numDevices; i++) {
HIP_CHECK(hipSetDevice(i));
hipGraphLaunch_test();
}
} else {
SUCCEED("Skipped the testcase as there is no device to test.");
}
}
tests/catch/unit/graph/hipGraphMemcpyNodeSetParams1D.cc
+185
ファイルの表示
@@ -0,0 +1,185 @@
/*
Copyright (c) 2022 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.
*/
/**
Testcase Scenarios :
Functional-
1) Create a graph, add Memcpy node to graph, update the Memcpy node params with set and make sure they are taking effect.
Negative-
1) Pass pGraphNode as nullptr and check if api returns error.
2) Pass destination ptr is nullptr, api expected to return error code.
3) Pass source ptr is nullptr, api expected to return error code.
4) Pass count as zero, api expected to return error code.
5) Pass same pointer as source ptr and destination ptr, api expected to return error code.
6) Pass overlap memory as source ptr and destination ptr where source ptr is ahead of destination ptr, api expected to return error code.
7) Pass overlap memory as source ptr and destination ptr where destination ptr is ahead of source ptr, api expected to return error code.
8) If count is more than allocated size for source and destination ptr, api should return error code.
9) If count is less than allocated size for source and destination ptr, api should return error code.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
/* Test verifies hipGraphMemcpyNodeSetParams1D API Negative scenarios.
*/
TEST_CASE("Unit_hipGraphMemcpyNodeSetParams1D_Negative") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(int);
int *A_d, *A_h;
hipGraphNode_t memcpyNode{};
hipError_t ret;
HIP_CHECK(hipMalloc(&A_d, Nbytes));
HIP_CHECK(hipMalloc(&A_h, Nbytes));
hipGraph_t graph;
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
SECTION("Pass pGraphNode as nullptr") {
ret = hipGraphMemcpyNodeSetParams1D(nullptr, A_d, A_h, Nbytes,
hipMemcpyHostToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass destination ptr is nullptr") {
ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, nullptr, A_h, Nbytes,
hipMemcpyHostToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass source ptr is nullptr") {
ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d, nullptr, Nbytes,
hipMemcpyHostToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass count as zero") {
ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d, A_h, 0,
hipMemcpyHostToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
#if HT_AMD
SECTION("Pass same pointer as source ptr and destination ptr") {
ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d, A_d, Nbytes,
hipMemcpyHostToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass overlap memory where destination ptr is ahead of source ptr") {
ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d, A_d-5, Nbytes,
hipMemcpyHostToDevice);
REQUIRE(hipSuccess == ret);
}
#endif
SECTION("Pass overlap memory where source ptr is ahead of destination ptr") {
ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d+5, A_d, Nbytes-5,
hipMemcpyHostToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Copy more than allocated memory") {
ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d, A_h, Nbytes+8,
hipMemcpyHostToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Copy less than allocated memory") {
ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d, A_h, Nbytes-8,
hipMemcpyHostToDevice);
REQUIRE(hipSuccess == ret);
}
SECTION("Change the kind from H2D to D2H") {
ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d, A_h, Nbytes,
hipMemcpyDeviceToHost);
REQUIRE(hipSuccess == ret);
}
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(A_h));
HIP_CHECK(hipGraphDestroy(graph));
}
/* Test verifies hipGraphMemcpyNodeSetParams1D API Functional scenarios.
*/
TEST_CASE("Unit_hipGraphMemcpyNodeSetParams1D_Functional") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(int);
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
size_t NElem{N};
int *hData = reinterpret_cast<int*>(malloc(Nbytes));
REQUIRE(hData != nullptr);
memset(hData, 0, Nbytes);
hipGraphNode_t memcpyH2D_A, memcpyH2D_B, memcpyD2H_C;
hipGraphNode_t kernel_vecAdd;
hipKernelNodeParams kernelNodeParams{};
hipGraph_t graph;
hipGraphExec_t graphExec;
hipStream_t streamForGraph;
HIP_CHECK(hipStreamCreate(&streamForGraph));
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, C_h, C_d,
Nbytes, hipMemcpyDeviceToHost));
HIP_CHECK(hipGraphMemcpyNodeSetParams1D(memcpyD2H_C, hData, C_d, Nbytes,
hipMemcpyDeviceToHost));
void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kernelNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
kernelNodeParams.gridDim = dim3(blocks);
kernelNodeParams.blockDim = dim3(threadsPerBlock);
kernelNodeParams.sharedMemBytes = 0;
kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs2);
kernelNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0,
&kernelNodeParams));
// Create dependencies
HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &kernel_vecAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1));
// Instantiate and launch the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
// Verify graph execution result
HipTest::checkVectorADD(A_h, B_h, hData, N);
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipStreamDestroy(streamForGraph));
HIP_CHECK(hipGraphDestroy(graph));
free(hData);
}
tests/catch/unit/graph/hipGraphMemcpyNodeSetParamsFromSymbol.cc
+1
ファイルの表示
@@ -228,6 +228,7 @@ void hipGraphMemcpyNodeSetParamsFromSymbol_GlobalMem(bool useConstDeviceVar) {
// Instantiate and launch the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, 0));
HIP_CHECK(hipStreamSynchronize(0));
// Validating the result
for (int i = 0; i < SIZE; i++) {
tests/catch/unit/memory/hipHostUnregister.cc
-4
ファイルの表示
@@ -65,10 +65,6 @@ TEST_CASE("Unit_hipHostUnregister_MemoryNotAccessableAfterUnregister") {
}
TEST_CASE("Unit_hipHostUnregister_NullPtr") {
#if HT_AMD
HipTest::HIP_SKIP_TEST("EXSWCPHIPT-90");
return;
#endif
HIP_CHECK_ERROR(hipHostUnregister(nullptr), hipErrorInvalidValue);
}
tests/catch/unit/memory/hipMallocArray.cc
-4
ファイルの表示
@@ -494,10 +494,6 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_MaxTexture_Default", "", uint, int4, ush
// Arrays with channels of different size are not allowed.
TEST_CASE("Unit_hipMallocArray_Negative_DifferentChannelSizes") {
#if HT_AMD
HipTest::HIP_SKIP_TEST("EXSWCPHIPT-129");
return;
#endif
const int bitsX = GENERATE(8, 16, 32);
const int bitsY = GENERATE(8, 16, 32);
const int bitsZ = GENERATE(8, 16, 32);
tests/catch/unit/memory/hipMemsetNegative.cc
-11
ファイルの表示
@@ -71,9 +71,6 @@ TEST_CASE("Unit_hipMemset_Negative_InvalidPtr") {
TEST_CASE("Unit_hipMemset_Negative_OutOfBoundsSize") {
#if HT_AMD
HipTest::HIP_SKIP_TEST("EXSWCPHIPT-20");
#endif
#if !HT_AMD
void* dst;
@@ -112,10 +109,6 @@ TEST_CASE("Unit_hipMemset2D_Negative_InvalidPtr") {
}
TEST_CASE("Unit_hipMemset2D_Negative_InvalidSizes") {
#if HT_AMD
HipTest::HIP_SKIP_TEST("EXSWCPHIPT-52");
#endif
void* dst;
size_t realPitch;
HIP_CHECK(hipMallocPitch(&dst, &realPitch, width, height));
@@ -181,10 +174,6 @@ TEST_CASE("Unit_hipMemset3D_Negative_ModifiedPtr") {
}
TEST_CASE("Unit_hipMemset3D_Negative_InvalidSizes") {
#if HT_AMD
HipTest::HIP_SKIP_TEST("EXSWCPHIPT-52");
#endif
hipPitchedPtr pitchedDevPtr;
HIP_CHECK(hipMalloc3D(&pitchedDevPtr, validExtent));
hipExtent invalidExtent{validExtent};
tests/catch/unit/multiThread/CMakeLists.txt
+10
ファイルの表示
@@ -0,0 +1,10 @@
# Common Tests - Test independent of all platforms
set(TEST_SRC
hipMultiThreadDevice.cc
hipMultiThreadStreams1.cc
hipMultiThreadStreams2.cc
)
hip_add_exe_to_target(NAME MultiThreadTest
TEST_SRC ${TEST_SRC}
TEST_TARGET_NAME build_tests)
tests/catch/unit/multiThread/hipMultiThreadDevice.cc
+112
ファイルの表示
@@ -0,0 +1,112 @@
/* HIT_START
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
* TEST_NAMED: %t hipMultiThreadDevice-serial --tests 0x1
* TEST_NAMED: %t hipMultiThreadDevice-pyramid --tests 0x4
* TEST_NAMED: %t hipMultiThreadDevice-nearzero --tests 0x10
* HIT_END
*/
#include "hip/hip_runtime_api.h"
#include <hip_test_common.hh>
#ifdef _WIN32
#define MAX_BURST_SIZE 40
#else
#define MAX_BURST_SIZE 100
#endif
// Create a lot of streams and then destroy 'em.
void createThenDestroyStreams(int iterations, int burstSize) {
hipStream_t* streams = new hipStream_t[burstSize];
for (int i = 0; i < iterations; i++) {
for (int j = 0; j < burstSize; j++) {
HIPCHECK(hipStreamCreate(&streams[j]));
}
for (int j = 0; j < burstSize; j++) {
HIPCHECK(hipStreamDestroy(streams[j]));
}
}
delete[] streams;
}
void waitStreams(int iterations) {
// Repeatedly sync and wait for all streams to complete.
// TO make this interesting, the test has other threads repeatedly adding and removing streams
// to the device.
for (int i = 0; i < iterations; i++) {
HIPCHECK(hipDeviceSynchronize());
}
}
// Create 3 streams, all creating and destroying streams on the same device.
// Some create many queue, some not many.
//
void multiThread_pyramid(bool serialize, int iters) {
std::thread t1(createThenDestroyStreams, iters * 1, MAX_BURST_SIZE);
if (serialize) {
t1.join();
}
std::thread t2(createThenDestroyStreams, iters * 10, 10);
if (serialize) {
t2.join();
}
std::thread t3(createThenDestroyStreams, iters * 100, 1);
if (serialize) {
t3.join();
}
if (!serialize) {
t1.join();
t2.join();
t3.join();
}
}
// Create 3 streams, all creating and destroying streams on the same device.
// Try to keep number of streams near zero, to cause problems.
void multiThread_nearzero(bool serialize, int iters) {
std::thread t1(createThenDestroyStreams, iters, 1);
if (serialize) {
t1.join();
}
std::thread t2(createThenDestroyStreams, iters, 1);
if (serialize) {
t2.join();
}
std::thread t3(waitStreams, iters * 50);
if (serialize) {
t3.join();
}
if (!serialize) {
t1.join();
t2.join();
t3.join();
}
}
TEST_CASE("Unit_hipMultiThreadDevice_Streams") {
// Serial version, just call once:
createThenDestroyStreams(10, 10);
}
TEST_CASE("Unit_hipMultiThreadDevice_SerialPyramid") {
multiThread_pyramid(true, 3);
}
TEST_CASE("Unit_hipMultiThreadDevice_ParallelPyramid") {
multiThread_pyramid(false, 3);
}
TEST_CASE("Unit_hipMultiThreadDevice_NearZero") {
multiThread_nearzero(false, 1000);
}
tests/catch/unit/multiThread/hipMultiThreadStreams1.cc
+145
ファイルの表示
@@ -0,0 +1,145 @@
/*
Copyright (c) 2015 - 2021 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 NVCC_OPTIONS -std=c++11
* TEST: %t
* HIT_END
*/
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
int p_iters = 10;
int N = 8000000;
unsigned blocksPerCU = 6;
unsigned threadsPerBlock = 256;
//---
// Test simple H2D copies and back.
// Designed to stress a small number of simple smoke tests
template <typename T = float, class P = HipTest::Unpinned, class C = HipTest::Memcpy>
void simpleVectorAdd(size_t numElements, int iters, hipStream_t stream) {
using HipTest::MemTraits;
size_t Nbytes = numElements * sizeof(T);
T *A_d, *B_d, *C_d;
T *A_h, *B_h, *C_h;
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, P::isPinned);
for (size_t i = 0; i < numElements; i++) {
A_h[i] = 1000.0f;
B_h[i] = 2000.0f;
C_h[i] = -1;
}
MemTraits<C>::Copy(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream);
MemTraits<C>::Copy(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream);
MemTraits<C>::Copy(C_d, C_h, Nbytes, hipMemcpyHostToDevice, stream);
HIPCHECK(hipDeviceSynchronize());
for (size_t i = 0; i < numElements; i++) {
A_h[i] = 1.0f;
B_h[i] = 2.0f;
C_h[i] = -1;
}
for (int i = 0; i < iters; i++) {
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements);
MemTraits<C>::Copy(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream);
MemTraits<C>::Copy(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream);
hipLaunchKernelGGL(HipTest::vectorADDReverse, dim3(blocks), dim3(threadsPerBlock), 0, 0,
static_cast<const T*>(A_d), static_cast<const T*>(B_d), C_d, numElements);
MemTraits<C>::Copy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream);
HIPCHECK(hipDeviceSynchronize());
HipTest::checkVectorADD(A_h, B_h, C_h, numElements);
}
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, P::isPinned);
HIPCHECK(hipDeviceSynchronize());
}
template <typename T, class C>
void test_multiThread_1(hipStream_t stream0, hipStream_t stream1, bool serialize) {
size_t numElements = N;
// Test 2 threads operating on same stream:
std::thread t1(simpleVectorAdd<T, HipTest::Pinned, C>, numElements, p_iters /*iters*/, stream0);
if (serialize) {
t1.join();
}
std::thread t2(simpleVectorAdd<T, HipTest::Pinned, C>, numElements, p_iters /*iters*/, stream1);
if (serialize) {
t2.join();
}
if (!serialize) {
t1.join();
t2.join();
}
HIPCHECK(hipDeviceSynchronize());
};
TEST_CASE("Unit_hipMultiThreadStreams1_AsyncSync") {
hipStream_t stream;
HIPCHECK(hipStreamCreate(&stream));
simpleVectorAdd<float, HipTest::Pinned, HipTest::MemcpyAsync>(N /*mb*/, 10 /*iters*/, stream);
simpleVectorAdd<float, HipTest::Pinned, HipTest::Memcpy>(N /*mb*/, 10 /*iters*/, stream);
HIPCHECK(hipStreamDestroy(stream));
}
TEST_CASE("Unit_hipMultiThreadStreams1_AsyncAsync") {
hipStream_t stream0, stream1;
HIPCHECK(hipStreamCreate(&stream0));
HIPCHECK(hipStreamCreate(&stream1));
// Easy tests to verify the test works - these don't allow overlap between the threads:
test_multiThread_1<float, HipTest::MemcpyAsync>(NULL, NULL, true);
test_multiThread_1<float, HipTest::MemcpyAsync>(stream0, stream1, true);
HIPCHECK(hipStreamDestroy(stream0));
HIPCHECK(hipStreamDestroy(stream1));
}
TEST_CASE("Unit_hipMultiThreadStreams1_AsyncSame") {
hipStream_t stream;
HIPCHECK(hipStreamCreate(&stream));
// test_multiThread_1<float, HipTest::MemcpyAsync> ("Multithread with NULL stream", NULL,
// NULL, false); test_multiThread_1<float, HipTest::MemcpyAsync> ("Multithread with two
// streams", stream0, stream1, false);
test_multiThread_1<float, HipTest::MemcpyAsync>(stream, stream, false);
HIPCHECK(hipStreamDestroy(stream));
}
tests/catch/unit/multiThread/hipMultiThreadStreams2.cc
+145
ファイルの表示
@@ -0,0 +1,145 @@
/*
Copyright (c) 2015 - 2021 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 NVCC_OPTIONS -std=c++11
* TEST: %t
* HIT_END
*/
//This file is a port from hiprocclrtests (hipMultiThreadStreams2)
#include <iostream>
#include <hip_test_common.hh>
#include <thread>
#define N 1000
template <typename T>
__global__ void Inc(T* Array) {
int tx = threadIdx.x + blockIdx.x * blockDim.x;
Array[tx] = Array[tx] + T(1);
}
void run1(size_t size, hipStream_t stream) {
float *Ah, *Bh, *Cd, *Dd, *Eh;
float *snap = (float *) malloc(size);
HIPCHECK(hipHostMalloc((void**)&Ah, size, hipHostMallocDefault));
HIPCHECK(hipHostMalloc((void**)&Bh, size, hipHostMallocDefault));
HIPCHECK(hipMalloc(&Cd, size));
HIPCHECK(hipMalloc(&Dd, size));
HIPCHECK(hipHostMalloc((void**)&Eh, size, hipHostMallocDefault));
for (int i = 0; i < N; i++) {
Ah[i] = 1.0f;
}
HIPCHECK(hipMemcpyAsync(Bh, Ah, size, hipMemcpyHostToHost, stream));
HIPCHECK(hipMemcpyAsync(Cd, Bh, size, hipMemcpyHostToDevice, stream));
hipLaunchKernelGGL(HIP_KERNEL_NAME(Inc), dim3(N / 500), dim3(500), 0, stream, Cd);
HIPCHECK(hipMemcpyAsync(Dd, Cd, size, hipMemcpyDeviceToDevice, stream));
HIPCHECK(hipMemcpyAsync(Eh, Dd, size, hipMemcpyDeviceToHost, stream));
HIPCHECK(hipDeviceSynchronize());
memcpy(snap, Eh, size);
for (int i = 0; i < N; i++) {
HIPASSERT(snap[i] == Ah[i] + 1.0f);
}
free(snap);
HIPCHECK(hipHostFree(Ah));
HIPCHECK(hipHostFree(Bh));
HIPCHECK(hipHostFree(Eh));
HIPCHECK(hipFree(Cd));
HIPCHECK(hipFree(Dd));
}
void run(size_t size, hipStream_t stream1, hipStream_t stream2) {
float *Ah, *Bh, *Cd, *Dd, *Eh;
float *Ahh, *Bhh, *Cdd, *Ddd, *Ehh;
float *snap, *snapp;
snap = (float *) malloc(size);
snapp = (float *) malloc(size);
HIPCHECK(hipHostMalloc((void**)&Ah, size, hipHostMallocDefault));
HIPCHECK(hipHostMalloc((void**)&Bh, size, hipHostMallocDefault));
HIPCHECK(hipMalloc(&Cd, size));
HIPCHECK(hipMalloc(&Dd, size));
HIPCHECK(hipHostMalloc((void**)&Eh, size, hipHostMallocDefault));
HIPCHECK(hipHostMalloc((void**)&Ahh, size, hipHostMallocDefault));
HIPCHECK(hipHostMalloc((void**)&Bhh, size, hipHostMallocDefault));
HIPCHECK(hipMalloc(&Cdd, size));
HIPCHECK(hipMalloc(&Ddd, size));
HIPCHECK(hipHostMalloc((void**)&Ehh, size, hipHostMallocDefault));
HIPCHECK(hipMemcpyAsync(Bh, Ah, size, hipMemcpyHostToHost, stream1));
HIPCHECK(hipMemcpyAsync(Bhh, Ahh, size, hipMemcpyHostToHost, stream2));
HIPCHECK(hipMemcpyAsync(Cd, Bh, size, hipMemcpyHostToDevice, stream1));
HIPCHECK(hipMemcpyAsync(Cdd, Bhh, size, hipMemcpyHostToDevice, stream2));
hipLaunchKernelGGL(HIP_KERNEL_NAME(Inc), dim3(N / 500), dim3(500), 0, stream1, Cd);
hipLaunchKernelGGL(HIP_KERNEL_NAME(Inc), dim3(N / 500), dim3(500), 0, stream2, Cdd);
HIPCHECK(hipMemcpyAsync(Dd, Cd, size, hipMemcpyDeviceToDevice, stream1));
HIPCHECK(hipMemcpyAsync(Ddd, Cdd, size, hipMemcpyDeviceToDevice, stream2));
HIPCHECK(hipMemcpyAsync(Eh, Dd, size, hipMemcpyDeviceToHost, stream1));
HIPCHECK(hipMemcpyAsync(Ehh, Ddd, size, hipMemcpyDeviceToHost, stream2));
HIPCHECK(hipDeviceSynchronize());
memcpy(snap, Eh, size);
memcpy(snapp, Ehh, size);
for (int i = 0; i < N; i++) {
HIPASSERT(snap[i] == Ah[i] + 1.0f);
HIPASSERT(snapp[i] == Ahh[i] + 1.0f);
}
free(snap);
free(snapp);
HIPCHECK(hipHostFree(Ah));
HIPCHECK(hipHostFree(Bh));
HIPCHECK(hipHostFree(Eh));
HIPCHECK(hipHostFree(Ahh));
HIPCHECK(hipHostFree(Bhh));
HIPCHECK(hipHostFree(Ehh));
HIPCHECK(hipFree(Cd));
HIPCHECK(hipFree(Dd));
HIPCHECK(hipFree(Cdd));
HIPCHECK(hipFree(Ddd));
}
TEST_CASE("Unit_hipMultiThreadStreams2") {
int iterations = 100;
hipStream_t stream[3];
for (int i = 0; i < 3; i++) {
HIPCHECK(hipStreamCreate(&stream[i]));
}
const size_t size = N * sizeof(float);
for (int i = 0; i < iterations; i++) {
std::thread t1(run1, size, stream[0]);
std::thread t2(run1, size, stream[0]);
std::thread t3(run, size, stream[1], stream[2]);
t1.join();
t2.join();
t3.join();
}
}
tests/catch/unit/rtc/customOptions.cc
+33
ファイルの表示
@@ -146,3 +146,36 @@ TEST_CASE("Unit_hiprtc_namehandling") {
hiprtcDestroyProgram(&prog);
REQUIRE(compileResult == HIPRTC_SUCCESS);
}
TEST_CASE("Unit_hiprtc_getloweredname") {
using namespace std;
hiprtcProgram prog;
hiprtcCreateProgram(&prog, // prog
template_kernel, // buffer
"template_kernel.cu", // name
0, nullptr, nullptr);
std::string name_expression = "my_sqrt<complex<double> >";
REQUIRE(HIPRTC_SUCCESS == hiprtcAddNameExpression(prog, name_expression.c_str()));
hiprtcResult compileResult{hiprtcCompileProgram(prog, 0, 0)};
size_t logSize;
HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize));
if (logSize) {
string log(logSize, '\0');
HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0]));
std::cout << log << '\n';
}
const char* mangled_instantiation_cstr;
// Verifies if hiprtcGetLoweredName successfully gets the lowered name for named expressions with space
REQUIRE(HIPRTC_SUCCESS == hiprtcGetLoweredName(prog, name_expression.c_str(), &mangled_instantiation_cstr));
std::string mangled_name_str = mangled_instantiation_cstr;
// Checks if the fetched lowered name is not empty
REQUIRE(mangled_name_str.size() > 0);
hiprtcDestroyProgram(&prog);
REQUIRE(compileResult == HIPRTC_SUCCESS);
}
tests/catch/unit/stream/hipStreamCreateWithPriority.cc
+21
ファイルの表示
@@ -481,6 +481,16 @@ bool validateStreamPrioritiesWithEvents() {
OP(high)
#undef OP
// destroy stream
#define OP(x) \
if (enable_priority_##x) { \
HIP_CHECK(hipStreamDestroy(stream_##x)); \
}
OP(low)
OP(normal)
OP(high)
#undef OP
// validate that stream priorities are working as expected
#define OP(x, y) \
if (enable_priority_##x && enable_priority_##y) { \
@@ -495,6 +505,17 @@ bool validateStreamPrioritiesWithEvents() {
OP(low, high)
#undef OP
// free host & device memory
#define OP(x) \
free(src_h_##x); \
free(dst_h_##x); \
hipFree(src_d_##x); \
hipFree(dst_d_##x);
OP(low)
OP(normal)
OP(high)
#undef OP
return true;
}
tests/catch/unit/texture/CMakeLists.txt
+6
ファイルの表示
@@ -37,6 +37,12 @@ set(TEST_SRC
hipGetChanDesc.cc
hipTexObjPitch.cc
hipTextureObj1DFetch.cc
hipBindTex2DPitch.cc
hipBindTexRef1DFetch.cc
hipTex1DFetchCheckModes.cc
hipTextureObj1DCheckModes.cc
hipTextureObj2DCheckModes.cc
hipTextureObj3DCheckModes.cc
)
hip_add_exe_to_target(NAME TextureTest
tests/catch/unit/texture/hipBindTex2DPitch.cc
+9 -5
ファイルの表示
@@ -28,15 +28,19 @@ texture<TYPE_t, 2, hipReadModeElementType> tex;
// texture object is a kernel argument
static __global__ void texture2dCopyKernel(TYPE_t* dst) {
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
if ( (x < SIZE_W) && (y < SIZE_H) ) {
dst[SIZE_W*y+x] = tex2D(tex, x, y);
}
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
if ( (x < SIZE_W) && (y < SIZE_H) ) {
dst[SIZE_W*y+x] = tex2D(tex, x, y);
}
#endif
}
TEST_CASE("Unit_hipBindTexture2D_Pitch") {
CHECK_IMAGE_SUPPORT
TYPE_t* B;
TYPE_t* A;
TYPE_t* devPtrB;
tests/catch/unit/texture/hipBindTexRef1DFetch.cc
+4 -1
ファイルの表示
@@ -24,14 +24,17 @@ THE SOFTWARE.
texture<float, 1, hipReadModeElementType> tex;
static __global__ void kernel(float *out) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
if (x < N) {
out[x] = tex1Dfetch(tex, x);
}
#endif
}
TEST_CASE("Unit_hipBindTexture_tex1DfetchVerification") {
CHECK_IMAGE_SUPPORT
float *texBuf;
float val[N], output[N];
size_t offset = 0;
tests/catch/unit/texture/hipCreateTextureObject_ArgValidation.cc
+1 -1
ファイルの表示
@@ -25,7 +25,7 @@ THE SOFTWARE.
* Validate argument list of texture object api.
*/
TEST_CASE("Unit_hipCreateTextureObject_ArgValidation") {
checkImageSupport();
CHECK_IMAGE_SUPPORT
float *texBuf;
hipError_t ret;
tests/catch/unit/texture/hipCreateTextureObject_Array.cc
+2 -2
ファイルの表示
@@ -23,7 +23,7 @@ THE SOFTWARE.
* Validates Array Resource texture object with negative/functional tests.
*/
TEST_CASE("Unit_hipCreateTextureObject_ArrayResource") {
checkImageSupport();
CHECK_IMAGE_SUPPORT
hipError_t ret;
hipResourceDesc resDesc;
@@ -49,7 +49,7 @@ TEST_CASE("Unit_hipCreateTextureObject_ArrayResource") {
* with negative/functional tests.
*/
TEST_CASE("Unit_hipCreateTextureObject_MmArrayResource") {
checkImageSupport();
CHECK_IMAGE_SUPPORT
hipError_t ret;
hipResourceDesc resDesc;
tests/catch/unit/texture/hipCreateTextureObject_Linear.cc
+1 -1
ファイルの表示
@@ -26,7 +26,7 @@ THE SOFTWARE.
* Validates Linear Resource texture object with negative/functional tests.
*/
TEST_CASE("Unit_hipCreateTextureObject_LinearResource") {
checkImageSupport();
CHECK_IMAGE_SUPPORT
float *texBuf;
hipError_t ret;
tests/catch/unit/texture/hipCreateTextureObject_Pitch2D.cc
+1 -1
ファイルの表示
@@ -29,7 +29,7 @@ THE SOFTWARE.
* Validates Pitch2D Resource texture object with negative and functional tests
*/
TEST_CASE("Unit_hipCreateTextureObject_Pitch2DResource") {
checkImageSupport();
CHECK_IMAGE_SUPPORT
hipError_t ret;
hipResourceDesc resDesc;
tests/catch/unit/texture/hipGetChanDesc.cc
+2 -10
ファイルの表示
@@ -24,19 +24,11 @@ THE SOFTWARE.
TEST_CASE("Unit_hipGetChannelDesc_CreateAndGet") {
CHECK_IMAGE_SUPPORT
hipChannelFormatDesc chan_test, chan_desc;
hipArray *hipArray;
#if HT_AMD
int imageSupport{};
HIP_CHECK(hipDeviceGetAttribute(&imageSupport,
hipDeviceAttributeImageSupport, 0));
if (!imageSupport) {
INFO("Texture is not supported on the device. Test is skipped");
return;
}
#endif
chan_desc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindSigned);
HIP_CHECK(hipMallocArray(&hipArray, &chan_desc, C, R, 0));
HIP_CHECK(hipGetChannelDesc(&chan_test, hipArray));
tests/catch/unit/texture/hipNormalizedFloatValueTex.cc
+2 -8
ファイルの表示
@@ -142,16 +142,10 @@ static void runTest_hipTextureFilterMode() {
}
TEST_CASE("Unit_hipNormalizedFloatValueTex_CheckModes") {
CHECK_IMAGE_SUPPORT
#if HT_AMD
int imageSupport{};
HIP_CHECK(hipDeviceGetAttribute(&imageSupport,
hipDeviceAttributeImageSupport, 0));
if (!imageSupport) {
INFO("Texture is not supported on the device. Test is skipped");
return;
}
hipDeviceProp_t props;
HIP_CHECK(hipSetDevice(0));
HIP_CHECK(hipGetDeviceProperties(&props, 0));
INFO("Device :: " << props.name);
INFO("Arch - AMD GPU :: " << props.gcnArch);
tests/catch/unit/texture/hipSimpleTexture2DLayered.cc
+4
ファイルの表示
@@ -27,13 +27,17 @@ texture<float, hipTextureType2DLayered> tex2DL;
__global__ void simpleKernelLayeredArray(T* outputData,
int width, int height, int layer) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
outputData[layer * width * height + y * width + x] = tex2DLayered(tex2DL,
x, y, layer);
#endif
}
TEST_CASE("Unit_hipSimpleTexture2DLayered_Check") {
CHECK_IMAGE_SUPPORT
constexpr int SIZE = 512;
constexpr int num_layers = 5;
constexpr unsigned int width = SIZE;
tests/catch/unit/texture/hipSimpleTexture3D.cc
+2 -9
ファイルの表示
@@ -110,15 +110,8 @@ static void runSimpleTexture3D_Check(int width, int height, int depth,
}
TEST_CASE("Unit_hipSimpleTexture3D_Check_DataTypes") {
#if HT_AMD
int imageSupport{};
HIP_CHECK(hipDeviceGetAttribute(&imageSupport,
hipDeviceAttributeImageSupport, 0));
if (!imageSupport) {
INFO("Texture is not supported on the device. Test is skipped");
return;
}
#endif
CHECK_IMAGE_SUPPORT
for ( int i = 1; i < 25; i++ ) {
runSimpleTexture3D_Check<float>(i, i, i, &texf);
runSimpleTexture3D_Check<int>(i+1, i, i, &texi);
tests/catch/unit/texture/hipTex1DFetchCheckModes.cc
+18 -13
ファイルの表示
@@ -23,9 +23,11 @@ THE SOFTWARE.
#define offset 3
static __global__ void tex1dKernel(float *val, hipTextureObject_t obj) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
int k = blockIdx.x * blockDim.x + threadIdx.x;
if (k < (N - offset))
val[k] = tex1Dfetch<float>(obj, k+offset);
#endif
}
@@ -84,7 +86,8 @@ static void runTest(hipTextureAddressMode addressMode,
for (int i = 0; i < (N - offset); i++) {
if (output[i] != val[i + offset]) {
INFO("Output not matching at index " << i);
INFO("Mismatch found at output[" << i << "]:" << output[i] <<
" val[" << i + offset << "]:" << val[i + offset]);
REQUIRE(false);
}
}
@@ -103,16 +106,18 @@ static void runTest(hipTextureAddressMode addressMode,
TEST_CASE("Unit_tex1Dfetch_CheckModes") {
SECTION("hipAddressModeClamp AND hipFilterModePoint") {
runTest(hipAddressModeClamp, hipFilterModePoint);
}
SECTION("hipAddressModeClamp AND hipFilterModeLinear") {
runTest(hipAddressModeClamp, hipFilterModeLinear);
}
SECTION("hipAddressModeWrap AND hipFilterModePoint") {
runTest(hipAddressModeWrap, hipFilterModePoint);
}
SECTION("hipAddressModeWrap AND hipFilterModeLinear") {
runTest(hipAddressModeWrap, hipFilterModeLinear);
}
CHECK_IMAGE_SUPPORT
SECTION("hipAddressModeClamp AND hipFilterModePoint") {
runTest(hipAddressModeClamp, hipFilterModePoint);
}
SECTION("hipAddressModeClamp AND hipFilterModeLinear") {
runTest(hipAddressModeClamp, hipFilterModeLinear);
}
SECTION("hipAddressModeWrap AND hipFilterModePoint") {
runTest(hipAddressModeWrap, hipFilterModePoint);
}
SECTION("hipAddressModeWrap AND hipFilterModeLinear") {
runTest(hipAddressModeWrap, hipFilterModeLinear);
}
}
tests/catch/unit/texture/hipTexObjPitch.cc
+2 -10
ファイルの表示
@@ -38,21 +38,13 @@ static __global__ void texture2dCopyKernel(hipTextureObject_t texObj,
TEMPLATE_TEST_CASE("Unit_hipTexObjPitch_texture2D", "", float, int,
unsigned char, int16_t, char, unsigned int) {
CHECK_IMAGE_SUPPORT
TestType* B;
TestType* A;
TestType* devPtrB;
TestType* devPtrA;
#if HT_AMD
int imageSupport{};
HIP_CHECK(hipDeviceGetAttribute(&imageSupport,
hipDeviceAttributeImageSupport, 0));
if (!imageSupport) {
INFO("Texture is not supported on the device. Test is skipped");
return;
}
#endif
B = new TestType[SIZE_H*SIZE_W];
A = new TestType[SIZE_H*SIZE_W];
for (size_t i=1; i <= (SIZE_H*SIZE_W); i++) {
tests/catch/unit/texture/hipTextureMipmapObj2D.cc
+2 -8
ファイルの表示
@@ -121,14 +121,8 @@ static void runMipMapTest(unsigned int width, unsigned int height, unsigned int
#endif
TEST_CASE("Unit_hipTextureMipmapObj2D_Check") {
#if HT_AMD
int imageSupport{};
HIP_CHECK(hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport, 0));
if (!imageSupport) {
INFO("Texture is not supported on the device. Test is skipped");
return;
}
#endif
CHECK_IMAGE_SUPPORT
#ifdef _WIN32
for (auto& hw : hw_vector) {
for (auto& mip : mip_vector) {
tests/catch/unit/texture/hipTextureObj1DCheckModes.cc
+148
ファイルの表示
@@ -0,0 +1,148 @@
/*
Copyright (c) 2022 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.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_texture_helper.hh>
template<bool normalizedCoords>
__global__ void tex1DKernel(float *outputData, hipTextureObject_t textureObject,
int width, float offsetX) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
outputData[x] = tex1D<float>(textureObject, normalizedCoords ? (x + offsetX) / width : x + offsetX);
#endif
}
template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode, bool normalizedCoords>
void runTest(const int width, const float offsetX) {
//printf("%s(addressMode=%d, filterMode=%d, normalizedCoords=%d, width=%d, offsetX=%f)\n", __FUNCTION__,
// addressMode, filterMode, normalizedCoords, width, offsetX);
unsigned int size = width * sizeof(float);
float *hData = (float*) malloc(size);
memset(hData, 0, size);
for (int j = 0; j < width; j++) {
hData[j] = j;
}
hipChannelFormatDesc channelDesc = hipCreateChannelDesc(
32, 0, 0, 0, hipChannelFormatKindFloat);
hipArray *hipArray;
hipMallocArray(&hipArray, &channelDesc, width);
HIP_CHECK(hipMemcpy2DToArray(hipArray, 0, 0, hData, width * sizeof(float), width * sizeof(float), 1, hipMemcpyHostToDevice));
hipResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = hipResourceTypeArray;
resDesc.res.array.array = hipArray;
// Specify texture object parameters
hipTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
texDesc.addressMode[0] = addressMode;
texDesc.filterMode = filterMode;
texDesc.readMode = hipReadModeElementType;
texDesc.normalizedCoords = normalizedCoords;
// Create texture object
hipTextureObject_t textureObject = 0;
HIP_CHECK(hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL));
float *dData = nullptr;
hipMalloc((void**) &dData, size);
dim3 dimBlock(16, 1, 1);
dim3 dimGrid((width + dimBlock.x - 1)/ dimBlock.x, 1, 1);
hipLaunchKernelGGL(tex1DKernel<normalizedCoords>, dimGrid, dimBlock, 0, 0, dData,
textureObject, width, offsetX);
hipDeviceSynchronize();
float *hOutputData = (float*) malloc(size);
memset(hOutputData, 0, size);
hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost);
bool result = true;
for (int j = 0; j < width; j++) {
float expectedValue = getExpectedValue<addressMode, filterMode>(width, offsetX + j, hData);
if (!hipTextureSamplingVerify<float, filterMode>(hOutputData[j], expectedValue)) {
INFO("Mismatch at " << offsetX + j << ":" << hOutputData[j] <<
" expected:" << expectedValue);
result = false;
break;
}
}
hipDestroyTextureObject(textureObject);
hipFree(dData);
hipFreeArray(hipArray);
free(hData);
free(hOutputData);
REQUIRE(result);
}
TEST_CASE("Unit_hipTextureObj1DCheckModes") {
CHECK_IMAGE_SUPPORT
#ifdef _WIN32
INFO("Unit_hipTextureObj1DCheckModes skipped on Windows");
return;
#endif
SECTION("hipAddressModeClamp, hipFilterModePoint, regularCoords") {
runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, -3);
runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 4);
}
SECTION("hipAddressModeBorder, hipFilterModePoint, regularCoords") {
runTest<hipAddressModeBorder, hipFilterModePoint, false>(256, -8.5);
runTest<hipAddressModeBorder, hipFilterModePoint, false>(256, 12.5);
}
SECTION("hipAddressModeClamp, hipFilterModeLinear, regularCoords") {
runTest<hipAddressModeClamp, hipFilterModeLinear, false>(256, -3);
runTest<hipAddressModeClamp, hipFilterModeLinear, false>(256, 4);
}
SECTION("hipAddressModeBorder, hipFilterModeLinear, regularCoords") {
runTest<hipAddressModeBorder, hipFilterModeLinear, false>(256, -8.5);
runTest<hipAddressModeBorder, hipFilterModeLinear, false>(256, 12.5);
}
SECTION("hipAddressModeClamp, hipFilterModePoint, normalizedCoords") {
runTest<hipAddressModeClamp, hipFilterModePoint, true>(256, -3);
runTest<hipAddressModeClamp, hipFilterModePoint, true>(256, 4);
}
SECTION("hipAddressModeBorder, hipFilterModePoint, normalizedCoords") {
runTest<hipAddressModeBorder, hipFilterModePoint, true>(256, -8.5);
runTest<hipAddressModeBorder, hipFilterModePoint, true>(256, 12.5);
}
SECTION("hipAddressModeClamp, hipFilterModeLinear, normalizedCoords") {
runTest<hipAddressModeClamp, hipFilterModeLinear, true>(256, -3);
runTest<hipAddressModeClamp, hipFilterModeLinear, true>(256, 4);
}
SECTION("hipAddressModeBorder, hipFilterModeLinear, normalizedCoords") {
runTest<hipAddressModeBorder, hipFilterModeLinear, true>(256, -8.5);
runTest<hipAddressModeBorder, hipFilterModeLinear, true>(256, 12.5);
}
}
tests/catch/unit/texture/hipTextureObj1DFetch.cc
+8 -4
ファイルの表示
@@ -22,14 +22,18 @@ THE SOFTWARE.
#define N 512
static __global__ void tex1dKernel(float *val, hipTextureObject_t obj) {
int k = blockIdx.x * blockDim.x + threadIdx.x;
if (k < N) {
val[k] = tex1Dfetch<float>(obj, k);
}
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
int k = blockIdx.x * blockDim.x + threadIdx.x;
if (k < N) {
val[k] = tex1Dfetch<float>(obj, k);
}
#endif
}
TEST_CASE("Unit_hipCreateTextureObject_tex1DfetchVerification") {
CHECK_IMAGE_SUPPORT
// Allocating the required buffer on gpu device
float *texBuf, *texBufOut;
float val[N], output[N];
tests/catch/unit/texture/hipTextureObj2D.cc
+5 -12
ファイルの表示
@@ -22,22 +22,15 @@ THE SOFTWARE.
__global__ void tex2DKernel(float* outputData,
hipTextureObject_t textureObject, int width) {
#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);
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
}
TEST_CASE("Unit_hipTextureObj2D_Check") {
#if HT_AMD
int imageSupport{};
HIP_CHECK(hipDeviceGetAttribute(&imageSupport,
hipDeviceAttributeImageSupport, 0));
if (!imageSupport) {
INFO("Texture is not supported on the device. Test is skipped");
return;
}
#endif
CHECK_IMAGE_SUPPORT
constexpr int SIZE = 256;
constexpr unsigned int width = SIZE;
constexpr unsigned int height = SIZE;
tests/catch/unit/texture/hipTextureObj2DCheckModes.cc
+160
ファイルの表示
@@ -0,0 +1,160 @@
/*
Copyright (c) 2022 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.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_texture_helper.hh>
template<bool normalizedCoords>
__global__ void tex2DKernel(float *outputData, hipTextureObject_t textureObject,
int width, int height, float offsetX,
float offsetY) {
#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,
normalizedCoords ? (x + offsetX) / width : x + offsetX,
normalizedCoords ? (y + offsetY) / height : y + offsetY);
#endif
}
template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode, bool normalizedCoords>
void runTest(const int width, const int height, const float offsetX, const float offsetY) {
//printf("%s(addressMode=%d, filterMode=%d, normalizedCoords=%d, width=%d, height=%d, offsetX=%f, offsetY=%f)\n",
// __FUNCTION__, addressMode, filterMode, normalizedCoords, width, height, offsetX, offsetY);
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++) {
int index = i * width + j;
hData[index] = index;
}
}
hipChannelFormatDesc channelDesc = hipCreateChannelDesc(
32, 0, 0, 0, hipChannelFormatKindFloat);
hipArray *hipArray;
hipMallocArray(&hipArray, &channelDesc, width, height);
HIP_CHECK(hipMemcpy2DToArray(hipArray, 0, 0, hData, width * sizeof(float), width * sizeof(float), height, hipMemcpyHostToDevice));
hipResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = hipResourceTypeArray;
resDesc.res.array.array = hipArray;
// Specify texture object parameters
hipTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
texDesc.addressMode[0] = addressMode;
texDesc.addressMode[1] = addressMode;
texDesc.filterMode = filterMode;
texDesc.readMode = hipReadModeElementType;
texDesc.normalizedCoords = normalizedCoords;
// Create texture object
hipTextureObject_t textureObject = 0;
HIP_CHECK(hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL));
float *dData = nullptr;
hipMalloc((void**) &dData, size);
dim3 dimBlock(16, 16, 1);
dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y, 1);
hipLaunchKernelGGL(tex2DKernel<normalizedCoords>, dimGrid, dimBlock, 0, 0, dData,
textureObject, width, height, offsetX, offsetY);
hipDeviceSynchronize();
float *hOutputData = (float*) malloc(size);
memset(hOutputData, 0, size);
hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost);
bool result = true;
for (int i = 0; i < height; i++) {
for (int j = 0; j < width; j++) {
int index = i * width + j;
float expectedValue = getExpectedValue<addressMode, filterMode>(width, height,
offsetX + j, offsetY + i, hData);
if (!hipTextureSamplingVerify<float, filterMode>(hOutputData[index], expectedValue)) {
INFO("Mismatch at (" << offsetX + j << ", " << offsetY + i << "):" <<
hOutputData[index] << " expected:" << expectedValue);
result = false;
goto line1;
}
}
}
line1:
hipDestroyTextureObject(textureObject);
hipFree(dData);
hipFreeArray(hipArray);
free(hData);
free(hOutputData);
REQUIRE(result);
}
TEST_CASE("Unit_hipTextureObj2DCheckModes") {
CHECK_IMAGE_SUPPORT
#ifdef _WIN32
INFO("Unit_hipTextureObj2DCheckModes skipped on Windows");
return;
#endif
SECTION("hipAddressModeClamp, hipFilterModePoint, regularCoords") {
runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 256, -3.9, 6.1);
runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 256, 4.4, -7.0);
}
SECTION("hipAddressModeBorder, hipFilterModePoint, regularCoords") {
runTest<hipAddressModeBorder, hipFilterModePoint, false>(256, 256, -8.5, 2.9);
runTest<hipAddressModeBorder, hipFilterModePoint, false>(256, 256, 12.5, 6.7);
}
SECTION("hipAddressModeClamp, hipFilterModePoint, regularCoords") {
runTest<hipAddressModeClamp, hipFilterModeLinear, false>(256, 256, -0.4, -0.4);
runTest<hipAddressModeClamp, hipFilterModeLinear, false>(256, 256, 4, 14.6);
}
SECTION("hipAddressModeBorder, hipFilterModeLinear, regularCoords") {
runTest<hipAddressModeBorder, hipFilterModeLinear, false>(256, 256, -0.4, 0.4);
runTest<hipAddressModeBorder, hipFilterModeLinear, false>(256, 256, 12.5, 23.7);
}
SECTION("hipAddressModeClamp, hipFilterModePoint, normalizedCoords") {
runTest<hipAddressModeClamp, hipFilterModePoint, true>(256, 256, -3, 8.9);
runTest<hipAddressModeClamp, hipFilterModePoint, true>(256, 256, 4, -0.1);
}
SECTION("hipAddressModeBorder, hipFilterModePoint, normalizedCoords") {
runTest<hipAddressModeBorder, hipFilterModePoint, true>(256, 256, -8.5, 15.9);
runTest<hipAddressModeBorder, hipFilterModePoint, true>(256, 256, 12.5, -17.9);
}
SECTION("hipAddressModeClamp, hipFilterModeLinear, normalizedCoords") {
runTest<hipAddressModeClamp, hipFilterModeLinear, true>(256, 256, -3, 5.8);
runTest<hipAddressModeClamp, hipFilterModeLinear, true>(256, 256, 4, 9.1);
}
SECTION("hipAddressModeBorder, hipFilterModeLinear, normalizedCoords") {
runTest<hipAddressModeBorder, hipFilterModeLinear, true>(256, 256, -8.5, 6.6);
runTest<hipAddressModeBorder, hipFilterModeLinear, true>(256, 256, 12.5, 0.01);
}
}
tests/catch/unit/texture/hipTextureObj3DCheckModes.cc
+218
ファイルの表示
@@ -0,0 +1,218 @@
/*
Copyright (c) 2022 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.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_texture_helper.hh>
bool isGfx90a = false;
template<bool normalizedCoords>
__global__ void tex3DKernel(float *outputData, hipTextureObject_t textureObject,
int width, int height, int depth, float offsetX,
float offsetY, float offsetZ) {
#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;
int z = blockIdx.z * blockDim.z + threadIdx.z;
outputData[z * width * height + y * width + x] = tex3D<float>(textureObject,
normalizedCoords ? (x + offsetX) / width : x + offsetX,
normalizedCoords ? (y + offsetY) / height : y + offsetY,
normalizedCoords ? (z + offsetZ) / depth : z + offsetZ);
#endif
}
template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode, bool normalizedCoords>
void runTest(const int width, const int height, const int depth, const float offsetX, const float offsetY,
const float offsetZ) {
//printf("%s(addressMode=%d, filterMode=%d, normalizedCoords=%d, width=%d, height=%d, depth=%d, offsetX=%f, offsetY=%f, offsetZ=%f)\n",
// __FUNCTION__, addressMode, filterMode, normalizedCoords, width, height,
// depth, offsetX, offsetY, offsetZ);
bool result = true;
unsigned int size = width * height * depth * sizeof(float);
float *hData = (float*) malloc(size);
memset(hData, 0, size);
for (int i = 0; i < depth; i++) {
for (int j = 0; j < height; j++) {
for (int k = 0; k < width; k++) {
int index = i * width * height + j * width + k;
hData[index] = index;
}
}
}
// Allocate array and copy image data
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<float>();
hipArray *arr;
HIP_CHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, depth), hipArrayDefault));
hipMemcpy3DParms myparms;
memset(&myparms, 0, sizeof(myparms));
myparms.srcPos = make_hipPos(0,0,0);
myparms.dstPos = make_hipPos(0,0,0);
myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(float), width, height);
myparms.dstArray = arr;
myparms.extent = make_hipExtent(width, height, depth);
myparms.kind = hipMemcpyHostToDevice;
HIP_CHECK(hipMemcpy3D(&myparms));
hipResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = hipResourceTypeArray;
resDesc.res.array.array = arr;
// Specify texture object parameters
hipTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
texDesc.addressMode[0] = addressMode;
texDesc.addressMode[1] = addressMode;
texDesc.addressMode[2] = addressMode;
texDesc.filterMode = filterMode;
texDesc.readMode = hipReadModeElementType;
texDesc.normalizedCoords = normalizedCoords;
// Create texture object
hipTextureObject_t textureObject = 0;
hipError_t res = hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL);
if (res != hipSuccess) {
hipFreeArray(arr);
free(hData);
if (res == hipErrorNotSupported && isGfx90a) {
printf("gfx90a doesn't support 3D linear filter! Skipped!\n");
} else {
result = false;
}
REQUIRE(result);
return;
}
float *dData = nullptr;
hipMalloc((void**) &dData, size);
hipMemset(dData, 0, size);
dim3 dimBlock(8, 8, 8); // 512 threads
dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y,
(depth + dimBlock.z - 1) / dimBlock.z);
hipLaunchKernelGGL(tex3DKernel<normalizedCoords>, dimGrid, dimBlock, 0, 0, dData,
textureObject, width, height, depth, offsetX, offsetY, offsetZ);
hipDeviceSynchronize();
float *hOutputData = (float*) malloc(size);
memset(hOutputData, 0, size);
hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost);
for (int i = 0; i < depth; i++) {
for (int j = 0; j < height; j++) {
for (int k = 0; k < width; k++) {
int index = i * width * height + j * width + k;
float expectedValue = getExpectedValue<addressMode, filterMode>(
width, height, depth, offsetX + k, offsetY + j, offsetZ + i, hData);
if (!hipTextureSamplingVerify<float, filterMode>(hOutputData[index], expectedValue)) {
INFO("Mismatch at (" << offsetX + k << ", " << offsetY + j << ", " << offsetZ + i << "):" <<
hOutputData[index] << " expected:" << expectedValue);
result = false;
goto line1;
}
}
}
}
line1:
hipDestroyTextureObject(textureObject);
free(hOutputData);
hipFree(dData);
hipFreeArray(arr);
free(hData);
REQUIRE(result);
}
TEST_CASE("Unit_hipTextureObj3DCheckModes") {
CHECK_IMAGE_SUPPORT
#ifdef _WIN32
INFO("Unit_hipTextureObj3DCheckModes skipped on Windows");
return;
#endif
int device = 0;
hipDeviceProp_t props;
HIPCHECK(hipGetDeviceProperties(&props, device));
if (!strncmp(props.gcnArchName, "gfx90a", strlen("gfx90a"))) {
isGfx90a = true;
}
SECTION("hipAddressModeClamp, hipFilterModePoint, regularCoords") {
runTest<hipAddressModeClamp, hipFilterModePoint, false>
(256, 256, 256, -3.9, 6.1, 9.5);
runTest<hipAddressModeClamp, hipFilterModePoint, false>
(256, 256, 256, 4.4, -7.0, 5.3);
}
SECTION("hipAddressModeBorder, hipFilterModePoint, regularCoords") {
runTest<hipAddressModeBorder, hipFilterModePoint, false>
(256, 256, 256, -8.5, 2.9, 5.8);
runTest<hipAddressModeBorder, hipFilterModePoint, false>
(256, 256, 256, 12.5, 6.7, 11.4);
}
SECTION("hipAddressModeClamp, hipFilterModeLinear, regularCoords") {
runTest<hipAddressModeClamp, hipFilterModeLinear, false>
(256, 256, 256, -0.4, -0.4, -0.4);
runTest<hipAddressModeClamp, hipFilterModeLinear, false>
(256, 256, 256, 4, 14.6, -0.3);
}
SECTION("hipAddressModeBorder, hipFilterModeLinear, regularCoords") {
runTest<hipAddressModeBorder, hipFilterModeLinear, false>
(256, 256, 256, 6.9, 7.4, 0.4);
runTest<hipAddressModeBorder, hipFilterModeLinear, false>
(256, 256, 256, 12.5, 23.7, 0.34);
}
SECTION("hipAddressModeClamp, hipFilterModePoint, normalizedCoords") {
runTest<hipAddressModeClamp, hipFilterModePoint, true>
(256, 256, 256, -3, 8.9, -4);
runTest<hipAddressModeClamp, hipFilterModePoint, true>
(256, 256, 256, 4, -0.1, 8.2);
}
SECTION("hipAddressModeBorder, hipFilterModePoint, normalizedCoords") {
runTest<hipAddressModeBorder, hipFilterModePoint, true>
(256, 256, 256, -8.5, 15.9, 0.1);
runTest<hipAddressModeBorder, hipFilterModePoint, true>
(256, 256, 256, 12.5, -17.9, -0.35);
}
SECTION("hipAddressModeClamp, hipFilterModeLinear, normalizedCoords") {
runTest<hipAddressModeClamp, hipFilterModeLinear, true>
(256, 256, 256, -3, 5.8, 0.89);
runTest<hipAddressModeClamp, hipFilterModeLinear, true>
(256, 256, 256, 4, 9.1, 2.08);
}
SECTION("hipAddressModeBorder, hipFilterModeLinear, normalizedCoords") {
runTest<hipAddressModeBorder, hipFilterModeLinear, true>
(256, 256, 256, -8.5, 6.6, 3.67);
runTest<hipAddressModeBorder, hipFilterModeLinear, true>
(256, 256, 256, 12.5, 0.01, -9.9);
}
}
tests/catch/unit/texture/hipTextureObjFetchVector.cc
+23 -26
ファイルの表示
@@ -138,14 +138,12 @@ static inline bool isEqual(const T &val0, const T &val1) {
}
template<typename T>
bool runTest(const char *description) {
bool runTest() {
const int N = 1024;
bool testResult = true;
// Allocating the required buffer on gpu device
T *texBuf, *texBufOut;
T val[N], output[N];
printf("%s<%s>(): size: %zu, %zu\n", __FUNCTION__, description,
sizeof(T), sizeof(decltype(T::x)));
memset(output, 0, sizeof(output));
std::srand(std::time(nullptr)); // use current time as seed for random generator
@@ -199,46 +197,45 @@ bool runTest(const char *description) {
HIP_CHECK(hipFree(texBuf));
HIP_CHECK(hipFree(texBufOut));
printf(": %s\n", testResult ? "succeeded" : "failed");
REQUIRE(testResult == true);
return testResult;
}
TEST_CASE("Unit_hipTextureFetch_vector") {
checkImageSupport();
CHECK_IMAGE_SUPPORT
// test for char
runTest<char1>("char1");
runTest<char2>("char2");
runTest<char4>("char4");
runTest<char1>();
runTest<char2>();
runTest<char4>();
// test for uchar
runTest<uchar1>("uchar1");
runTest<uchar2>("uchar2");
runTest<uchar4>("uchar4");
runTest<uchar1>();
runTest<uchar2>();
runTest<uchar4>();
// test for short
runTest<short1>("short1");
runTest<short2>("short2");
runTest<short4>("short4");
runTest<short1>();
runTest<short2>();
runTest<short4>();
// test for ushort
runTest<ushort1>("ushort1");
runTest<ushort2>("ushort2");
runTest<ushort4>("ushort4");
runTest<ushort1>();
runTest<ushort2>();
runTest<ushort4>();
// test for int
runTest<int1>("int1");
runTest<int2>("int2");
runTest<int4>("int4");
runTest<int1>();
runTest<int2>();
runTest<int4>();
// test for unsigned int
runTest<uint1>("uint1");
runTest<uint2>("uint2");
runTest<uint4>("uint4");
runTest<uint1>();
runTest<uint2>();
runTest<uint4>();
// test for float
runTest<float1>("float1");
runTest<float2>("float2");
runTest<float4>("float4");
runTest<float1>();
runTest<float2>();
runTest<float4>();
}
tests/catch/unit/texture/hipTextureRef2D.cc
+4
ファイルの表示
@@ -22,12 +22,16 @@ THE SOFTWARE.
texture<float, 2, hipReadModeElementType> tex;
__global__ void tex2DKernel(float* outputData, int width) {
#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(tex, x, y);
#endif
}
TEST_CASE("Unit_hipTextureRef2D_Check") {
CHECK_IMAGE_SUPPORT
constexpr int SIZE = 256;
constexpr unsigned int width = SIZE;
constexpr unsigned int height = SIZE;
tests/src/texture/hipTextureObj3DCheckModes.cpp
+23 -7
ファイルの表示
@@ -11,6 +11,7 @@
#include "test_common.h"
#include "hipTextureHelper.hpp"
bool isGfx90a = false;
template<bool normalizedCoords>
__global__ void tex3DKernel(float *outputData, hipTextureObject_t textureObject,
@@ -20,7 +21,7 @@ __global__ void tex3DKernel(float *outputData, hipTextureObject_t textureObject,
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
outputData[z * width * depth + y * width + x] = tex3D<float>(textureObject,
outputData[z * width * height + y * width + x] = tex3D<float>(textureObject,
normalizedCoords ? (x + offsetX) / width : x + offsetX,
normalizedCoords ? (y + offsetY) / height : y + offsetY,
normalizedCoords ? (z + offsetZ) / depth : z + offsetZ);
@@ -40,7 +41,7 @@ bool runTest(const int width, const int height, const int depth, const float off
for (int i = 0; i < depth; i++) {
for (int j = 0; j < height; j++) {
for (int k = 0; k < width; k++) {
int index = i * width * depth + j * width + k;
int index = i * width * height + j * width + k;
hData[index] = index;
}
}
@@ -78,8 +79,17 @@ bool runTest(const int width, const int height, const int depth, const float off
// Create texture object
hipTextureObject_t textureObject = 0;
hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL);
hipError_t res = hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL);
if (res != hipSuccess) {
hipFreeArray(arr);
free(hData);
if (res == hipErrorNotSupported && isGfx90a) {
printf("gfx90a doesn't support 3D linear filter! Skipped!\n");
} else {
testResult = false;
}
return testResult;
}
float *dData = NULL;
hipMalloc((void**) &dData, size);
hipMemset(dData, 0, size);
@@ -98,7 +108,7 @@ bool runTest(const int width, const int height, const int depth, const float off
for (int i = 0; i < depth; i++) {
for (int j = 0; j < height; j++) {
for (int k = 0; k < width; k++) {
int index = i * width * depth + j * width + k;
int index = i * width * height + j * width + k;
float expectedValue = getExpectedValue<addressMode, filterMode>(
width, height, depth, offsetX + k, offsetY + j, offsetZ + i, hData);
@@ -112,11 +122,10 @@ bool runTest(const int width, const int height, const int depth, const float off
}
line1:
hipDestroyTextureObject(textureObject);
free(hOutputData);
hipFree(dData);
hipFreeArray(arr);
free(hData);
free(hOutputData);
printf("%s %s\n", __FUNCTION__, testResult ? "succeeded":"failed");
return testResult;
}
@@ -125,6 +134,13 @@ int main(int argc, char **argv) {
bool testResult = true;
int device = 0;
hipDeviceProp_t props;
HIPCHECK(hipGetDeviceProperties(&props, device));
if (!strncmp(props.gcnArchName, "gfx90a", strlen("gfx90a"))) {
isGfx90a = true;
}
testResult = testResult && runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 256, 256, -3.9, 6.1, 9.5);
testResult = testResult && runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 256, 256, 4.4, -7.0, 5.3);