SWDEV-299127 - Merge 'develop' into 'amd-staging'
Change-Id: Iaccdfd0e2d82bb989ab3cc933bdc7acc23c59e9a
[ROCm/hip commit: 8c04036f47]
This commit is contained in:
Jenkins
@@ -488,7 +488,7 @@ foreach $arg (@ARGV)
|
||||
my $fileType = `file $obj`;
|
||||
my $isObj = ($fileType =~ m/ELF/ or $fileType =~ m/COFF/);
|
||||
if ($fileType =~ m/ELF/) {
|
||||
my $sections = `readelf -e -W $obj`;
|
||||
my $sections = `$HIP_CLANG_PATH/llvm-readelf -e -W $obj`;
|
||||
$isObj = !($sections =~ m/__CLANG_OFFLOAD_BUNDLE__/);
|
||||
}
|
||||
$allIsObj = ($allIsObj and $isObj);
|
||||
|
||||
@@ -28,12 +28,34 @@ THE SOFTWARE.
|
||||
#if (defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && !(defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
|
||||
|
||||
typedef enum hipDataType {
|
||||
HIP_R_16F = 2,
|
||||
HIP_R_32F = 0,
|
||||
HIP_R_64F = 1,
|
||||
HIP_C_16F = 6,
|
||||
HIP_C_32F = 4,
|
||||
HIP_C_64F = 5
|
||||
HIP_R_32F = 0,
|
||||
HIP_R_64F = 1,
|
||||
HIP_R_16F = 2,
|
||||
HIP_R_8I = 3,
|
||||
HIP_C_32F = 4,
|
||||
HIP_C_64F = 5,
|
||||
HIP_C_16F = 6,
|
||||
HIP_C_8I = 7,
|
||||
HIP_R_8U = 8,
|
||||
HIP_C_8U = 9,
|
||||
HIP_R_32I = 10,
|
||||
HIP_C_32I = 11,
|
||||
HIP_R_32U = 12,
|
||||
HIP_C_32U = 13,
|
||||
HIP_R_16BF = 14,
|
||||
HIP_C_16BF = 15,
|
||||
HIP_R_4I = 16,
|
||||
HIP_C_4I = 17,
|
||||
HIP_R_4U = 18,
|
||||
HIP_C_4U = 19,
|
||||
HIP_R_16I = 20,
|
||||
HIP_C_16I = 21,
|
||||
HIP_R_16U = 22,
|
||||
HIP_C_16U = 23,
|
||||
HIP_R_64I = 24,
|
||||
HIP_C_64I = 25,
|
||||
HIP_R_64U = 26,
|
||||
HIP_C_64U = 27
|
||||
} hipDataType;
|
||||
|
||||
typedef enum hipLibraryPropertyType {
|
||||
|
||||
@@ -2,6 +2,8 @@
|
||||
"DisabledTests":
|
||||
[
|
||||
"Unit_hipStreamPerThread_DeviceReset_1",
|
||||
"Unit_hipMallocManaged_OverSubscription"
|
||||
"Unit_hipMallocManaged_OverSubscription",
|
||||
"Unit_hipDeviceGetCacheConfig_Positive_Basic",
|
||||
"Unit_hipDeviceGetCacheConfig_Positive_Threaded"
|
||||
]
|
||||
}
|
||||
|
||||
@@ -79,6 +79,8 @@
|
||||
"Unit_hipStreamQuery_WithPendingWork",
|
||||
"Unit_hipStreamWaitEvent_DifferentStreams",
|
||||
"Unit_hipStreamQuery_WithFinishedWork",
|
||||
"Unit_hipDeviceGetCacheConfig_Positive_Basic",
|
||||
"Unit_hipDeviceGetCacheConfig_Positive_Threaded",
|
||||
"Unit_hipStreamValue_Wait32_Blocking_Mask_Gte",
|
||||
"Unit_hipStreamValue_Wait32_Blocking_Mask_Eq_1",
|
||||
"Unit_hipStreamValue_Wait32_Blocking_Mask_Eq_2",
|
||||
|
||||
@@ -31,6 +31,8 @@ add_subdirectory(texture)
|
||||
add_subdirectory(streamperthread)
|
||||
add_subdirectory(kernel)
|
||||
add_subdirectory(multiThread)
|
||||
add_subdirectory(compiler)
|
||||
add_subdirectory(errorHandling)
|
||||
if(HIP_PLATFORM STREQUAL "amd")
|
||||
add_subdirectory(clock)
|
||||
endif()
|
||||
|
||||
@@ -0,0 +1,8 @@
|
||||
# Common Tests - Test independent of all platforms
|
||||
set(TEST_SRC
|
||||
hipClassKernel.cc
|
||||
)
|
||||
|
||||
hip_add_exe_to_target(NAME CompilerTest
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME build_tests)
|
||||
@@ -0,0 +1,220 @@
|
||||
/*
|
||||
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
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
#include "hipClassKernel.h"
|
||||
|
||||
__global__ void
|
||||
ovrdClassKernel(bool* result_ecd){
|
||||
int tid = threadIdx.x + blockIdx.x * blockDim.x;
|
||||
testOvrD tobj1;
|
||||
result_ecd[tid] = (tobj1.ovrdFunc1() == 30);
|
||||
}
|
||||
|
||||
__global__ void
|
||||
ovldClassKernel(bool* result_ecd){
|
||||
int tid = threadIdx.x + blockIdx.x * blockDim.x;
|
||||
testFuncOvld tfo1;
|
||||
result_ecd[tid] = (tfo1.func1(10) == 20)
|
||||
&& (tfo1.func1(10,10) == 30);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipClassKernel_Overload_Override") {
|
||||
bool *result_ecd, *result_ech;
|
||||
result_ech = AllocateHostMemory();
|
||||
result_ecd = AllocateDeviceMemory();
|
||||
|
||||
hipLaunchKernelGGL(ovrdClassKernel,
|
||||
dim3(BLOCKS),
|
||||
dim3(THREADS_PER_BLOCK),
|
||||
0,
|
||||
0,
|
||||
result_ecd);
|
||||
|
||||
VerifyResult(result_ech,result_ecd);
|
||||
FreeMem(result_ech,result_ecd);
|
||||
|
||||
result_ech = AllocateHostMemory();
|
||||
result_ecd = AllocateDeviceMemory();
|
||||
|
||||
hipLaunchKernelGGL(ovldClassKernel,
|
||||
dim3(BLOCKS),
|
||||
dim3(THREADS_PER_BLOCK),
|
||||
0,
|
||||
0,
|
||||
result_ecd);
|
||||
|
||||
VerifyResult(result_ech,result_ecd);
|
||||
FreeMem(result_ech,result_ecd);
|
||||
}
|
||||
|
||||
// check for friend
|
||||
__global__ void
|
||||
friendClassKernel(bool* result_ecd){
|
||||
int tid = threadIdx.x + blockIdx.x * blockDim.x;
|
||||
testFrndB tfb1;
|
||||
result_ecd[tid] = (tfb1.showA() == 10);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipClassKernel_Friend") {
|
||||
bool *result_ecd;
|
||||
result_ecd = AllocateDeviceMemory();
|
||||
hipLaunchKernelGGL(friendClassKernel,
|
||||
dim3(BLOCKS),
|
||||
dim3(THREADS_PER_BLOCK),
|
||||
0,
|
||||
0,
|
||||
result_ecd);
|
||||
}
|
||||
|
||||
// check sizeof empty class is 1
|
||||
__global__ void
|
||||
emptyClassKernel(bool* result_ecd) {
|
||||
int tid = threadIdx.x + blockIdx.x * blockDim.x;
|
||||
testClassEmpty ob1,ob2;
|
||||
result_ecd[tid] = (sizeof(testClassEmpty) == 1)
|
||||
&& (&ob1 != &ob2);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipClassKernel_Empty") {
|
||||
bool *result_ecd, *result_ech;
|
||||
result_ech = AllocateHostMemory();
|
||||
result_ecd = AllocateDeviceMemory();
|
||||
|
||||
hipLaunchKernelGGL(emptyClassKernel,
|
||||
dim3(BLOCKS),
|
||||
dim3(THREADS_PER_BLOCK),
|
||||
0,
|
||||
0,
|
||||
result_ecd);
|
||||
|
||||
VerifyResult(result_ech,result_ecd);
|
||||
FreeMem(result_ech,result_ecd);
|
||||
}
|
||||
|
||||
// tests for classes >8 bytes
|
||||
__global__ void
|
||||
sizeClassBKernel(bool* result_ecd) {
|
||||
int tid = threadIdx.x + blockIdx.x * blockDim.x;
|
||||
result_ecd[tid] = (sizeof(testSizeB) == 12)
|
||||
&& (sizeof(testSizeC) == 16)
|
||||
&& (sizeof(testSizeP1) == 6)
|
||||
&& (sizeof(testSizeP2) == 13)
|
||||
&& (sizeof(testSizeP3) == 8);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipClassKernel_BSize") {
|
||||
bool *result_ecd, *result_ech;
|
||||
result_ech = AllocateHostMemory();
|
||||
result_ecd = AllocateDeviceMemory();
|
||||
|
||||
hipLaunchKernelGGL(sizeClassBKernel,
|
||||
dim3(BLOCKS),
|
||||
dim3(THREADS_PER_BLOCK),
|
||||
0,
|
||||
0,
|
||||
result_ecd);
|
||||
|
||||
VerifyResult(result_ech,result_ecd);
|
||||
FreeMem(result_ech,result_ecd);
|
||||
}
|
||||
|
||||
__global__ void
|
||||
sizeClassKernel(bool* result_ecd) {
|
||||
int tid = threadIdx.x + blockIdx.x * blockDim.x;
|
||||
result_ecd[tid] = (sizeof(testSizeA) == 16)
|
||||
&& (sizeof(testSizeDerived) == 24)
|
||||
&& (sizeof(testSizeDerived2) == 20);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipClassKernel_Size") {
|
||||
bool *result_ecd, *result_ech;
|
||||
result_ech = AllocateHostMemory();
|
||||
result_ecd = AllocateDeviceMemory();
|
||||
|
||||
hipLaunchKernelGGL(sizeClassKernel,
|
||||
dim3(BLOCKS),
|
||||
dim3(THREADS_PER_BLOCK),
|
||||
0,
|
||||
0,
|
||||
result_ecd);
|
||||
|
||||
VerifyResult(result_ech,result_ecd);
|
||||
FreeMem(result_ech,result_ecd);
|
||||
}
|
||||
|
||||
__global__ void
|
||||
sizeVirtualClassKernel(bool* result_ecd) {
|
||||
int tid = threadIdx.x + blockIdx.x * blockDim.x;
|
||||
result_ecd[tid] = (sizeof(testSizeDV) == 16)
|
||||
&& (sizeof(testSizeDerivedDV) == 16)
|
||||
&& (sizeof(testSizeVirtDerPack) == 24)
|
||||
&& (sizeof(testSizeVirtDer) == 24)
|
||||
&& (sizeof(testSizeDerMulti) == 48) ;
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipClassKernel_Virtual") {
|
||||
bool *result_ecd, *result_ech;
|
||||
result_ech = AllocateHostMemory();
|
||||
result_ecd = AllocateDeviceMemory();
|
||||
|
||||
hipLaunchKernelGGL(sizeVirtualClassKernel,
|
||||
dim3(BLOCKS),
|
||||
dim3(THREADS_PER_BLOCK),
|
||||
0,
|
||||
0,
|
||||
result_ecd);
|
||||
|
||||
VerifyResult(result_ech,result_ecd);
|
||||
FreeMem(result_ech,result_ecd);
|
||||
}
|
||||
|
||||
// check pass by value
|
||||
__global__ void
|
||||
passByValueKernel(testPassByValue obj, bool* result_ecd) {
|
||||
int tid = threadIdx.x + blockIdx.x * blockDim.x;
|
||||
result_ecd[tid] = (obj.exI == 10)
|
||||
&& (obj.exC == 'C');
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipClassKernel_Value") {
|
||||
bool *result_ecd,*result_ech;
|
||||
result_ech = AllocateHostMemory();
|
||||
result_ecd = AllocateDeviceMemory();
|
||||
|
||||
testPassByValue exObj;
|
||||
exObj.exI = 10;
|
||||
exObj.exC = 'C';
|
||||
hipLaunchKernelGGL(passByValueKernel,
|
||||
dim3(BLOCKS),
|
||||
dim3(THREADS_PER_BLOCK),
|
||||
0,
|
||||
0,
|
||||
exObj,
|
||||
result_ecd);
|
||||
|
||||
VerifyResult(result_ech,result_ecd);
|
||||
FreeMem(result_ech,result_ecd);
|
||||
}
|
||||
@@ -0,0 +1,234 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
#ifndef _COMPILER_HIPCLASSKERNEL_H_
|
||||
#define _COMPILER_HIPCLASSKERNEL_H_
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
static const int BLOCKS = 512;
|
||||
static const int THREADS_PER_BLOCK = 1;
|
||||
size_t NBOOL = BLOCKS * sizeof(bool);
|
||||
|
||||
class testFuncOvld{
|
||||
public:
|
||||
int __host__ __device__ func1(int a){
|
||||
return a + 10;
|
||||
}
|
||||
|
||||
int __host__ __device__ func1(int a , int b){
|
||||
return a + b + 10;
|
||||
}
|
||||
|
||||
};
|
||||
|
||||
class testOvrB{
|
||||
public:
|
||||
int __host__ __device__ ovrdFunc1(){
|
||||
return 10;
|
||||
}
|
||||
|
||||
|
||||
};
|
||||
|
||||
class testOvrD: public testOvrB{
|
||||
public:
|
||||
int __host__ __device__ ovrdFunc1(){
|
||||
int x = testOvrB::ovrdFunc1();
|
||||
return x + 20;
|
||||
}
|
||||
|
||||
};
|
||||
|
||||
class testFrndA{
|
||||
private:
|
||||
int fa = 10;
|
||||
public:
|
||||
friend class testFrndB;
|
||||
};
|
||||
|
||||
class testFrndB{
|
||||
public:
|
||||
__host__ __device__ int showA(){
|
||||
testFrndA x;
|
||||
return x.fa;
|
||||
}
|
||||
};
|
||||
|
||||
class testClassEmpty {};
|
||||
|
||||
class testPassByValue{
|
||||
public:
|
||||
int exI;
|
||||
char exC;
|
||||
};
|
||||
|
||||
class testSizeA {
|
||||
public:
|
||||
float xa;
|
||||
int ia;
|
||||
double da;
|
||||
static char ca;
|
||||
};
|
||||
|
||||
class testSizeDerived : testSizeA {
|
||||
public:
|
||||
float fd;
|
||||
};
|
||||
|
||||
#pragma pack(push,4)
|
||||
class testSizeDerived2 : testSizeA {
|
||||
public:
|
||||
float fd;
|
||||
};
|
||||
#pragma pack(pop)
|
||||
|
||||
class testSizeB {
|
||||
public:
|
||||
char ab;
|
||||
int ib;
|
||||
char cb;
|
||||
};
|
||||
|
||||
class testSizeVirtDer : public virtual testSizeB {
|
||||
public:
|
||||
int ivd;
|
||||
};
|
||||
|
||||
class testSizeVirtDer1 : public virtual testSizeB {
|
||||
public:
|
||||
int ivd1;
|
||||
};
|
||||
|
||||
class testSizeDerMulti : public testSizeVirtDer, public testSizeVirtDer1 {
|
||||
public:
|
||||
int ivd2;
|
||||
};
|
||||
|
||||
#pragma pack(push,4)
|
||||
class testSizeVirtDerPack : public virtual testSizeB {
|
||||
public:
|
||||
int ivd;
|
||||
};
|
||||
#pragma pack(pop)
|
||||
|
||||
class testSizeC {
|
||||
public:
|
||||
char ac;
|
||||
int ic;
|
||||
int bc[2];
|
||||
};
|
||||
|
||||
class testSizeDV {
|
||||
public:
|
||||
virtual void __host__ __device__ func1();
|
||||
private:
|
||||
int iDV;
|
||||
|
||||
};
|
||||
|
||||
class testSizeDerivedDV : testSizeDV {
|
||||
public:
|
||||
virtual void __host__ __device__ funcD1();
|
||||
private:
|
||||
int iDDV;
|
||||
};
|
||||
|
||||
#pragma pack(push, 1)
|
||||
class testSizeP1 {
|
||||
public:
|
||||
char ap;
|
||||
int ip;
|
||||
char cp;
|
||||
};
|
||||
#pragma pack(pop)
|
||||
|
||||
#pragma pack(push, 1)
|
||||
class testSizeP2 {
|
||||
public:
|
||||
char ap1;
|
||||
int ip1;
|
||||
int bp1[2];
|
||||
};
|
||||
#pragma pack(pop)
|
||||
|
||||
#pragma pack(push, 2)
|
||||
class testSizeP3 {
|
||||
public:
|
||||
char ap2;
|
||||
int ip2;
|
||||
char cp2;
|
||||
};
|
||||
#pragma pack(pop)
|
||||
|
||||
class testDeviceClass {
|
||||
public:
|
||||
int iVar;
|
||||
__host__ __device__ testDeviceClass();
|
||||
__host__ __device__ testDeviceClass(int a);
|
||||
__host__ __device__ ~testDeviceClass();
|
||||
};
|
||||
|
||||
__host__ __device__
|
||||
testDeviceClass::testDeviceClass() {
|
||||
iVar = 5;
|
||||
}
|
||||
|
||||
__host__ __device__
|
||||
testDeviceClass::testDeviceClass(int a) {
|
||||
iVar = a;
|
||||
}
|
||||
|
||||
bool* AllocateHostMemory(void){
|
||||
bool *result_ech;
|
||||
HIPCHECK(hipHostMalloc(&result_ech,
|
||||
NBOOL,
|
||||
hipHostMallocDefault));
|
||||
return result_ech;
|
||||
}
|
||||
|
||||
bool* AllocateDeviceMemory(void){
|
||||
bool* result_ecd;
|
||||
HIPCHECK(hipMalloc(&result_ecd,
|
||||
NBOOL));
|
||||
HIPCHECK(hipMemset(result_ecd,
|
||||
false,
|
||||
NBOOL));
|
||||
return result_ecd;
|
||||
}
|
||||
|
||||
void VerifyResult(bool* result_ech, bool* result_ecd){
|
||||
HIPCHECK(hipMemcpy(result_ech,
|
||||
result_ecd,
|
||||
BLOCKS*sizeof(bool),
|
||||
hipMemcpyDeviceToHost));
|
||||
// validation on host side
|
||||
for (int i = 0; i < BLOCKS; i++) {
|
||||
HIPASSERT(result_ech[i] == true);
|
||||
}
|
||||
}
|
||||
|
||||
void FreeMem(bool* result_ech, bool* result_ecd){
|
||||
HIPCHECK(hipHostFree(result_ech));
|
||||
HIPCHECK(hipFree(result_ecd));
|
||||
}
|
||||
|
||||
#endif // _HIPCLASSKERNEL_H_
|
||||
@@ -17,8 +17,10 @@ set(TEST_SRC
|
||||
hipSetGetDevice.cc
|
||||
hipDeviceGetUuid.cc
|
||||
hipDeviceGetP2PAttribute.cc
|
||||
hipDeviceGetDefaultMemPool.cc
|
||||
hipExtGetLinkTypeAndHopCount.cc
|
||||
hipDeviceSetLimit.cc
|
||||
hipDeviceSetGetMemPool.cc
|
||||
)
|
||||
|
||||
set_source_files_properties(hipGetDeviceCount.cc PROPERTIES COMPILE_FLAGS -std=c++17)
|
||||
|
||||
@@ -0,0 +1,57 @@
|
||||
/*
|
||||
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip/hip_runtime_api.h>
|
||||
|
||||
TEST_CASE("Unit_hipDeviceGetDefaultMemPool_Positive_Basic") {
|
||||
const int device = GENERATE(range(0, HipTest::getDeviceCount()));
|
||||
|
||||
int mem_pool_support = 0;
|
||||
HIP_CHECK(
|
||||
hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, device));
|
||||
if (!mem_pool_support) {
|
||||
HipTest::HIP_SKIP_TEST("Test only runs on devices with memory pool support");
|
||||
return;
|
||||
}
|
||||
|
||||
hipMemPool_t mem_pool;
|
||||
HIP_CHECK(hipDeviceGetDefaultMemPool(&mem_pool, device));
|
||||
REQUIRE(mem_pool != nullptr);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipDeviceGetDefaultMemPool_Negative_Parameters") {
|
||||
hipMemPool_t mem_pool;
|
||||
|
||||
SECTION("mem_pool == nullptr") {
|
||||
HIP_CHECK_ERROR(hipDeviceGetDefaultMemPool(nullptr, 0), hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("device < 0") {
|
||||
HIP_CHECK_ERROR(hipDeviceGetDefaultMemPool(&mem_pool, -1), hipErrorInvalidDevice);
|
||||
}
|
||||
|
||||
SECTION("device ordinance too large") {
|
||||
HIP_CHECK_ERROR(hipDeviceGetDefaultMemPool(&mem_pool, HipTest::getDeviceCount()),
|
||||
hipErrorInvalidDevice);
|
||||
}
|
||||
}
|
||||
@@ -1,45 +1,108 @@
|
||||
/*
|
||||
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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, INNCLUDING 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 ANNY 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 IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
#include <hip_test_common.hh>
|
||||
/**
|
||||
* hipDeviceGetCacheConfig tests
|
||||
* Scenario1: Validates if pConfig = nullptr returns hip error code.
|
||||
* Scenario2: Validates if the value returned by hipDeviceGetCacheConfig is valid.
|
||||
*/
|
||||
TEST_CASE("Unit_hipDeviceGetCacheConfig_NegTst") {
|
||||
// Scenario1
|
||||
REQUIRE_FALSE(hipSuccess == hipDeviceGetCacheConfig(nullptr));
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipDeviceGetCacheConfig_FuncTst") {
|
||||
hipFuncCache_t cacheConfig;
|
||||
// Scenario2
|
||||
HIP_CHECK(hipDeviceGetCacheConfig(&cacheConfig));
|
||||
REQUIRE_FALSE(((cacheConfig != hipFuncCachePreferNone) &&
|
||||
(cacheConfig != hipFuncCachePreferShared) &&
|
||||
(cacheConfig != hipFuncCachePreferL1) &&
|
||||
(cacheConfig != hipFuncCachePreferEqual)));
|
||||
// This code exists to test the dummy implementation of
|
||||
// hipDeviceSetCacheConfig.
|
||||
#if HT_NVIDIA
|
||||
HIP_CHECK(hipDeviceSetCacheConfig(cacheConfig));
|
||||
#else
|
||||
REQUIRE_FALSE(hipSuccess == hipDeviceSetCacheConfig(cacheConfig));
|
||||
#include <array>
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <threaded_zig_zag_test.hh>
|
||||
|
||||
namespace {
|
||||
constexpr std::array<hipFuncCache_t, 4> kCacheConfigs{
|
||||
hipFuncCachePreferNone, hipFuncCachePreferShared, hipFuncCachePreferL1,
|
||||
hipFuncCachePreferEqual};
|
||||
} // anonymous namespace
|
||||
|
||||
|
||||
TEST_CASE("Unit_hipDeviceSetCacheConfig_Positive_Basic") {
|
||||
const auto device = GENERATE(range(0, HipTest::getDeviceCount()));
|
||||
HIP_CHECK(hipSetDevice(device));
|
||||
INFO("Current device is: " << device);
|
||||
|
||||
const auto cache_config =
|
||||
GENERATE(from_range(std::begin(kCacheConfigs), std::end(kCacheConfigs)));
|
||||
#if HT_AMD
|
||||
HIP_CHECK_ERROR(hipDeviceSetCacheConfig(cache_config), hipErrorNotSupported);
|
||||
#elif HT_NVIDIA
|
||||
HIP_CHECK(hipDeviceSetCacheConfig(cache_config));
|
||||
#endif
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipDeviceSetCacheConfig_Negative_Parameters") {
|
||||
#if HT_AMD
|
||||
HIP_CHECK_ERROR(hipDeviceSetCacheConfig(static_cast<hipFuncCache_t>(-1)), hipErrorNotSupported);
|
||||
#elif HT_NVIDIA
|
||||
HIP_CHECK_ERROR(hipDeviceSetCacheConfig(static_cast<hipFuncCache_t>(-1)), hipErrorInvalidValue);
|
||||
#endif
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipDeviceGetCacheConfig_Positive_Default") {
|
||||
const auto device = GENERATE(range(0, HipTest::getDeviceCount()));
|
||||
HIP_CHECK(hipSetDevice(device));
|
||||
INFO("Current device is: " << device);
|
||||
|
||||
hipFuncCache_t cache_config;
|
||||
HIP_CHECK(hipDeviceGetCacheConfig(&cache_config));
|
||||
REQUIRE(cache_config == hipFuncCachePreferNone);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipDeviceGetCacheConfig_Positive_Basic") {
|
||||
const auto device = GENERATE(range(0, HipTest::getDeviceCount()));
|
||||
HIP_CHECK(hipSetDevice(device));
|
||||
INFO("Current device is: " << device);
|
||||
|
||||
const auto cache_config =
|
||||
GENERATE(from_range(std::begin(kCacheConfigs), std::end(kCacheConfigs)));
|
||||
|
||||
HIP_CHECK(hipDeviceSetCacheConfig(cache_config));
|
||||
hipFuncCache_t returned_cache_config;
|
||||
HIP_CHECK(hipDeviceGetCacheConfig(&returned_cache_config));
|
||||
|
||||
REQUIRE(returned_cache_config == cache_config);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipDeviceGetCacheConfig_Positive_Threaded") {
|
||||
class HipDeviceSetGetCacheConfigTest : public ThreadedZigZagTest<HipDeviceSetGetCacheConfigTest> {
|
||||
public:
|
||||
HipDeviceSetGetCacheConfigTest(const hipFuncCache_t cache_config)
|
||||
: cache_config_{cache_config} {}
|
||||
|
||||
void TestPart2() { HIP_CHECK_THREAD(hipDeviceSetCacheConfig(cache_config_)); }
|
||||
|
||||
void TestPart3() {
|
||||
hipFuncCache_t returned_cache_config;
|
||||
HIP_CHECK(hipDeviceGetCacheConfig(&returned_cache_config));
|
||||
REQUIRE(returned_cache_config == cache_config_);
|
||||
}
|
||||
|
||||
private:
|
||||
const hipFuncCache_t cache_config_;
|
||||
};
|
||||
|
||||
const auto cache_config =
|
||||
GENERATE(from_range(std::begin(kCacheConfigs), std::end(kCacheConfigs)));
|
||||
|
||||
HipDeviceSetGetCacheConfigTest test(cache_config);
|
||||
test.run();
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_HipDeviceGetCacheConfig_Negative_Parameters") {
|
||||
HIP_CHECK_ERROR(hipDeviceGetCacheConfig(nullptr), hipErrorInvalidValue);
|
||||
}
|
||||
@@ -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.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip/hip_runtime_api.h>
|
||||
#include <threaded_zig_zag_test.hh>
|
||||
|
||||
static inline bool CheckMemPoolSupport(const int device) {
|
||||
int mem_pool_support = 0;
|
||||
HIP_CHECK(
|
||||
hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, device));
|
||||
if (!mem_pool_support) {
|
||||
HipTest::HIP_SKIP_TEST("Test only runs on devices with memory pool support");
|
||||
return false;
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
||||
static inline hipMemPool_t CreateMemPool(const int device) {
|
||||
hipMemPoolProps kPoolProps;
|
||||
kPoolProps.allocType = hipMemAllocationTypePinned;
|
||||
kPoolProps.handleTypes = hipMemHandleTypeNone;
|
||||
kPoolProps.location.type = hipMemLocationTypeDevice;
|
||||
kPoolProps.location.id = device;
|
||||
kPoolProps.win32SecurityAttributes = nullptr;
|
||||
memset(kPoolProps.reserved, 0, sizeof(kPoolProps.reserved));
|
||||
|
||||
hipMemPool_t mem_pool;
|
||||
HIP_CHECK(hipMemPoolCreate(&mem_pool, &kPoolProps));
|
||||
|
||||
return mem_pool;
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipDeviceSetMemPool_Positive_Basic") {
|
||||
const int device = GENERATE(range(0, HipTest::getDeviceCount()));
|
||||
|
||||
if (!CheckMemPoolSupport(device)) {
|
||||
return;
|
||||
}
|
||||
|
||||
hipMemPool_t mem_pool = CreateMemPool(device);
|
||||
HIP_CHECK(hipDeviceSetMemPool(device, mem_pool));
|
||||
|
||||
HIP_CHECK(hipMemPoolDestroy(mem_pool));
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipDeviceSetMemPool_Negative_Parameters") {
|
||||
hipMemPool_t mem_pool;
|
||||
HIP_CHECK(hipDeviceGetDefaultMemPool(&mem_pool, 0));
|
||||
|
||||
SECTION("mem_pool == nullptr") {
|
||||
HIP_CHECK_ERROR(hipDeviceSetMemPool(0, nullptr), hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("device < 0") {
|
||||
HIP_CHECK_ERROR(hipDeviceSetMemPool(-1, mem_pool), hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("device ordinance too large") {
|
||||
HIP_CHECK_ERROR(hipDeviceSetMemPool(HipTest::getDeviceCount(), mem_pool), hipErrorInvalidValue);
|
||||
}
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipDeviceGetMemPool_Positive_Default") {
|
||||
const int device = GENERATE(range(0, HipTest::getDeviceCount()));
|
||||
|
||||
if (!CheckMemPoolSupport(device)) {
|
||||
return;
|
||||
}
|
||||
|
||||
hipMemPool_t default_mem_pool;
|
||||
HIP_CHECK(hipDeviceGetDefaultMemPool(&default_mem_pool, device));
|
||||
|
||||
hipMemPool_t mem_pool;
|
||||
HIP_CHECK(hipDeviceGetMemPool(&mem_pool, device));
|
||||
|
||||
REQUIRE(mem_pool == default_mem_pool);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipDeviceGetMemPool_Positive_Basic") {
|
||||
const int device = GENERATE(range(0, HipTest::getDeviceCount()));
|
||||
|
||||
if (!CheckMemPoolSupport(device)) {
|
||||
return;
|
||||
}
|
||||
|
||||
hipMemPool_t mem_pool = CreateMemPool(device);
|
||||
HIP_CHECK(hipDeviceSetMemPool(device, mem_pool));
|
||||
|
||||
hipMemPool_t returned_mem_pool;
|
||||
HIP_CHECK(hipDeviceGetMemPool(&returned_mem_pool, device));
|
||||
|
||||
REQUIRE(returned_mem_pool == mem_pool);
|
||||
|
||||
HIP_CHECK(hipMemPoolDestroy(mem_pool));
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipDeviceGetMemPool_Positive_Threaded") {
|
||||
class HipDeviceGetMemPoolTest : public ThreadedZigZagTest<HipDeviceGetMemPoolTest> {
|
||||
public:
|
||||
void TestPart2() {
|
||||
mem_pool_ = CreateMemPool(0);
|
||||
HIP_CHECK_THREAD(hipDeviceSetMemPool(0, mem_pool_));
|
||||
}
|
||||
void TestPart3() {
|
||||
hipMemPool_t returned_mem_pool;
|
||||
HIP_CHECK(hipDeviceGetMemPool(&returned_mem_pool, 0));
|
||||
|
||||
REQUIRE(returned_mem_pool == mem_pool_);
|
||||
}
|
||||
void TestPart4() { HIP_CHECK_THREAD(hipMemPoolDestroy(mem_pool_)); }
|
||||
|
||||
private:
|
||||
hipMemPool_t mem_pool_;
|
||||
};
|
||||
|
||||
if (!CheckMemPoolSupport(0)) {
|
||||
return;
|
||||
}
|
||||
|
||||
HipDeviceGetMemPoolTest test;
|
||||
test.run();
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipDeviceGetMemPool_Negative_Parameters") {
|
||||
hipMemPool_t mem_pool;
|
||||
|
||||
SECTION("mem_pool == nullptr") {
|
||||
HIP_CHECK_ERROR(hipDeviceGetMemPool(nullptr, 0), hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("device < 0") {
|
||||
HIP_CHECK_ERROR(hipDeviceGetMemPool(&mem_pool, -1), hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("device ordinance too large") {
|
||||
HIP_CHECK_ERROR(hipDeviceGetMemPool(&mem_pool, HipTest::getDeviceCount()),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
}
|
||||
@@ -1,19 +1,22 @@
|
||||
/*
|
||||
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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, INNCLUDING 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 ANNY 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 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.
|
||||
*/
|
||||
|
||||
@@ -41,6 +44,36 @@ static __global__ void Iter(int* Ad, int num) {
|
||||
}
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipDeviceSynchronize_Positive_Empty_Streams") {
|
||||
const auto device = GENERATE(range(0, HipTest::getDeviceCount()));
|
||||
HIP_CHECK(hipSetDevice(device));
|
||||
INFO("Current device: " << device);
|
||||
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipDeviceSynchronize_Positive_Nullstream") {
|
||||
const auto device = GENERATE(range(0, HipTest::getDeviceCount()));
|
||||
HIP_CHECK(hipSetDevice(device));
|
||||
INFO("Current device: " << device);
|
||||
|
||||
int *A_h = nullptr, *A_d = nullptr;
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), _SIZE, hipHostMallocDefault));
|
||||
A_h[0] = 1;
|
||||
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&A_d), _SIZE));
|
||||
|
||||
HIP_CHECK(hipMemcpyAsync(A_d, A_h, _SIZE, hipMemcpyHostToDevice, NULL));
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(Iter), dim3(1), dim3(1), 0, NULL, A_d, 1 << 30);
|
||||
HIP_CHECK(hipMemcpyAsync(A_h, A_d, _SIZE, hipMemcpyDeviceToHost, NULL));
|
||||
|
||||
CHECK(1 << 30 != A_h[0] - 1);
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
CHECK(1 << 30 == A_h[0] - 1);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipDeviceSynchronize_Functional") {
|
||||
int* A[NUM_STREAMS];
|
||||
int* Ad[NUM_STREAMS];
|
||||
|
||||
@@ -0,0 +1,12 @@
|
||||
# Common Tests - Test independent of all platforms
|
||||
set(TEST_SRC
|
||||
hipGetErrorName.cc
|
||||
hipGetErrorString.cc
|
||||
hipGetLastError.cc
|
||||
hipPeekAtLastError.cc
|
||||
)
|
||||
|
||||
hip_add_exe_to_target(NAME ErrorHandlingTest
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME build_tests
|
||||
COMPILE_OPTIONS -std=c++14)
|
||||
@@ -0,0 +1,99 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <hip/hip_runtime_api.h>
|
||||
|
||||
constexpr hipError_t kErrorEnumerators[] = {hipSuccess,
|
||||
hipErrorInvalidValue,
|
||||
hipErrorOutOfMemory,
|
||||
hipErrorNotInitialized,
|
||||
hipErrorDeinitialized,
|
||||
hipErrorProfilerDisabled,
|
||||
hipErrorProfilerNotInitialized,
|
||||
hipErrorProfilerAlreadyStarted,
|
||||
hipErrorProfilerAlreadyStopped,
|
||||
hipErrorInvalidConfiguration,
|
||||
hipErrorInvalidPitchValue,
|
||||
hipErrorInvalidSymbol,
|
||||
hipErrorInvalidDevicePointer,
|
||||
hipErrorInvalidMemcpyDirection,
|
||||
hipErrorInsufficientDriver,
|
||||
hipErrorMissingConfiguration,
|
||||
hipErrorPriorLaunchFailure,
|
||||
hipErrorInvalidDeviceFunction,
|
||||
hipErrorNoDevice,
|
||||
hipErrorInvalidDevice,
|
||||
hipErrorInvalidImage,
|
||||
hipErrorInvalidContext,
|
||||
hipErrorContextAlreadyCurrent,
|
||||
hipErrorMapFailed,
|
||||
hipErrorUnmapFailed,
|
||||
hipErrorArrayIsMapped,
|
||||
hipErrorAlreadyMapped,
|
||||
hipErrorNoBinaryForGpu,
|
||||
hipErrorAlreadyAcquired,
|
||||
hipErrorNotMapped,
|
||||
hipErrorNotMappedAsArray,
|
||||
hipErrorNotMappedAsPointer,
|
||||
hipErrorECCNotCorrectable,
|
||||
hipErrorUnsupportedLimit,
|
||||
hipErrorContextAlreadyInUse,
|
||||
hipErrorPeerAccessUnsupported,
|
||||
hipErrorInvalidKernelFile,
|
||||
hipErrorInvalidGraphicsContext,
|
||||
hipErrorInvalidSource,
|
||||
hipErrorFileNotFound,
|
||||
hipErrorSharedObjectSymbolNotFound,
|
||||
hipErrorSharedObjectInitFailed,
|
||||
hipErrorOperatingSystem,
|
||||
hipErrorInvalidHandle,
|
||||
hipErrorIllegalState,
|
||||
hipErrorNotFound,
|
||||
hipErrorNotReady,
|
||||
hipErrorIllegalAddress,
|
||||
hipErrorLaunchOutOfResources,
|
||||
hipErrorLaunchTimeOut,
|
||||
hipErrorPeerAccessAlreadyEnabled,
|
||||
hipErrorPeerAccessNotEnabled,
|
||||
hipErrorSetOnActiveProcess,
|
||||
hipErrorContextIsDestroyed,
|
||||
hipErrorAssert,
|
||||
hipErrorHostMemoryAlreadyRegistered,
|
||||
hipErrorHostMemoryNotRegistered,
|
||||
hipErrorLaunchFailure,
|
||||
hipErrorCooperativeLaunchTooLarge,
|
||||
hipErrorNotSupported,
|
||||
hipErrorStreamCaptureUnsupported,
|
||||
hipErrorStreamCaptureInvalidated,
|
||||
hipErrorStreamCaptureMerge,
|
||||
hipErrorStreamCaptureUnmatched,
|
||||
hipErrorStreamCaptureUnjoined,
|
||||
hipErrorStreamCaptureIsolation,
|
||||
hipErrorStreamCaptureImplicit,
|
||||
hipErrorCapturedEvent,
|
||||
hipErrorStreamCaptureWrongThread,
|
||||
hipErrorGraphExecUpdateFailure,
|
||||
hipErrorUnknown,
|
||||
hipErrorRuntimeMemory,
|
||||
hipErrorRuntimeOther};
|
||||
@@ -0,0 +1,47 @@
|
||||
/*
|
||||
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 "errorEnumerators.h"
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip/hip_runtime_api.h>
|
||||
|
||||
TEST_CASE("Unit_hipGetErrorName_Positive_Basic") {
|
||||
const char* error_string = nullptr;
|
||||
const auto enumerator =
|
||||
GENERATE(from_range(std::begin(kErrorEnumerators), std::end(kErrorEnumerators)));
|
||||
|
||||
error_string = hipGetErrorName(enumerator);
|
||||
|
||||
REQUIRE(error_string != nullptr);
|
||||
REQUIRE(strlen(error_string) > 0);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipGetErrorName_Negative_Parameters") {
|
||||
const char* error_string = hipGetErrorName(static_cast<hipError_t>(-1));
|
||||
REQUIRE(error_string != nullptr);
|
||||
#if HT_NVIDIA
|
||||
REQUIRE_THAT(error_string, Catch::Equals("cudaErrorUnknown"));
|
||||
#elif HT_AMD
|
||||
REQUIRE_THAT(error_string, Catch::Equals("hipErrorUnknown"));
|
||||
#endif
|
||||
}
|
||||
@@ -0,0 +1,41 @@
|
||||
/*
|
||||
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 "errorEnumerators.h"
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip/hip_runtime_api.h>
|
||||
|
||||
TEST_CASE("Unit_hipGetErrorString_Positive_Basic") {
|
||||
const char* error_string = nullptr;
|
||||
const auto enumerator =
|
||||
GENERATE(from_range(std::begin(kErrorEnumerators), std::end(kErrorEnumerators)));
|
||||
|
||||
error_string = hipGetErrorString(enumerator);
|
||||
|
||||
REQUIRE(error_string != nullptr);
|
||||
REQUIRE(strlen(error_string) > 0);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipGetErrorString_Negative_Parameters") {
|
||||
const char* error_string = hipGetErrorString(static_cast<hipError_t>(-1));
|
||||
REQUIRE(error_string != nullptr);
|
||||
}
|
||||
@@ -0,0 +1,44 @@
|
||||
/*
|
||||
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip/hip_runtime_api.h>
|
||||
#include <threaded_zig_zag_test.hh>
|
||||
|
||||
TEST_CASE("Unit_hipGetLastError_Positive_Basic") {
|
||||
HIP_CHECK(hipGetLastError());
|
||||
HIP_CHECK_ERROR(hipMalloc(nullptr, 1), hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidValue);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipGetLastError_Positive_Threaded") {
|
||||
class HipGetLastErrorThreadedTest : public ThreadedZigZagTest<HipGetLastErrorThreadedTest> {
|
||||
public:
|
||||
void TestPart2() { REQUIRE_THREAD(hipMalloc(nullptr, 1) == hipErrorInvalidValue); }
|
||||
void TestPart3() { HIP_CHECK(hipGetLastError()); }
|
||||
void TestPart4() { REQUIRE_THREAD(hipGetLastError() == hipErrorInvalidValue); }
|
||||
};
|
||||
|
||||
HipGetLastErrorThreadedTest test;
|
||||
test.run();
|
||||
}
|
||||
@@ -0,0 +1,48 @@
|
||||
/*
|
||||
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip/hip_runtime_api.h>
|
||||
#include <threaded_zig_zag_test.hh>
|
||||
|
||||
TEST_CASE("Unit_hipPeekAtLastError_Positive_Basic") {
|
||||
HIP_CHECK(hipPeekAtLastError());
|
||||
HIP_CHECK_ERROR(hipMalloc(nullptr, 1), hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipPeekAtLastError(), hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidValue);
|
||||
HIP_CHECK(hipPeekAtLastError());
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipPeekAtLastError_Positive_Threaded") {
|
||||
class HipPeekAtLastErrorTest : public ThreadedZigZagTest<HipPeekAtLastErrorTest> {
|
||||
public:
|
||||
void TestPart2() { REQUIRE_THREAD(hipMalloc(nullptr, 1) == hipErrorInvalidValue); }
|
||||
void TestPart3() {
|
||||
HIP_CHECK(hipPeekAtLastError());
|
||||
HIP_CHECK(hipGetLastError());
|
||||
}
|
||||
void TestPart4() { REQUIRE_THREAD(hipPeekAtLastError() == hipErrorInvalidValue); }
|
||||
};
|
||||
|
||||
HipPeekAtLastErrorTest test;
|
||||
test.run();
|
||||
}
|
||||
@@ -6,6 +6,7 @@ set(TEST_SRC
|
||||
Unit_hipEventRecord.cc
|
||||
Unit_hipEventIpc.cc
|
||||
hipEventDestroy.cc
|
||||
hipEventCreate.cc
|
||||
hipEventCreateWithFlags.cc
|
||||
hipEventSynchronize.cc
|
||||
)
|
||||
|
||||
@@ -0,0 +1,38 @@
|
||||
/*
|
||||
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 :
|
||||
Unit_hipEventCreate_Positive - Test simple event creation with hipEventCreate api
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
TEST_CASE("Unit_hipEventCreate_Positive") {
|
||||
int id = GENERATE(range(0, HipTest::getDeviceCount()));
|
||||
HIP_CHECK(hipSetDevice(id));
|
||||
|
||||
hipEvent_t event;
|
||||
HIP_CHECK(hipEventCreate(&event));
|
||||
REQUIRE(event != nullptr);
|
||||
|
||||
HIP_CHECK(hipEventDestroy(event));
|
||||
}
|
||||
@@ -29,8 +29,7 @@ Unit_hipEventCreateWithFlags_Positive - Test simple event creation with hipEvent
|
||||
TEST_CASE("Unit_hipEventCreateWithFlags_Positive") {
|
||||
|
||||
#if HT_AMD
|
||||
// On AMD platform, hipEventInterprocess support is under development. Use of this flag will return an error. Omitted
|
||||
const unsigned int flagUnderTest = GENERATE(hipEventDefault, hipEventBlockingSync, hipEventDisableTiming, hipEventReleaseToDevice, hipEventReleaseToSystem);
|
||||
const unsigned int flagUnderTest = GENERATE(hipEventDefault, hipEventBlockingSync, hipEventDisableTiming, hipEventInterprocess | hipEventDisableTiming, hipEventReleaseToDevice, hipEventReleaseToSystem);
|
||||
#else
|
||||
// On Non-AMD platforms hipEventReleaseToDevice / hipEventReleaseToSystem are not defined
|
||||
const unsigned int flagUnderTest = GENERATE(hipEventDefault, hipEventBlockingSync, hipEventDisableTiming, hipEventInterprocess | hipEventDisableTiming);
|
||||
|
||||
@@ -23,6 +23,7 @@ set(TEST_SRC
|
||||
hipGraphAddEmptyNode.cc
|
||||
hipGraphAddDependencies.cc
|
||||
hipGraphAddEventRecordNode.cc
|
||||
hipGraphAddEventWaitNode.cc
|
||||
hipGraph.cc
|
||||
hipSimpleGraphWithKernel.cc
|
||||
hipGraphAddMemcpyNode.cc
|
||||
@@ -78,6 +79,7 @@ set(TEST_SRC
|
||||
hipGraphNodeGetDependencies.cc
|
||||
hipGraphHostNodeGetParams.cc
|
||||
hipGraphExecChildGraphNodeSetParams.cc
|
||||
hipStreamGetCaptureInfo_v2.cc
|
||||
)
|
||||
|
||||
hip_add_exe_to_target(NAME GraphsTest
|
||||
|
||||
@@ -0,0 +1,303 @@
|
||||
/*
|
||||
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 :
|
||||
1) Simple Scenario: Create an event record node and then create an event
|
||||
wait node using the same event and add it to graph. Instantiate and Launch
|
||||
the Graph. Wait for the graph to complete. The operation must succeed without
|
||||
any failures.
|
||||
2) Create a graph 1 with memcpyh2d, event record node (event A), kernel1
|
||||
and memcpyd2h nodes. Create a graph 2 with Event Wait (event A) , kernel2
|
||||
and memcpyd2h nodes. Instantiate and launch graph1 on stream1 and graph2 on
|
||||
stream2. Wait for both graph1 and graph2 to complete. Validate the result of
|
||||
both graphs.
|
||||
3) Execute graph1 and graph2 in scenario 2 multiple times in a loop
|
||||
(100 times).
|
||||
4) Execute scenario 2 with stream1 = stream2.
|
||||
5) Repeat scenario 2 for different event flags.
|
||||
6) Negative Scenarios
|
||||
- Pass input node parameter as nullptr.
|
||||
- Pass input graph parameter as nullptr.
|
||||
- Pass input dependency parameter as nullptr.
|
||||
- Pass input event parameter as nullptr.
|
||||
- Pass uninitialized input graph parameter.
|
||||
- Pass uninitialized input event parameter.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <hip_test_kernels.hh>
|
||||
|
||||
#define LEN 512
|
||||
|
||||
/**
|
||||
* Scenario 1
|
||||
*/
|
||||
TEST_CASE("Unit_hipGraphAddEventWaitNode_Functional_Simple") {
|
||||
hipGraph_t graph;
|
||||
hipStream_t streamForGraph;
|
||||
hipGraphExec_t graphExec;
|
||||
HIP_CHECK(hipStreamCreate(&streamForGraph));
|
||||
HIP_CHECK(hipGraphCreate(&graph, 0));
|
||||
hipEvent_t event;
|
||||
HIP_CHECK(hipEventCreate(&event));
|
||||
hipGraphNode_t event_rec_node, event_wait_node;
|
||||
// Create a event record node in graph
|
||||
HIP_CHECK(hipGraphAddEventRecordNode(&event_rec_node, graph, nullptr, 0,
|
||||
event));
|
||||
// Create a event wait node in graph
|
||||
HIP_CHECK(hipGraphAddEventWaitNode(&event_wait_node, graph, nullptr, 0,
|
||||
event));
|
||||
HIP_CHECK(hipGraphAddDependencies(graph, &event_rec_node,
|
||||
&event_wait_node, 1));
|
||||
// Instantiate and launch the graph
|
||||
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
|
||||
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
|
||||
HIP_CHECK(hipStreamSynchronize(streamForGraph));
|
||||
HIP_CHECK(hipGraphExecDestroy(graphExec));
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
HIP_CHECK(hipEventDestroy(event));
|
||||
HIP_CHECK(hipStreamDestroy(streamForGraph));
|
||||
}
|
||||
|
||||
/**
|
||||
* Local Function
|
||||
*/
|
||||
static void validate_hipGraphAddEventWaitNode_internodedep(int test,
|
||||
int nstep, unsigned flag = hipEventDefault) {
|
||||
size_t memsize = LEN * sizeof(int);
|
||||
constexpr auto blocksPerCU = 6; // to hide latency
|
||||
constexpr auto threadsPerBlock = 256;
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, LEN);
|
||||
size_t NElem{LEN};
|
||||
hipGraph_t graph1, graph2;
|
||||
hipStream_t streamForGraph1, streamForGraph2;
|
||||
hipGraphExec_t graphExec1, graphExec2;
|
||||
HIP_CHECK(hipStreamCreate(&streamForGraph1));
|
||||
HIP_CHECK(hipGraphCreate(&graph1, 0));
|
||||
HIP_CHECK(hipGraphCreate(&graph2, 0));
|
||||
if (0 == test) {
|
||||
HIP_CHECK(hipStreamCreate(&streamForGraph2));
|
||||
} else if (1 == test) {
|
||||
streamForGraph2 = streamForGraph1;
|
||||
}
|
||||
hipEvent_t event1;
|
||||
HIP_CHECK(hipEventCreateWithFlags(&event1, flag));
|
||||
hipGraphNode_t event_rec_node, event_wait_node;
|
||||
int *inp_h, *inp_d, *out_h_g1, *out_d_g1, *out_h_g2, *out_d_g2;
|
||||
// Allocate host buffers
|
||||
inp_h = reinterpret_cast<int*>(malloc(memsize));
|
||||
REQUIRE(inp_h != nullptr);
|
||||
out_h_g1 = reinterpret_cast<int*>(malloc(memsize));
|
||||
REQUIRE(out_h_g1 != nullptr);
|
||||
out_h_g2 = reinterpret_cast<int*>(malloc(memsize));
|
||||
REQUIRE(out_h_g2 != nullptr);
|
||||
// Allocate device buffers
|
||||
HIP_CHECK(hipMalloc(&inp_d, memsize));
|
||||
HIP_CHECK(hipMalloc(&out_d_g1, memsize));
|
||||
HIP_CHECK(hipMalloc(&out_d_g2, memsize));
|
||||
// Initialize host buffer
|
||||
for (uint32_t i = 0; i < LEN; i++) {
|
||||
inp_h[i] = i;
|
||||
out_h_g1[i] = 0;
|
||||
out_h_g2[i] = 0;
|
||||
}
|
||||
// Graph1 creation ...........
|
||||
// Create event1 record node in graph1
|
||||
HIP_CHECK(hipGraphAddEventRecordNode(&event_rec_node, graph1, nullptr, 0,
|
||||
event1));
|
||||
|
||||
// Create memcpy and kernel nodes for graph1
|
||||
hipGraphNode_t memcpyH2D, memcpyD2H_1, kernelnode_1;
|
||||
hipKernelNodeParams kernelNodeParams1{};
|
||||
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D, graph1, nullptr, 0, inp_d,
|
||||
inp_h, memsize, hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_1, graph1, nullptr, 0,
|
||||
out_h_g1, out_d_g1, memsize, hipMemcpyDeviceToHost));
|
||||
|
||||
void* kernelArgs1[] = {&inp_d, &out_d_g1, reinterpret_cast<void *>(&NElem)};
|
||||
kernelNodeParams1.func =
|
||||
reinterpret_cast<void *>(HipTest::vector_square<int>);
|
||||
kernelNodeParams1.gridDim = dim3(blocks);
|
||||
kernelNodeParams1.blockDim = dim3(threadsPerBlock);
|
||||
kernelNodeParams1.sharedMemBytes = 0;
|
||||
kernelNodeParams1.kernelParams = reinterpret_cast<void**>(kernelArgs1);
|
||||
kernelNodeParams1.extra = nullptr;
|
||||
HIP_CHECK(hipGraphAddKernelNode(&kernelnode_1, graph1, nullptr, 0,
|
||||
&kernelNodeParams1));
|
||||
// Create dependencies for graph1
|
||||
HIP_CHECK(hipGraphAddDependencies(graph1, &memcpyH2D,
|
||||
&event_rec_node, 1));
|
||||
HIP_CHECK(hipGraphAddDependencies(graph1, &event_rec_node,
|
||||
&kernelnode_1, 1));
|
||||
HIP_CHECK(hipGraphAddDependencies(graph1, &kernelnode_1,
|
||||
&memcpyD2H_1, 1));
|
||||
|
||||
// Graph2 creation ...........
|
||||
// Create event1 record node in graph2
|
||||
HIP_CHECK(hipGraphAddEventWaitNode(&event_wait_node, graph2, nullptr, 0,
|
||||
event1));
|
||||
|
||||
// Create memcpy and kernel nodes for graph2
|
||||
hipGraphNode_t memcpyD2H_2, kernelnode_2;
|
||||
hipKernelNodeParams kernelNodeParams2{};
|
||||
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_2, graph2, nullptr, 0,
|
||||
out_h_g2, out_d_g2, memsize, hipMemcpyDeviceToHost));
|
||||
|
||||
void* kernelArgs2[] = {&inp_d, &out_d_g2, reinterpret_cast<void *>(&NElem)};
|
||||
kernelNodeParams2.func =
|
||||
reinterpret_cast<void *>(HipTest::vector_cubic<int>);
|
||||
kernelNodeParams2.gridDim = dim3(blocks);
|
||||
kernelNodeParams2.blockDim = dim3(threadsPerBlock);
|
||||
kernelNodeParams2.sharedMemBytes = 0;
|
||||
kernelNodeParams2.kernelParams = reinterpret_cast<void**>(kernelArgs2);
|
||||
kernelNodeParams2.extra = nullptr;
|
||||
HIP_CHECK(hipGraphAddKernelNode(&kernelnode_2, graph2, nullptr, 0,
|
||||
&kernelNodeParams2));
|
||||
// Create dependencies for graph2
|
||||
HIP_CHECK(hipGraphAddDependencies(graph2, &event_wait_node,
|
||||
&kernelnode_2, 1));
|
||||
HIP_CHECK(hipGraphAddDependencies(graph2, &kernelnode_2,
|
||||
&memcpyD2H_2, 1));
|
||||
|
||||
// Instantiate and launch the graphs
|
||||
HIP_CHECK(hipGraphInstantiate(&graphExec1, graph1, nullptr, nullptr, 0));
|
||||
HIP_CHECK(hipGraphInstantiate(&graphExec2, graph2, nullptr, nullptr, 0));
|
||||
for (int istep = 0; istep < nstep; istep++) {
|
||||
HIP_CHECK(hipGraphLaunch(graphExec1, streamForGraph1));
|
||||
HIP_CHECK(hipGraphLaunch(graphExec2, streamForGraph2));
|
||||
HIP_CHECK(hipStreamSynchronize(streamForGraph1));
|
||||
HIP_CHECK(hipStreamSynchronize(streamForGraph2));
|
||||
// Validate output
|
||||
bool btestPassed1 = true;
|
||||
for (uint32_t i = 0; i < LEN; i++) {
|
||||
if (out_h_g1[i] != (inp_h[i]*inp_h[i])) {
|
||||
btestPassed1 = false;
|
||||
break;
|
||||
}
|
||||
}
|
||||
REQUIRE(btestPassed1 == true);
|
||||
bool btestPassed2 = true;
|
||||
for (uint32_t i = 0; i < LEN; i++) {
|
||||
if (out_h_g2[i] != (inp_h[i]*inp_h[i]*inp_h[i])) {
|
||||
btestPassed2 = false;
|
||||
break;
|
||||
}
|
||||
}
|
||||
REQUIRE(btestPassed2 == true);
|
||||
}
|
||||
// Destroy all resources
|
||||
HIP_CHECK(hipFree(inp_d));
|
||||
HIP_CHECK(hipFree(out_d_g1));
|
||||
HIP_CHECK(hipFree(out_d_g2));
|
||||
free(inp_h);
|
||||
free(out_h_g1);
|
||||
free(out_h_g2);
|
||||
HIP_CHECK(hipGraphExecDestroy(graphExec1));
|
||||
HIP_CHECK(hipGraphExecDestroy(graphExec2));
|
||||
HIP_CHECK(hipGraphDestroy(graph1));
|
||||
HIP_CHECK(hipGraphDestroy(graph2));
|
||||
HIP_CHECK(hipEventDestroy(event1));
|
||||
HIP_CHECK(hipStreamDestroy(streamForGraph1));
|
||||
if (0 == test) {
|
||||
HIP_CHECK(hipStreamDestroy(streamForGraph2));
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Scenario 2
|
||||
*/
|
||||
TEST_CASE("Unit_hipGraphAddEventWaitNode_MultGraphMultStrmDependency") {
|
||||
validate_hipGraphAddEventWaitNode_internodedep(0, 1);
|
||||
}
|
||||
|
||||
/**
|
||||
* Scenario 3
|
||||
*/
|
||||
TEST_CASE("Unit_hipGraphAddEventWaitNode_MultipleRun") {
|
||||
validate_hipGraphAddEventWaitNode_internodedep(0, 100);
|
||||
}
|
||||
|
||||
/**
|
||||
* Scenario 4
|
||||
*/
|
||||
TEST_CASE("Unit_hipGraphAddEventWaitNode_MultGraphOneStrmDependency") {
|
||||
validate_hipGraphAddEventWaitNode_internodedep(1, 1);
|
||||
}
|
||||
|
||||
/**
|
||||
* Scenario 5
|
||||
*/
|
||||
TEST_CASE("Unit_hipGraphAddEventWaitNode_differentFlags") {
|
||||
SECTION("flag = hipEventBlockingSync") {
|
||||
validate_hipGraphAddEventWaitNode_internodedep(0, 1,
|
||||
hipEventBlockingSync);
|
||||
}
|
||||
SECTION("graph = hipEventDisableTiming") {
|
||||
validate_hipGraphAddEventWaitNode_internodedep(0, 1,
|
||||
hipEventDisableTiming);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Scenario 6
|
||||
*/
|
||||
TEST_CASE("Unit_hipGraphAddEventWaitNode_Negative") {
|
||||
hipGraph_t graph;
|
||||
HIP_CHECK(hipGraphCreate(&graph, 0));
|
||||
hipEvent_t event;
|
||||
HIP_CHECK(hipEventCreate(&event));
|
||||
hipGraphNode_t eventwait;
|
||||
|
||||
SECTION("pGraphNode = nullptr") {
|
||||
REQUIRE(hipErrorInvalidValue == hipGraphAddEventWaitNode(nullptr,
|
||||
graph, nullptr, 0, event));
|
||||
}
|
||||
|
||||
SECTION("graph = nullptr") {
|
||||
REQUIRE(hipErrorInvalidValue == hipGraphAddEventWaitNode(&eventwait,
|
||||
nullptr, nullptr, 0, event));
|
||||
}
|
||||
|
||||
SECTION("pDependencies = nullptr") {
|
||||
REQUIRE(hipErrorInvalidValue == hipGraphAddEventWaitNode(&eventwait,
|
||||
graph, nullptr, 1, event));
|
||||
}
|
||||
|
||||
SECTION("event = nullptr") {
|
||||
REQUIRE(hipErrorInvalidValue == hipGraphAddEventWaitNode(&eventwait,
|
||||
graph, nullptr, 0, nullptr));
|
||||
}
|
||||
|
||||
SECTION("graph is uninitialized") {
|
||||
hipGraph_t graph_uninit{};
|
||||
REQUIRE(hipErrorInvalidValue == hipGraphAddEventWaitNode(&eventwait,
|
||||
graph_uninit, nullptr, 0, event));
|
||||
}
|
||||
|
||||
SECTION("event is uninitialized") {
|
||||
hipEvent_t event_uninit{};
|
||||
REQUIRE(hipErrorInvalidValue == hipGraphAddEventWaitNode(&eventwait,
|
||||
graph, nullptr, 0, event_uninit));
|
||||
}
|
||||
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
HIP_CHECK(hipEventDestroy(event));
|
||||
}
|
||||
@@ -25,6 +25,8 @@ Functional ::
|
||||
2) Pass nodes as nullptr and verify numNodes returns actual number of nodes added to graph.
|
||||
3) If numNodes passed is greater than the actual number of nodes, the remaining entries in nodes
|
||||
will be set to NULL, and the number of nodes actually obtained will be returned in numNodes.
|
||||
4) Begin stream capture and push operations to stream. Verify nodes of created graph are matching the
|
||||
operations pushed.
|
||||
|
||||
Argument Validation ::
|
||||
1) Pass graph as nullptr and verify api returns error code.
|
||||
@@ -139,6 +141,110 @@ TEST_CASE("Unit_hipGraphGetNodes_Functional") {
|
||||
free(nodes);
|
||||
}
|
||||
|
||||
/**
|
||||
* Begin stream capture and push operations to stream.
|
||||
* Verify nodes of created graph are matching the operations pushed.
|
||||
*/
|
||||
TEST_CASE("Unit_hipGraphGetNodes_CapturedStream") {
|
||||
hipGraph_t graph{nullptr};
|
||||
hipGraphExec_t graphExec{nullptr};
|
||||
constexpr unsigned blocks = 512;
|
||||
constexpr unsigned threadsPerBlock = 256;
|
||||
constexpr size_t N = 1000000;
|
||||
size_t Nbytes = N * sizeof(float);
|
||||
constexpr int numMemcpy{2}, numKernel{1}, numMemset{1};
|
||||
int cntMemcpy{}, cntKernel{}, cntMemset{};
|
||||
hipStream_t stream, streamForGraph;
|
||||
hipGraphNodeType nodeType;
|
||||
float *A_d, *C_d;
|
||||
float *A_h, *C_h;
|
||||
|
||||
A_h = reinterpret_cast<float*>(malloc(Nbytes));
|
||||
C_h = reinterpret_cast<float*>(malloc(Nbytes));
|
||||
REQUIRE(A_h != nullptr);
|
||||
REQUIRE(C_h != nullptr);
|
||||
HIP_CHECK(hipMalloc(&A_d, Nbytes));
|
||||
HIP_CHECK(hipMalloc(&C_d, Nbytes));
|
||||
REQUIRE(A_d != nullptr);
|
||||
REQUIRE(C_d != nullptr);
|
||||
|
||||
HIP_CHECK(hipStreamCreate(&streamForGraph));
|
||||
// Initialize input buffer
|
||||
for (size_t i = 0; i < N; ++i) {
|
||||
A_h[i] = 3.146f + i; // Pi
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
|
||||
HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
|
||||
HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream));
|
||||
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
|
||||
dim3(threadsPerBlock), 0, stream, A_d, C_d, N);
|
||||
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream));
|
||||
HIP_CHECK(hipStreamEndCapture(stream, &graph));
|
||||
REQUIRE(graph != nullptr);
|
||||
|
||||
size_t numNodes{};
|
||||
HIP_CHECK(hipGraphGetNodes(graph, nullptr, &numNodes));
|
||||
INFO("Num of nodes returned by GetNodes : " << numNodes);
|
||||
REQUIRE(numNodes == numMemcpy + numKernel + numMemset);
|
||||
|
||||
int numBytes = sizeof(hipGraphNode_t) * numNodes;
|
||||
hipGraphNode_t* nodes = reinterpret_cast<hipGraphNode_t *>(malloc(numBytes));
|
||||
REQUIRE(nodes != nullptr);
|
||||
|
||||
HIP_CHECK(hipGraphGetNodes(graph, nodes, &numNodes));
|
||||
for (size_t i = 0; i < numNodes; i++) {
|
||||
HIP_CHECK(hipGraphNodeGetType(nodes[i], &nodeType));
|
||||
|
||||
switch (nodeType) {
|
||||
case hipGraphNodeTypeMemcpy:
|
||||
cntMemcpy++;
|
||||
break;
|
||||
|
||||
case hipGraphNodeTypeKernel:
|
||||
cntKernel++;
|
||||
break;
|
||||
|
||||
case hipGraphNodeTypeMemset:
|
||||
cntMemset++;
|
||||
break;
|
||||
|
||||
default:
|
||||
INFO("Unexpected nodetype returned : " << nodeType);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
|
||||
REQUIRE(cntMemcpy == numMemcpy);
|
||||
REQUIRE(cntKernel == numKernel);
|
||||
REQUIRE(cntMemset == numMemset);
|
||||
|
||||
// Instantiate and launch the graph
|
||||
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
|
||||
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
|
||||
HIP_CHECK(hipStreamSynchronize(streamForGraph));
|
||||
|
||||
// Validate the computation
|
||||
for (size_t i = 0; i < N; i++) {
|
||||
if (C_h[i] != A_h[i] * A_h[i]) {
|
||||
INFO("A and C not matching at " << i << " C_h[i] " << C_h[i]
|
||||
<< " A_h[i] " << A_h[i]);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamDestroy(streamForGraph));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
HIP_CHECK(hipGraphExecDestroy(graphExec));
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
free(A_h);
|
||||
free(C_h);
|
||||
free(nodes);
|
||||
HIP_CHECK(hipFree(A_d));
|
||||
HIP_CHECK(hipFree(C_d));
|
||||
}
|
||||
|
||||
/**
|
||||
* Test performs api parameter validation by passing various values
|
||||
* as input and output parameters and validates the behavior.
|
||||
|
||||
@@ -26,6 +26,8 @@ Functional ::
|
||||
2) Pass nodes as nullptr and verify api returns actual number of root nodes added to graph.
|
||||
3) If NumRootNodes passed is greater than the actual number of root nodes, the remaining entries in
|
||||
nodes list will be set to NULL, and the number of nodes actually obtained will be returned in NumRootNodes.
|
||||
4) Create a graph with stream capture done on multiple dependent streams.
|
||||
Verify root nodes of created graph are matching the operations pushed which doesn't have dependencies.
|
||||
|
||||
Argument Validation ::
|
||||
1) Pass graph as nullptr and verify api returns error code.
|
||||
@@ -50,6 +52,7 @@ TEST_CASE("Unit_hipGraphGetRootNodes_Functional") {
|
||||
constexpr auto addlEntries = 5;
|
||||
hipGraph_t graph;
|
||||
|
||||
|
||||
hipGraphNode_t memcpyNode, kernelNode;
|
||||
hipKernelNodeParams kernelNodeParams{};
|
||||
hipStream_t streamForGraph;
|
||||
@@ -141,6 +144,112 @@ TEST_CASE("Unit_hipGraphGetRootNodes_Functional") {
|
||||
free(rootnodes);
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* Create a graph with stream capture done on multiple dependent streams. Verify root nodes
|
||||
* of created graph are matching the operations pushed which doesn't have dependencies.
|
||||
*/
|
||||
TEST_CASE("Unit_hipGraphGetRootNodes_CapturedStream") {
|
||||
hipStream_t stream1{nullptr}, stream2{nullptr}, mstream{nullptr};
|
||||
hipStream_t streamForGraph{nullptr};
|
||||
hipEvent_t memsetEvent1, memsetEvent2, forkStreamEvent;
|
||||
hipGraph_t graph{nullptr};
|
||||
hipGraphExec_t graphExec{nullptr};
|
||||
constexpr unsigned blocks = 512;
|
||||
constexpr unsigned threadsPerBlock = 256;
|
||||
constexpr size_t N = 1000000;
|
||||
constexpr int numMemsetNodes = 2;
|
||||
size_t Nbytes = N * sizeof(float), numRootNodes{};
|
||||
float *A_d, *C_d;
|
||||
float *A_h, *C_h;
|
||||
A_h = reinterpret_cast<float*>(malloc(Nbytes));
|
||||
C_h = reinterpret_cast<float*>(malloc(Nbytes));
|
||||
REQUIRE(A_h != nullptr);
|
||||
REQUIRE(C_h != nullptr);
|
||||
HIP_CHECK(hipMalloc(&A_d, Nbytes));
|
||||
HIP_CHECK(hipMalloc(&C_d, Nbytes));
|
||||
REQUIRE(A_d != nullptr);
|
||||
REQUIRE(C_d != nullptr);
|
||||
|
||||
HIP_CHECK(hipStreamCreate(&streamForGraph));
|
||||
|
||||
// Initialize input buffer
|
||||
for (size_t i = 0; i < N; ++i) {
|
||||
A_h[i] = 3.146f + i; // Pi
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamCreate(&stream1));
|
||||
HIP_CHECK(hipStreamCreate(&stream2));
|
||||
HIP_CHECK(hipStreamCreate(&mstream));
|
||||
HIP_CHECK(hipEventCreate(&memsetEvent1));
|
||||
HIP_CHECK(hipEventCreate(&memsetEvent2));
|
||||
HIP_CHECK(hipEventCreate(&forkStreamEvent));
|
||||
HIP_CHECK(hipStreamBeginCapture(mstream, hipStreamCaptureModeGlobal));
|
||||
HIP_CHECK(hipEventRecord(forkStreamEvent, mstream));
|
||||
HIP_CHECK(hipStreamWaitEvent(stream1, forkStreamEvent, 0));
|
||||
HIP_CHECK(hipStreamWaitEvent(stream2, forkStreamEvent, 0));
|
||||
HIP_CHECK(hipMemsetAsync(A_d, 0, Nbytes, stream1));
|
||||
HIP_CHECK(hipEventRecord(memsetEvent1, stream1));
|
||||
HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream2));
|
||||
HIP_CHECK(hipEventRecord(memsetEvent2, stream2));
|
||||
HIP_CHECK(hipStreamWaitEvent(mstream, memsetEvent1, 0));
|
||||
HIP_CHECK(hipStreamWaitEvent(mstream, memsetEvent2, 0));
|
||||
HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mstream));
|
||||
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
|
||||
dim3(threadsPerBlock), 0, mstream, A_d, C_d, N);
|
||||
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mstream));
|
||||
HIP_CHECK(hipStreamEndCapture(mstream, &graph));
|
||||
|
||||
// Verify numof root nodes
|
||||
HIP_CHECK(hipGraphGetRootNodes(graph, nullptr, &numRootNodes));
|
||||
REQUIRE(numRootNodes == numMemsetNodes);
|
||||
INFO("Num of nodes returned by GetRootNodes : " << numRootNodes);
|
||||
|
||||
int numBytes = sizeof(hipGraphNode_t) * numRootNodes;
|
||||
hipGraphNode_t* nodes = reinterpret_cast<hipGraphNode_t *>(malloc(numBytes));
|
||||
REQUIRE(nodes != nullptr);
|
||||
|
||||
hipGraphNodeType nodeType;
|
||||
HIP_CHECK(hipGraphGetRootNodes(graph, nodes, &numRootNodes));
|
||||
REQUIRE(numRootNodes == numMemsetNodes);
|
||||
|
||||
// Verify root nodes returned are memset nodes.
|
||||
HIP_CHECK(hipGraphNodeGetType(nodes[0], &nodeType));
|
||||
REQUIRE(nodeType == hipGraphNodeTypeMemset);
|
||||
HIP_CHECK(hipGraphNodeGetType(nodes[1], &nodeType));
|
||||
REQUIRE(nodeType == hipGraphNodeTypeMemset);
|
||||
|
||||
// Instantiate and launch the graph
|
||||
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
|
||||
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
|
||||
HIP_CHECK(hipStreamSynchronize(streamForGraph));
|
||||
|
||||
// Validate the computation
|
||||
for (size_t i = 0; i < N; i++) {
|
||||
if (C_h[i] != A_h[i] * A_h[i]) {
|
||||
INFO("A and C not matching at " << i << " C_h[i] " << C_h[i]
|
||||
<< " A_h[i] " << A_h[i]);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipGraphExecDestroy(graphExec));
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
HIP_CHECK(hipStreamDestroy(streamForGraph));
|
||||
HIP_CHECK(hipStreamDestroy(mstream));
|
||||
HIP_CHECK(hipStreamDestroy(stream1));
|
||||
HIP_CHECK(hipStreamDestroy(stream2));
|
||||
HIP_CHECK(hipEventDestroy(forkStreamEvent));
|
||||
HIP_CHECK(hipEventDestroy(memsetEvent1));
|
||||
HIP_CHECK(hipEventDestroy(memsetEvent2));
|
||||
free(A_h);
|
||||
free(C_h);
|
||||
free(nodes);
|
||||
HIP_CHECK(hipFree(A_d));
|
||||
HIP_CHECK(hipFree(C_d));
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* Test performs api parameter validation by passing various values
|
||||
* as input and output parameters and validates the behavior.
|
||||
|
||||
@@ -0,0 +1,280 @@
|
||||
/*
|
||||
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
|
||||
------------------
|
||||
Functional:
|
||||
1) Start stream capture and get capture info v2. Verify api is success, capture status is hipStreamCaptureStatusActive,
|
||||
identifier returned is valid/non-zero, graph object is returned.
|
||||
2) When stream capture is in progress, create dependent nodes by creating multistream dependencies and verify the api returns
|
||||
valid dependent nodes.
|
||||
3) End stream capture and get capture info. Verify api is success, capture status is hipStreamCaptureStatusNone and
|
||||
identifier/graph/nodes are not returned by api.
|
||||
4) When optional parameters are not passed, make sure api still returns capture status of stream.
|
||||
5) Begin capture on hipStreamPerThread, get capture info v2 and validate results.
|
||||
6) Perform multiple captures and verify the identifier returned is unique.
|
||||
|
||||
Parameter Validation/Negative:
|
||||
1) Capture Status location as nullptr and verify api returns error code.
|
||||
2) Stream as nullptr and verify api returns error code.
|
||||
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <hip_test_kernels.hh>
|
||||
|
||||
constexpr size_t N = 1000000;
|
||||
constexpr int LAUNCH_ITERS = 1;
|
||||
|
||||
/**
|
||||
* Validates stream capture infov2, launches graph and verifies results
|
||||
*/
|
||||
void validateStreamCaptureInfoV2(hipStream_t mstream) {
|
||||
hipStream_t stream1{nullptr}, stream2{nullptr}, streamForLaunch{nullptr};
|
||||
hipEvent_t memcpyEvent1, memsetEvent2, forkStreamEvent;
|
||||
hipGraph_t graph{nullptr}, capInfoGraph{nullptr};
|
||||
hipGraphExec_t graphExec{nullptr};
|
||||
constexpr unsigned blocks = 512;
|
||||
constexpr unsigned threadsPerBlock = 256;
|
||||
const hipGraphNode_t *nodelist{};
|
||||
size_t Nbytes = N * sizeof(float), numDependencies;
|
||||
float *A_d, *C_d;
|
||||
float *A_h, *C_h;
|
||||
A_h = reinterpret_cast<float*>(malloc(Nbytes));
|
||||
C_h = reinterpret_cast<float*>(malloc(Nbytes));
|
||||
REQUIRE(A_h != nullptr);
|
||||
REQUIRE(C_h != nullptr);
|
||||
HIP_CHECK(hipMalloc(&A_d, Nbytes));
|
||||
HIP_CHECK(hipMalloc(&C_d, Nbytes));
|
||||
REQUIRE(A_d != nullptr);
|
||||
REQUIRE(C_d != nullptr);
|
||||
HIP_CHECK(hipStreamCreate(&streamForLaunch));
|
||||
|
||||
// Initialize input buffer
|
||||
for (size_t i = 0; i < N; ++i) {
|
||||
A_h[i] = 3.146f + i; // Pi
|
||||
}
|
||||
|
||||
// Create cross stream dependencies.
|
||||
// memset/memcpy operations are done on stream1 and stream2
|
||||
// and they are joined back to mainstream
|
||||
HIP_CHECK(hipStreamCreate(&stream1));
|
||||
HIP_CHECK(hipStreamCreate(&stream2));
|
||||
HIP_CHECK(hipEventCreate(&memcpyEvent1));
|
||||
HIP_CHECK(hipEventCreate(&memsetEvent2));
|
||||
HIP_CHECK(hipEventCreate(&forkStreamEvent));
|
||||
|
||||
HIP_CHECK(hipStreamBeginCapture(mstream, hipStreamCaptureModeGlobal));
|
||||
HIP_CHECK(hipEventRecord(forkStreamEvent, mstream));
|
||||
HIP_CHECK(hipStreamWaitEvent(stream1, forkStreamEvent, 0));
|
||||
HIP_CHECK(hipStreamWaitEvent(stream2, forkStreamEvent, 0));
|
||||
HIP_CHECK(hipMemsetAsync(A_d, 0, Nbytes, stream1));
|
||||
HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream1));
|
||||
HIP_CHECK(hipEventRecord(memcpyEvent1, stream1));
|
||||
HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream2));
|
||||
HIP_CHECK(hipEventRecord(memsetEvent2, stream2));
|
||||
HIP_CHECK(hipStreamWaitEvent(mstream, memcpyEvent1, 0));
|
||||
HIP_CHECK(hipStreamWaitEvent(mstream, memsetEvent2, 0));
|
||||
|
||||
hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone};
|
||||
unsigned long long capSequenceID = 0; // NOLINT
|
||||
constexpr int numDepsCreated = 2; // Num of dependencies created
|
||||
|
||||
HIP_CHECK(hipStreamGetCaptureInfo_v2(mstream, &captureStatus,
|
||||
&capSequenceID, &capInfoGraph, &nodelist, &numDependencies));
|
||||
|
||||
// verify capture status is active, sequence id is valid, graph is returned,
|
||||
REQUIRE(captureStatus == hipStreamCaptureStatusActive);
|
||||
REQUIRE(capSequenceID > 0);
|
||||
REQUIRE(capInfoGraph != nullptr);
|
||||
REQUIRE(numDependencies == numDepsCreated);
|
||||
|
||||
// verify dependency nodes list returned is the one we created.
|
||||
hipGraphNodeType type(hipGraphNodeTypeEmpty);
|
||||
|
||||
HIP_CHECK(hipGraphNodeGetType(nodelist[0], &type));
|
||||
if ((type != hipGraphNodeTypeMemset) && (type != hipGraphNodeTypeMemcpy)) {
|
||||
INFO("Type0 returned as " << type);
|
||||
REQUIRE(false);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipGraphNodeGetType(nodelist[1], &type));
|
||||
if ((type != hipGraphNodeTypeMemset) && (type != hipGraphNodeTypeMemcpy)) {
|
||||
INFO("Type1 returned as " << type);
|
||||
REQUIRE(false);
|
||||
}
|
||||
|
||||
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
|
||||
dim3(threadsPerBlock), 0, mstream, A_d, C_d, N);
|
||||
|
||||
// End capture and verify graph is returned
|
||||
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mstream));
|
||||
HIP_CHECK(hipStreamEndCapture(mstream, &graph));
|
||||
REQUIRE(graph != nullptr);
|
||||
|
||||
// verify capture status is inactive and other params are not updated
|
||||
capSequenceID = 0;
|
||||
capInfoGraph = nullptr;
|
||||
numDependencies = 0;
|
||||
nodelist = nullptr;
|
||||
HIP_CHECK(hipStreamGetCaptureInfo_v2(mstream, &captureStatus,
|
||||
&capSequenceID, &capInfoGraph, &nodelist, &numDependencies));
|
||||
REQUIRE(captureStatus == hipStreamCaptureStatusNone);
|
||||
REQUIRE(capSequenceID == 0);
|
||||
REQUIRE(capInfoGraph == nullptr);
|
||||
REQUIRE(nodelist == nullptr);
|
||||
REQUIRE(numDependencies == 0);
|
||||
|
||||
// Verify api still returns capture status when optional args are not passed
|
||||
HIP_CHECK(hipStreamGetCaptureInfo_v2(mstream, &captureStatus));
|
||||
REQUIRE(captureStatus == hipStreamCaptureStatusNone);
|
||||
|
||||
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
|
||||
REQUIRE(graphExec != nullptr);
|
||||
|
||||
// Replay the recorded sequence multiple times
|
||||
for (int i = 0; i < LAUNCH_ITERS; i++) {
|
||||
HIP_CHECK(hipGraphLaunch(graphExec, streamForLaunch));
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamSynchronize(streamForLaunch));
|
||||
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
HIP_CHECK(hipStreamDestroy(streamForLaunch));
|
||||
HIP_CHECK(hipStreamDestroy(stream1));
|
||||
HIP_CHECK(hipStreamDestroy(stream2));
|
||||
HIP_CHECK(hipEventDestroy(forkStreamEvent));
|
||||
HIP_CHECK(hipEventDestroy(memcpyEvent1));
|
||||
HIP_CHECK(hipEventDestroy(memsetEvent2));
|
||||
HIP_CHECK(hipFree(A_d));
|
||||
HIP_CHECK(hipFree(C_d));
|
||||
|
||||
// Validate the computation
|
||||
for (size_t i = 0; i < N; i++) {
|
||||
if (C_h[i] != A_h[i] * A_h[i]) {
|
||||
INFO("A and C not matching at " << i << " C_h[i] " << C_h[i]
|
||||
<< " A_h[i] " << A_h[i]);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
free(A_h);
|
||||
free(C_h);
|
||||
}
|
||||
|
||||
/**
|
||||
* Basic Functional Test for stream capture and getting capture info V2.
|
||||
* Regular/custom stream is used for stream capture.
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamGetCaptureInfo_v2_BasicFunctional") {
|
||||
hipStream_t streamForCapture;
|
||||
|
||||
HIP_CHECK(hipStreamCreate(&streamForCapture));
|
||||
validateStreamCaptureInfoV2(streamForCapture);
|
||||
HIP_CHECK(hipStreamDestroy(streamForCapture));
|
||||
}
|
||||
|
||||
/**
|
||||
* Test performs stream capture on hipStreamPerThread and validates
|
||||
* capture info V2.
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamGetCaptureInfo_v2_hipStreamPerThread") {
|
||||
validateStreamCaptureInfoV2(hipStreamPerThread);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test starts stream capture on multiple streams and verifies uniqueness of
|
||||
* identifiers returned from capture Info V2.
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamGetCaptureInfo_v2_UniqueID") {
|
||||
constexpr int numStreams = 100;
|
||||
hipStream_t streams[numStreams]{};
|
||||
hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone};
|
||||
std::vector<int> idlist;
|
||||
unsigned long long capSequenceID{}; // NOLINT
|
||||
hipGraph_t graph{nullptr};
|
||||
|
||||
for (int i = 0; i < numStreams; i++) {
|
||||
HIP_CHECK(hipStreamCreate(&streams[i]));
|
||||
HIP_CHECK(hipStreamBeginCapture(streams[i], hipStreamCaptureModeGlobal));
|
||||
HIP_CHECK(hipStreamGetCaptureInfo_v2(streams[i], &captureStatus,
|
||||
&capSequenceID, nullptr, nullptr, nullptr));
|
||||
REQUIRE(captureStatus == hipStreamCaptureStatusActive);
|
||||
REQUIRE(capSequenceID > 0);
|
||||
idlist.push_back(capSequenceID);
|
||||
}
|
||||
|
||||
for (int i = 0; i < numStreams; i++) {
|
||||
for (int j = i+1; j < numStreams; j++) {
|
||||
if (idlist[i] == idlist[j]) {
|
||||
INFO("Same identifier returned for stream "
|
||||
<< i << " and stream " << j);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < numStreams; i++) {
|
||||
HIP_CHECK(hipStreamEndCapture(streams[i], &graph));
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
HIP_CHECK(hipStreamDestroy(streams[i]));
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Parameter validation/Negative tests for api
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamGetCaptureInfo_v2_ParamValidation") {
|
||||
hipError_t ret;
|
||||
hipStream_t stream;
|
||||
float *A_d;
|
||||
hipGraph_t graph{}, capInfoGraph{};
|
||||
hipStreamCaptureStatus captureStatus;
|
||||
unsigned long long capSequenceID; // NOLINT
|
||||
size_t numDependencies;
|
||||
const hipGraphNode_t *nodelist{};
|
||||
constexpr int numBytes{100};
|
||||
|
||||
HIP_CHECK(hipMalloc(&A_d, numBytes));
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
|
||||
HIP_CHECK(hipMemsetAsync(A_d, 0, numBytes, stream));
|
||||
|
||||
SECTION("Capture Status location as nullptr") {
|
||||
ret = hipStreamGetCaptureInfo_v2(stream, nullptr,
|
||||
&capSequenceID, &capInfoGraph, &nodelist, &numDependencies);
|
||||
REQUIRE(ret == hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("Stream as nullptr") {
|
||||
ret = hipStreamGetCaptureInfo_v2(nullptr, &captureStatus,
|
||||
&capSequenceID, &capInfoGraph, &nodelist, &numDependencies);
|
||||
if ((ret != hipErrorUnknown) && (ret != hipErrorStreamCaptureImplicit)) {
|
||||
INFO("Ret : " << ret);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamEndCapture(stream, &graph));
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
HIP_CHECK(hipFree(A_d));
|
||||
}
|
||||
@@ -676,8 +676,14 @@ TEST_CASE("Unit_hipMemAdvise_TstAlignedAllocMem") {
|
||||
int stat = 0;
|
||||
if (fork() == 0) {
|
||||
// The below part should be inside fork
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
int managedMem = 0, pageMemAccess = 0;
|
||||
HIP_CHECK(hipDeviceGetAttribute(&pageMemAccess,
|
||||
hipDeviceAttributePageableMemoryAccess, 0));
|
||||
WARN("hipDeviceAttributePageableMemoryAccess:" << pageMemAccess);
|
||||
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managedMem, hipDeviceAttributeManagedMemory, 0));
|
||||
WARN("hipDeviceAttributeManagedMemory: " << managedMem);
|
||||
if ((managedMem == 1) && (pageMemAccess == 1)) {
|
||||
int *Mllc = nullptr, MemSz = 4096 * 4, NumElms = 4096, InitVal = 123;
|
||||
// Mllc = reinterpret_cast<(int *)>(aligned_alloc(4096, MemSz));
|
||||
Mllc = reinterpret_cast<int*>(aligned_alloc(4096, 4096*4));
|
||||
@@ -706,19 +712,27 @@ TEST_CASE("Unit_hipMemAdvise_TstAlignedAllocMem") {
|
||||
exit(10); // 10 for Pass result
|
||||
}
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
SUCCEED("GPU 0 doesn't support ManagedMemory with hipDeviceAttributePageableMemoryAccess "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
exit(Catch::ResultDisposition::ContinueOnFailure);
|
||||
}
|
||||
} else {
|
||||
wait(&stat);
|
||||
int Result = WEXITSTATUS(stat);
|
||||
if (Result != 10) {
|
||||
REQUIRE(false);
|
||||
if (Result == Catch::ResultDisposition::ContinueOnFailure) {
|
||||
WARN("GPU 0 doesn't support ManagedMemory with hipDeviceAttributePageableMemoryAccess "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
} else {
|
||||
if (Result != 10) {
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
} else {
|
||||
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
|
||||
"skipping the testcase for this GPU " << device);
|
||||
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
|
||||
"skipping the testcase for this GPU " << device);
|
||||
WARN("Memory model feature is only supported for gfx90a, Hence"
|
||||
"skipping the testcase for this GPU " << device);
|
||||
}
|
||||
|
||||
}
|
||||
|
||||
@@ -16,7 +16,13 @@ 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 :
|
||||
Unit_hipStreamWaitEvent_Negative - Test unsuccessful hipStreamWaitEvent when either event or flags are invalid
|
||||
Unit_hipStreamWaitEvent_UninitializedStream_Negative - Test unsuccessful hipStreamWaitEvent when stream is uninitialized
|
||||
Unit_hipStreamWaitEvent_Default - Test simple waiting for an event with hipStreamWaitEvent api
|
||||
Unit_hipStreamWaitEvent_DifferentStreams - Test waiting for an event on a different stream with hipStreamWaitEvent api
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
@@ -59,6 +65,20 @@ TEST_CASE("Unit_hipStreamWaitEvent_Negative") {
|
||||
}
|
||||
}
|
||||
|
||||
/* Test removed for Nvidia devices because it returns unexpected error */
|
||||
#if !HT_NVIDIA
|
||||
TEST_CASE("Unit_hipStreamWaitEvent_UninitializedStream_Negative") {
|
||||
hipStream_t stream{reinterpret_cast<hipStream_t>(0xFFFF)};
|
||||
hipEvent_t event{nullptr};
|
||||
|
||||
HIP_CHECK(hipEventCreate(&event));
|
||||
|
||||
HIP_CHECK_ERROR(hipStreamWaitEvent(stream, event, 0), hipErrorContextIsDestroyed);
|
||||
|
||||
HIP_CHECK(hipEventDestroy(event));
|
||||
}
|
||||
#endif
|
||||
|
||||
// Since we can not use atomic*_system on every gpu, we will use wait based on clock rate.
|
||||
// This wont be accurate since clock rate of a GPU varies depending on many variables including
|
||||
// thermals, load, utilization
|
||||
@@ -74,6 +94,38 @@ __global__ void waitKernel(int clockRate, int seconds) {
|
||||
}
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipStreamWaitEvent_Default") {
|
||||
hipStream_t stream{nullptr};
|
||||
hipEvent_t waitEvent{nullptr};
|
||||
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipEventCreate(&waitEvent));
|
||||
|
||||
REQUIRE(stream != nullptr);
|
||||
REQUIRE(waitEvent != nullptr);
|
||||
|
||||
int deviceId {};
|
||||
HIP_CHECK(hipGetDevice(&deviceId));
|
||||
|
||||
hipDeviceProp_t prop{};
|
||||
HIP_CHECK(hipGetDeviceProperties(&prop, deviceId));
|
||||
auto clockRate = prop.clockRate;
|
||||
|
||||
waitKernel<<<1, 1, 0, stream>>>(clockRate, 2); // Wait for 2 seconds
|
||||
|
||||
HIP_CHECK(hipEventRecord(waitEvent, stream));
|
||||
|
||||
// Make sure stream is waiting for data to be set
|
||||
HIP_CHECK_ERROR(hipEventQuery(waitEvent), hipErrorNotReady);
|
||||
|
||||
HIP_CHECK(hipStreamWaitEvent(stream, waitEvent, 0));
|
||||
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
HIP_CHECK(hipEventDestroy(waitEvent));
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipStreamWaitEvent_DifferentStreams") {
|
||||
hipStream_t blockedStreamA{nullptr}, streamBlockedOnStreamA{nullptr}, unblockingStream{nullptr};
|
||||
hipEvent_t waitEvent{nullptr};
|
||||
@@ -87,12 +139,15 @@ TEST_CASE("Unit_hipStreamWaitEvent_DifferentStreams") {
|
||||
REQUIRE(streamBlockedOnStreamA != nullptr);
|
||||
REQUIRE(waitEvent != nullptr);
|
||||
|
||||
int deviceId {};
|
||||
HIP_CHECK(hipGetDevice(&deviceId));
|
||||
|
||||
hipDeviceProp_t prop{};
|
||||
HIP_CHECK(hipGetDeviceProperties(&prop, 0));
|
||||
HIP_CHECK(hipGetDeviceProperties(&prop, deviceId));
|
||||
auto clockRate = prop.clockRate;
|
||||
|
||||
waitKernel<<<1, 1, 0, blockedStreamA>>>(clockRate,
|
||||
3); // wait for 5 seconds
|
||||
3); // wait for 3 seconds
|
||||
HIP_CHECK(hipEventRecord(waitEvent, blockedStreamA));
|
||||
|
||||
// Make sure stream is waiting for data to be set
|
||||
@@ -102,12 +157,12 @@ TEST_CASE("Unit_hipStreamWaitEvent_DifferentStreams") {
|
||||
|
||||
waitKernel<<<1, 1, 0, streamBlockedOnStreamA>>>(clockRate, 2); // Wait for 2 seconds
|
||||
|
||||
// Make sure stream is waiting for event on blockedStreamA
|
||||
HIP_CHECK_ERROR(hipStreamQuery(streamBlockedOnStreamA), hipErrorNotReady);
|
||||
|
||||
HIP_CHECK(hipStreamSynchronize(unblockingStream));
|
||||
|
||||
HIP_CHECK(hipStreamSynchronize(blockedStreamA));
|
||||
|
||||
// Make sure streamBlockedOnStreamA waited for event on blockedStreamA
|
||||
HIP_CHECK_ERROR(hipStreamQuery(streamBlockedOnStreamA), hipErrorNotReady);
|
||||
HIP_CHECK(hipStreamSynchronize(streamBlockedOnStreamA));
|
||||
|
||||
// Check that both streams have finished
|
||||
|
||||
Reference in New Issue
Block a user