diff --git a/README.md b/README.md index 39165993cc..575fdd061a 100644 --- a/README.md +++ b/README.md @@ -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); diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index 868f0de44a..53d88f65aa 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -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. */ diff --git a/tests/catch/hipTestMain/config/config_amd_linux.json b/tests/catch/hipTestMain/config/config_amd_linux.json index 91c89ae7dc..60ef935548 100644 --- a/tests/catch/hipTestMain/config/config_amd_linux.json +++ b/tests/catch/hipTestMain/config/config_amd_linux.json @@ -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" ] - } diff --git a/tests/catch/hipTestMain/config/config_amd_windows.json b/tests/catch/hipTestMain/config/config_amd_windows.json index 547d92a4bf..c755d65649 100644 --- a/tests/catch/hipTestMain/config/config_amd_windows.json +++ b/tests/catch/hipTestMain/config/config_amd_windows.json @@ -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" ] } diff --git a/tests/catch/include/hip_test_checkers.hh b/tests/catch/include/hip_test_checkers.hh index fce61fe677..3e152f099b 100644 --- a/tests/catch/include/hip_test_checkers.hh +++ b/tests/catch/include/hip_test_checkers.hh @@ -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 +size_t checkVectorSUB(T* A_h, T* B_h, T* result_H, size_t N, bool expectMatch = true, + bool reportMismatch = true) { + return checkVectors( + A_h, B_h, result_H, N, [](T a, T b) { return a - b; }, expectMatch, reportMismatch); +} + template void checkTest(T* expected_H, T* result_H, size_t N, bool expectMatch = true) { checkVectors( diff --git a/tests/catch/include/hip_test_common.hh b/tests/catch/include/hip_test_common.hh index fe272f7ace..e97c99bd8f 100644 --- a/tests/catch/include/hip_test_common.hh +++ b/tests/catch/include/hip_test_common.hh @@ -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(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 +struct MemTraits; + + +template <> +struct MemTraits { + 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 { + 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; } diff --git a/tests/catch/include/hip_test_kernels.hh b/tests/catch/include/hip_test_kernels.hh index d59fd41b5a..68f452ba53 100644 --- a/tests/catch/include/hip_test_kernels.hh +++ b/tests/catch/include/hip_test_kernels.hh @@ -34,6 +34,14 @@ template __global__ void vectorADD(const T* A_d, const T* B_d, T* C } } +template __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 __global__ void vectorADDReverse(const T* A_d, const T* B_d, T* C_d, size_t NELEM) { diff --git a/tests/catch/include/hip_texture_helper.hh b/tests/catch/include/hip_texture_helper.hh new file mode 100644 index 0000000000..80d57f8a38 --- /dev/null +++ b/tests/catch/include/hip_texture_helper.hh @@ -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 +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 +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 +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::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 +float getExpectedValue(const int width, float x, const float *data) { + float result = std::numeric_limits::lowest(); + switch (filterMode) { + case hipFilterModePoint: { + int i1 = static_cast(floor(x)); + hipTextureGetAddress < addressMode > (i1, width); + result = hipTextureGetValue < addressMode > (data, i1, width); + } + break; + case hipFilterModeLinear: { + x -= 0.5; + int i1 = static_cast(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 +float getExpectedValue(const int width, const int height, float x, float y, const float *data) { + float result = std::numeric_limits::lowest(); + switch (filterMode) { + case hipFilterModePoint: { + int i1 = static_cast(floor(x)); + int j1 = static_cast(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(floor(x)); + int j1 = static_cast(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 +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::lowest(); + switch (filterMode) { + case hipFilterModePoint: { + int i1 = static_cast(floor(x)); + int j1 = static_cast(floor(y)); + int k1 = static_cast(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(floor(x)); + int j1 = static_cast(floor(y)); + int k1 = static_cast(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; +} \ No newline at end of file diff --git a/tests/catch/unit/CMakeLists.txt b/tests/catch/unit/CMakeLists.txt index 28c6fcb6aa..9874ce29d2 100644 --- a/tests/catch/unit/CMakeLists.txt +++ b/tests/catch/unit/CMakeLists.txt @@ -30,3 +30,4 @@ add_subdirectory(printf) add_subdirectory(texture) add_subdirectory(streamperthread) add_subdirectory(kernel) +add_subdirectory(multiThread) diff --git a/tests/catch/unit/device/hipSetGetDevice.cc b/tests/catch/unit/device/hipSetGetDevice.cc index 7294863ecf..011779a489 100644 --- a/tests/catch/unit/device/hipSetGetDevice.cc +++ b/tests/catch/unit/device/hipSetGetDevice.cc @@ -23,21 +23,91 @@ */ #include +#include 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 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); + } } diff --git a/tests/catch/unit/event/Unit_hipEvent_Negative.cc b/tests/catch/unit/event/Unit_hipEvent_Negative.cc index 76b75d6c06..e315710bbd 100644 --- a/tests/catch/unit/event/Unit_hipEvent_Negative.cc +++ b/tests/catch/unit/event/Unit_hipEvent_Negative.cc @@ -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 diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index 33a6ce3cd5..b07f15891b 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -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 diff --git a/tests/catch/unit/graph/hipGraphExecKernelNodeSetParams.cc b/tests/catch/unit/graph/hipGraphExecKernelNodeSetParams.cc new file mode 100644 index 0000000000..e1141fbe2a --- /dev/null +++ b/tests/catch/unit/graph/hipGraphExecKernelNodeSetParams.cc @@ -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 +#include +#include + +/** + * 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 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(&NElem)}; + kNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kNodeParams.gridDim = dim3(blocks); + kNodeParams.blockDim = dim3(threadsPerBlock); + kNodeParams.sharedMemBytes = 0; + kNodeParams.kernelParams = reinterpret_cast(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 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(&NElem)}; + kNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kNodeParams.gridDim = dim3(blocks); + kNodeParams.blockDim = dim3(threadsPerBlock); + kNodeParams.sharedMemBytes = 0; + kNodeParams.kernelParams = reinterpret_cast(kernelArgs); + kNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, dependencies.data(), + dependencies.size(), &kNodeParams)); + + memset(&kNodeParams1, 0, sizeof(kNodeParams1)); + kNodeParams1.func = reinterpret_cast(HipTest::vectorSUB); + kNodeParams1.gridDim = dim3(blocks); + kNodeParams1.blockDim = dim3(threadsPerBlock); + kNodeParams1.sharedMemBytes = 0; + kNodeParams1.kernelParams = reinterpret_cast(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(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)); +} diff --git a/tests/catch/unit/graph/hipGraphKernelNodeGetParams.cc b/tests/catch/unit/graph/hipGraphKernelNodeGetParams.cc new file mode 100644 index 0000000000..1ccdc62b59 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphKernelNodeGetParams.cc @@ -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 +#include + +#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(&NElem)}; + + kNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kNodeParams.gridDim = dim3(N / THREADS_PER_BLOCK, 1, 1); + kNodeParams.blockDim = dim3(THREADS_PER_BLOCK, 1, 1); + kNodeParams.sharedMemBytes = 0; + kNodeParams.kernelParams = reinterpret_cast(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(p1[i]) != *reinterpret_cast(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(&NElem)}; + + kNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kNodeParams.gridDim = dim3(N / THREADS_PER_BLOCK, 1, 1); + kNodeParams.blockDim = dim3(THREADS_PER_BLOCK, 1, 1); + kNodeParams.sharedMemBytes = 0; + kNodeParams.kernelParams = reinterpret_cast(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(HipTest::vectorADDReverse); + kNodeParams1.gridDim = dim3(N / THREADS_PER_BLOCK, 1, 1); + kNodeParams1.blockDim = dim3(THREADS_PER_BLOCK, 1, 1); + kNodeParams1.sharedMemBytes = 0; + kNodeParams1.kernelParams = reinterpret_cast(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)); +} diff --git a/tests/catch/unit/graph/hipGraphKernelNodeSetParams.cc b/tests/catch/unit/graph/hipGraphKernelNodeSetParams.cc new file mode 100644 index 0000000000..05a6ec345f --- /dev/null +++ b/tests/catch/unit/graph/hipGraphKernelNodeSetParams.cc @@ -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 +#include +#include + +/* 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(&NElem)}; + + kNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kNodeParams.gridDim = dim3(blocks); + kNodeParams.blockDim = dim3(threadsPerBlock); + kNodeParams.sharedMemBytes = 0; + kNodeParams.kernelParams = reinterpret_cast(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 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(&NElem)}; + kNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kNodeParams.gridDim = dim3(blocks); + kNodeParams.blockDim = dim3(threadsPerBlock); + kNodeParams.sharedMemBytes = 0; + kNodeParams.kernelParams = reinterpret_cast(kernelArgs); + kNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, dependencies.data(), + dependencies.size(), &kNodeParams)); + + kNodeParams1.func = reinterpret_cast(HipTest::vectorSUB); + kNodeParams1.gridDim = dim3(blocks); + kNodeParams1.blockDim = dim3(threadsPerBlock); + kNodeParams1.sharedMemBytes = 0; + kNodeParams1.kernelParams = reinterpret_cast(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(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)); +} diff --git a/tests/catch/unit/graph/hipGraphLaunch.cc b/tests/catch/unit/graph/hipGraphLaunch.cc new file mode 100644 index 0000000000..216ec421ef --- /dev/null +++ b/tests/catch/unit/graph/hipGraphLaunch.cc @@ -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 +#include + +/* 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(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(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(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(&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(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 dependencies; + dependencies.push_back(memsetNode); + + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(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(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(&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(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 dependencies; + dependencies.push_back(memsetNode); + + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(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(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."); + } +} diff --git a/tests/catch/unit/graph/hipGraphMemcpyNodeSetParams1D.cc b/tests/catch/unit/graph/hipGraphMemcpyNodeSetParams1D.cc new file mode 100644 index 0000000000..92f51ad3aa --- /dev/null +++ b/tests/catch/unit/graph/hipGraphMemcpyNodeSetParams1D.cc @@ -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 +#include +#include + +/* 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(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(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(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); +} + diff --git a/tests/catch/unit/graph/hipGraphMemcpyNodeSetParamsFromSymbol.cc b/tests/catch/unit/graph/hipGraphMemcpyNodeSetParamsFromSymbol.cc index 495685f4a0..1a6da7ff19 100644 --- a/tests/catch/unit/graph/hipGraphMemcpyNodeSetParamsFromSymbol.cc +++ b/tests/catch/unit/graph/hipGraphMemcpyNodeSetParamsFromSymbol.cc @@ -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++) { diff --git a/tests/catch/unit/memory/hipHostUnregister.cc b/tests/catch/unit/memory/hipHostUnregister.cc index 68a34122b9..2b54c10731 100644 --- a/tests/catch/unit/memory/hipHostUnregister.cc +++ b/tests/catch/unit/memory/hipHostUnregister.cc @@ -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); } diff --git a/tests/catch/unit/memory/hipMallocArray.cc b/tests/catch/unit/memory/hipMallocArray.cc index f1f6ebc9d2..24519f4b1a 100644 --- a/tests/catch/unit/memory/hipMallocArray.cc +++ b/tests/catch/unit/memory/hipMallocArray.cc @@ -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); diff --git a/tests/catch/unit/memory/hipMemsetNegative.cc b/tests/catch/unit/memory/hipMemsetNegative.cc index a06638904e..82a0c2879d 100644 --- a/tests/catch/unit/memory/hipMemsetNegative.cc +++ b/tests/catch/unit/memory/hipMemsetNegative.cc @@ -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}; diff --git a/tests/catch/unit/multiThread/CMakeLists.txt b/tests/catch/unit/multiThread/CMakeLists.txt new file mode 100644 index 0000000000..32abf0f5f3 --- /dev/null +++ b/tests/catch/unit/multiThread/CMakeLists.txt @@ -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) diff --git a/tests/catch/unit/multiThread/hipMultiThreadDevice.cc b/tests/catch/unit/multiThread/hipMultiThreadDevice.cc new file mode 100644 index 0000000000..7550d27a73 --- /dev/null +++ b/tests/catch/unit/multiThread/hipMultiThreadDevice.cc @@ -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 + +#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); +} diff --git a/tests/catch/unit/multiThread/hipMultiThreadStreams1.cc b/tests/catch/unit/multiThread/hipMultiThreadStreams1.cc new file mode 100644 index 0000000000..d4e3af63bd --- /dev/null +++ b/tests/catch/unit/multiThread/hipMultiThreadStreams1.cc @@ -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 +#include + + +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 +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::Copy(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream); + MemTraits::Copy(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream); + MemTraits::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::Copy(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream); + MemTraits::Copy(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream); + + hipLaunchKernelGGL(HipTest::vectorADDReverse, dim3(blocks), dim3(threadsPerBlock), 0, 0, + static_cast(A_d), static_cast(B_d), C_d, numElements); + + MemTraits::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 +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, numElements, p_iters /*iters*/, stream0); + if (serialize) { + t1.join(); + } + std::thread t2(simpleVectorAdd, 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(N /*mb*/, 10 /*iters*/, stream); + simpleVectorAdd(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(NULL, NULL, true); + test_multiThread_1(stream0, stream1, true); + + HIPCHECK(hipStreamDestroy(stream0)); + HIPCHECK(hipStreamDestroy(stream1)); +} +TEST_CASE("Unit_hipMultiThreadStreams1_AsyncSame") { + hipStream_t stream; + HIPCHECK(hipStreamCreate(&stream)); + + // test_multiThread_1 ("Multithread with NULL stream", NULL, + // NULL, false); test_multiThread_1 ("Multithread with two + // streams", stream0, stream1, false); + test_multiThread_1(stream, stream, false); + + HIPCHECK(hipStreamDestroy(stream)); +} diff --git a/tests/catch/unit/multiThread/hipMultiThreadStreams2.cc b/tests/catch/unit/multiThread/hipMultiThreadStreams2.cc new file mode 100644 index 0000000000..d6b6573cb1 --- /dev/null +++ b/tests/catch/unit/multiThread/hipMultiThreadStreams2.cc @@ -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 +#include +#include +#define N 1000 + +template +__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(); + } +} diff --git a/tests/catch/unit/rtc/customOptions.cc b/tests/catch/unit/rtc/customOptions.cc index 040b567100..bfedbd9658 100644 --- a/tests/catch/unit/rtc/customOptions.cc +++ b/tests/catch/unit/rtc/customOptions.cc @@ -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 >"; + 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); +} diff --git a/tests/catch/unit/stream/hipStreamCreateWithPriority.cc b/tests/catch/unit/stream/hipStreamCreateWithPriority.cc index fadc2966a8..300d78ac3c 100644 --- a/tests/catch/unit/stream/hipStreamCreateWithPriority.cc +++ b/tests/catch/unit/stream/hipStreamCreateWithPriority.cc @@ -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; } diff --git a/tests/catch/unit/texture/CMakeLists.txt b/tests/catch/unit/texture/CMakeLists.txt index e9c24a8226..686699a197 100644 --- a/tests/catch/unit/texture/CMakeLists.txt +++ b/tests/catch/unit/texture/CMakeLists.txt @@ -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 diff --git a/tests/catch/unit/texture/hipBindTex2DPitch.cc b/tests/catch/unit/texture/hipBindTex2DPitch.cc index c1d82ebfbe..2a7aab6fa6 100644 --- a/tests/catch/unit/texture/hipBindTex2DPitch.cc +++ b/tests/catch/unit/texture/hipBindTex2DPitch.cc @@ -28,15 +28,19 @@ texture 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; diff --git a/tests/catch/unit/texture/hipBindTexRef1DFetch.cc b/tests/catch/unit/texture/hipBindTexRef1DFetch.cc index e714958163..4e96abc806 100644 --- a/tests/catch/unit/texture/hipBindTexRef1DFetch.cc +++ b/tests/catch/unit/texture/hipBindTexRef1DFetch.cc @@ -24,14 +24,17 @@ THE SOFTWARE. texture 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; diff --git a/tests/catch/unit/texture/hipCreateTextureObject_ArgValidation.cc b/tests/catch/unit/texture/hipCreateTextureObject_ArgValidation.cc index c7570aea03..0ce2d51f2d 100644 --- a/tests/catch/unit/texture/hipCreateTextureObject_ArgValidation.cc +++ b/tests/catch/unit/texture/hipCreateTextureObject_ArgValidation.cc @@ -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; diff --git a/tests/catch/unit/texture/hipCreateTextureObject_Array.cc b/tests/catch/unit/texture/hipCreateTextureObject_Array.cc index da42c7878e..366547d655 100644 --- a/tests/catch/unit/texture/hipCreateTextureObject_Array.cc +++ b/tests/catch/unit/texture/hipCreateTextureObject_Array.cc @@ -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; diff --git a/tests/catch/unit/texture/hipCreateTextureObject_Linear.cc b/tests/catch/unit/texture/hipCreateTextureObject_Linear.cc index 0e2caf6a80..fb12586d37 100644 --- a/tests/catch/unit/texture/hipCreateTextureObject_Linear.cc +++ b/tests/catch/unit/texture/hipCreateTextureObject_Linear.cc @@ -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; diff --git a/tests/catch/unit/texture/hipCreateTextureObject_Pitch2D.cc b/tests/catch/unit/texture/hipCreateTextureObject_Pitch2D.cc index 7fc593708c..fff56d7dd7 100644 --- a/tests/catch/unit/texture/hipCreateTextureObject_Pitch2D.cc +++ b/tests/catch/unit/texture/hipCreateTextureObject_Pitch2D.cc @@ -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; diff --git a/tests/catch/unit/texture/hipGetChanDesc.cc b/tests/catch/unit/texture/hipGetChanDesc.cc index e5c37412ce..b7585059c8 100644 --- a/tests/catch/unit/texture/hipGetChanDesc.cc +++ b/tests/catch/unit/texture/hipGetChanDesc.cc @@ -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)); diff --git a/tests/catch/unit/texture/hipNormalizedFloatValueTex.cc b/tests/catch/unit/texture/hipNormalizedFloatValueTex.cc index 8113117a64..a6eceef38d 100644 --- a/tests/catch/unit/texture/hipNormalizedFloatValueTex.cc +++ b/tests/catch/unit/texture/hipNormalizedFloatValueTex.cc @@ -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); diff --git a/tests/catch/unit/texture/hipSimpleTexture2DLayered.cc b/tests/catch/unit/texture/hipSimpleTexture2DLayered.cc index 0c08c0f6d2..fdf28458d2 100644 --- a/tests/catch/unit/texture/hipSimpleTexture2DLayered.cc +++ b/tests/catch/unit/texture/hipSimpleTexture2DLayered.cc @@ -27,13 +27,17 @@ texture 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; diff --git a/tests/catch/unit/texture/hipSimpleTexture3D.cc b/tests/catch/unit/texture/hipSimpleTexture3D.cc index 1ff46a1f2e..3d59101d51 100644 --- a/tests/catch/unit/texture/hipSimpleTexture3D.cc +++ b/tests/catch/unit/texture/hipSimpleTexture3D.cc @@ -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(i, i, i, &texf); runSimpleTexture3D_Check(i+1, i, i, &texi); diff --git a/tests/catch/unit/texture/hipTex1DFetchCheckModes.cc b/tests/catch/unit/texture/hipTex1DFetchCheckModes.cc index eb3501191c..c3249411b6 100644 --- a/tests/catch/unit/texture/hipTex1DFetchCheckModes.cc +++ b/tests/catch/unit/texture/hipTex1DFetchCheckModes.cc @@ -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(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); + } } diff --git a/tests/catch/unit/texture/hipTexObjPitch.cc b/tests/catch/unit/texture/hipTexObjPitch.cc index 22a58e986d..727d8a64f2 100644 --- a/tests/catch/unit/texture/hipTexObjPitch.cc +++ b/tests/catch/unit/texture/hipTexObjPitch.cc @@ -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++) { diff --git a/tests/catch/unit/texture/hipTextureMipmapObj2D.cc b/tests/catch/unit/texture/hipTextureMipmapObj2D.cc index 7a55f87ca4..7545f2a073 100644 --- a/tests/catch/unit/texture/hipTextureMipmapObj2D.cc +++ b/tests/catch/unit/texture/hipTextureMipmapObj2D.cc @@ -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) { diff --git a/tests/catch/unit/texture/hipTextureObj1DCheckModes.cc b/tests/catch/unit/texture/hipTextureObj1DCheckModes.cc new file mode 100644 index 0000000000..7ff602ab79 --- /dev/null +++ b/tests/catch/unit/texture/hipTextureObj1DCheckModes.cc @@ -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 +#include +#include + +template +__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(textureObject, normalizedCoords ? (x + offsetX) / width : x + offsetX); +#endif +} + +template +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, 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(width, offsetX + j, hData); + if (!hipTextureSamplingVerify(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(256, -3); + runTest(256, 4); + } + + SECTION("hipAddressModeBorder, hipFilterModePoint, regularCoords") { + runTest(256, -8.5); + runTest(256, 12.5); + } + + SECTION("hipAddressModeClamp, hipFilterModeLinear, regularCoords") { + runTest(256, -3); + runTest(256, 4); + } + + SECTION("hipAddressModeBorder, hipFilterModeLinear, regularCoords") { + runTest(256, -8.5); + runTest(256, 12.5); + } + + SECTION("hipAddressModeClamp, hipFilterModePoint, normalizedCoords") { + runTest(256, -3); + runTest(256, 4); + } + + SECTION("hipAddressModeBorder, hipFilterModePoint, normalizedCoords") { + runTest(256, -8.5); + runTest(256, 12.5); + } + + SECTION("hipAddressModeClamp, hipFilterModeLinear, normalizedCoords") { + runTest(256, -3); + runTest(256, 4); + } + + SECTION("hipAddressModeBorder, hipFilterModeLinear, normalizedCoords") { + runTest(256, -8.5); + runTest(256, 12.5); + } +} diff --git a/tests/catch/unit/texture/hipTextureObj1DFetch.cc b/tests/catch/unit/texture/hipTextureObj1DFetch.cc index a52b8baaf4..a7e25c6a40 100644 --- a/tests/catch/unit/texture/hipTextureObj1DFetch.cc +++ b/tests/catch/unit/texture/hipTextureObj1DFetch.cc @@ -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(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(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]; diff --git a/tests/catch/unit/texture/hipTextureObj2D.cc b/tests/catch/unit/texture/hipTextureObj2D.cc index 1836981e41..48efb13bb1 100644 --- a/tests/catch/unit/texture/hipTextureObj2D.cc +++ b/tests/catch/unit/texture/hipTextureObj2D.cc @@ -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(textureObject, x, y); + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + outputData[y * width + x] = tex2D(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; diff --git a/tests/catch/unit/texture/hipTextureObj2DCheckModes.cc b/tests/catch/unit/texture/hipTextureObj2DCheckModes.cc new file mode 100644 index 0000000000..34696b32e3 --- /dev/null +++ b/tests/catch/unit/texture/hipTextureObj2DCheckModes.cc @@ -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 +#include +#include + +template +__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(textureObject, + normalizedCoords ? (x + offsetX) / width : x + offsetX, + normalizedCoords ? (y + offsetY) / height : y + offsetY); +#endif +} + +template +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, 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(width, height, + offsetX + j, offsetY + i, hData); + if (!hipTextureSamplingVerify(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(256, 256, -3.9, 6.1); + runTest(256, 256, 4.4, -7.0); + } + + SECTION("hipAddressModeBorder, hipFilterModePoint, regularCoords") { + runTest(256, 256, -8.5, 2.9); + runTest(256, 256, 12.5, 6.7); + } + + SECTION("hipAddressModeClamp, hipFilterModePoint, regularCoords") { + runTest(256, 256, -0.4, -0.4); + runTest(256, 256, 4, 14.6); + } + + SECTION("hipAddressModeBorder, hipFilterModeLinear, regularCoords") { + runTest(256, 256, -0.4, 0.4); + runTest(256, 256, 12.5, 23.7); + } + + SECTION("hipAddressModeClamp, hipFilterModePoint, normalizedCoords") { + runTest(256, 256, -3, 8.9); + runTest(256, 256, 4, -0.1); + } + + SECTION("hipAddressModeBorder, hipFilterModePoint, normalizedCoords") { + runTest(256, 256, -8.5, 15.9); + runTest(256, 256, 12.5, -17.9); + } + + SECTION("hipAddressModeClamp, hipFilterModeLinear, normalizedCoords") { + runTest(256, 256, -3, 5.8); + runTest(256, 256, 4, 9.1); + } + + SECTION("hipAddressModeBorder, hipFilterModeLinear, normalizedCoords") { + runTest(256, 256, -8.5, 6.6); + runTest(256, 256, 12.5, 0.01); + } +} diff --git a/tests/catch/unit/texture/hipTextureObj3DCheckModes.cc b/tests/catch/unit/texture/hipTextureObj3DCheckModes.cc new file mode 100644 index 0000000000..df5132972b --- /dev/null +++ b/tests/catch/unit/texture/hipTextureObj3DCheckModes.cc @@ -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 +#include +#include + +bool isGfx90a = false; + +template +__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(textureObject, + normalizedCoords ? (x + offsetX) / width : x + offsetX, + normalizedCoords ? (y + offsetY) / height : y + offsetY, + normalizedCoords ? (z + offsetZ) / depth : z + offsetZ); +#endif +} + +template +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(); + 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, 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( + width, height, depth, offsetX + k, offsetY + j, offsetZ + i, hData); + + if (!hipTextureSamplingVerify(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 + (256, 256, 256, -3.9, 6.1, 9.5); + runTest + (256, 256, 256, 4.4, -7.0, 5.3); + } + + SECTION("hipAddressModeBorder, hipFilterModePoint, regularCoords") { + runTest + (256, 256, 256, -8.5, 2.9, 5.8); + runTest + (256, 256, 256, 12.5, 6.7, 11.4); + } + + SECTION("hipAddressModeClamp, hipFilterModeLinear, regularCoords") { + runTest + (256, 256, 256, -0.4, -0.4, -0.4); + runTest + (256, 256, 256, 4, 14.6, -0.3); + } + + SECTION("hipAddressModeBorder, hipFilterModeLinear, regularCoords") { + runTest + (256, 256, 256, 6.9, 7.4, 0.4); + runTest + (256, 256, 256, 12.5, 23.7, 0.34); + } + + SECTION("hipAddressModeClamp, hipFilterModePoint, normalizedCoords") { + runTest + (256, 256, 256, -3, 8.9, -4); + runTest + (256, 256, 256, 4, -0.1, 8.2); + } + + SECTION("hipAddressModeBorder, hipFilterModePoint, normalizedCoords") { + runTest + (256, 256, 256, -8.5, 15.9, 0.1); + runTest + (256, 256, 256, 12.5, -17.9, -0.35); + } + + SECTION("hipAddressModeClamp, hipFilterModeLinear, normalizedCoords") { + runTest + (256, 256, 256, -3, 5.8, 0.89); + runTest + (256, 256, 256, 4, 9.1, 2.08); + } + + SECTION("hipAddressModeBorder, hipFilterModeLinear, normalizedCoords") { + runTest + (256, 256, 256, -8.5, 6.6, 3.67); + runTest + (256, 256, 256, 12.5, 0.01, -9.9); + } +} diff --git a/tests/catch/unit/texture/hipTextureObjFetchVector.cc b/tests/catch/unit/texture/hipTextureObjFetchVector.cc index d3faa5d5b0..616602038a 100644 --- a/tests/catch/unit/texture/hipTextureObjFetchVector.cc +++ b/tests/catch/unit/texture/hipTextureObjFetchVector.cc @@ -138,14 +138,12 @@ static inline bool isEqual(const T &val0, const T &val1) { } template -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"); - runTest("char2"); - runTest("char4"); + runTest(); + runTest(); + runTest(); // test for uchar - runTest("uchar1"); - runTest("uchar2"); - runTest("uchar4"); + runTest(); + runTest(); + runTest(); // test for short - runTest("short1"); - runTest("short2"); - runTest("short4"); + runTest(); + runTest(); + runTest(); // test for ushort - runTest("ushort1"); - runTest("ushort2"); - runTest("ushort4"); + runTest(); + runTest(); + runTest(); // test for int - runTest("int1"); - runTest("int2"); - runTest("int4"); + runTest(); + runTest(); + runTest(); // test for unsigned int - runTest("uint1"); - runTest("uint2"); - runTest("uint4"); + runTest(); + runTest(); + runTest(); // test for float - runTest("float1"); - runTest("float2"); - runTest("float4"); + runTest(); + runTest(); + runTest(); } diff --git a/tests/catch/unit/texture/hipTextureRef2D.cc b/tests/catch/unit/texture/hipTextureRef2D.cc index 1a6600d099..a4abcd4975 100644 --- a/tests/catch/unit/texture/hipTextureRef2D.cc +++ b/tests/catch/unit/texture/hipTextureRef2D.cc @@ -22,12 +22,16 @@ THE SOFTWARE. texture 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; diff --git a/tests/src/texture/hipTextureObj3DCheckModes.cpp b/tests/src/texture/hipTextureObj3DCheckModes.cpp index 25390f6c7e..b96cf86717 100644 --- a/tests/src/texture/hipTextureObj3DCheckModes.cpp +++ b/tests/src/texture/hipTextureObj3DCheckModes.cpp @@ -11,6 +11,7 @@ #include "test_common.h" #include "hipTextureHelper.hpp" +bool isGfx90a = false; template __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(textureObject, + outputData[z * width * height + y * width + x] = tex3D(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( 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(256, 256, 256, -3.9, 6.1, 9.5); testResult = testResult && runTest(256, 256, 256, 4.4, -7.0, 5.3);