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

SWDEV-337452 - Changing Clock64 to WallClock64 in tests for gfx11. (#78)

Browse Source 
Change-Id: I484fe9ff7cd56c70a37a3ac5a4a55812f8557259
This commit is contained in:
ROCm CI Service Account authored and GitHub committed 2023-01-06 15:05:21 -08:00
1 parent 70c61d55a4
commit 87fac87657
10 files changed
+339 -50
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

catch/unit/memory/hipMemPoolApi.cc
+93 -17
View File
@@ -123,6 +123,21 @@ __global__ void kernel500ms(float* hostRes, int clkRate) {
}
}
__global__ void kernel500ms_gfx11(float* hostRes, int clkRate) {
#if HT_AMD
int tid = threadIdx.x + blockIdx.x * blockDim.x;
hostRes[tid] = tid + 1;
__threadfence_system();
// expecting that the data is getting flushed to host here!
uint64_t start = wall_clock64()/clkRate, cur;
if (clkRate > 1) {
do { cur = wall_clock64()/clkRate-start;}while (cur < wait_ms);
} else {
do { cur = wall_clock64()/start;}while (cur < wait_ms);
}
#endif
}
TEST_CASE("Unit_hipMemPoolApi_BasicAlloc") {
int mem_pool_support = 0;
HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0));
@@ -147,9 +162,14 @@ TEST_CASE("Unit_hipMemPoolApi_BasicAlloc") {
int blocks = 1024;
int clkRate;
hipMemPoolAttr attr;
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
if (IsGfx11()) {
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0));
kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate);
} else {
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
}
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(B), stream));
@@ -229,9 +249,14 @@ TEST_CASE("Unit_hipMemPoolApi_BasicTrim") {
int blocks = 2;
int clkRate;
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
if (IsGfx11()) {
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0));
kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate);
} else {
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
}
hipMemPoolAttr attr;
attr = hipMemPoolAttrReleaseThreshold;
@@ -312,9 +337,15 @@ TEST_CASE("Unit_hipMemPoolApi_BasicReuse") {
int blocks = 2;
int clkRate;
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
if (IsGfx11()) {
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0));
kernel500ms_gfx11<<<32, blocks, 0, stream>>>(A, clkRate);
} else {
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
}
hipMemPoolAttr attr;
// Not a real free, since kernel isn't done
@@ -329,7 +360,11 @@ TEST_CASE("Unit_hipMemPoolApi_BasicReuse") {
HIP_CHECK(hipStreamSynchronize(stream));
// Second kernel launch with new memory
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
if (IsGfx11()) {
kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate);
} else {
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
}
HIP_CHECK(hipStreamSynchronize(stream));
@@ -369,7 +404,11 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") {
hipMemPoolAttr attr;
int blocks = 2;
int clkRate;
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
if (IsGfx11()) {
HIPCHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0));
} else {
HIPCHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
}
float *A, *B, *C;
hipStream_t stream, stream2;
@@ -395,7 +434,11 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") {
HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value));
// Run kernel for 500 ms in the first stream
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
if (IsGfx11()) {
kernel500ms_gfx11<<<32, blocks, 0, stream>>>(A, clkRate);
} else {
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
}
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(A), stream));
@@ -410,7 +453,11 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") {
REQUIRE(A != B);
// Run kernel with the new memory in the second stream
kernel500ms<<<32, blocks, 0, stream2>>>(B, clkRate);
if (IsGfx11()) {
kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate);
} else {
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
}
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipStreamSynchronize(stream2));
@@ -428,7 +475,13 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") {
HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value));
// Run kernel for 500 ms in the first stream
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
if (IsGfx11()) {
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0));
kernel500ms_gfx11<<<32, blocks, 0, stream>>>(A, clkRate);
} else {
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
}
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(A), stream));
@@ -443,7 +496,11 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") {
REQUIRE(A == B);
// Run kernel with the new memory in the second stream
kernel500ms<<<32, blocks, 0, stream2>>>(B, clkRate);
if (IsGfx11()) {
kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate);
} else {
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
}
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipStreamSynchronize(stream2));
@@ -461,7 +518,12 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") {
HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value));
// Run kernel for 500 ms in the first stream
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
if (IsGfx11()) {
kernel500ms_gfx11<<<32, blocks, 0, stream>>>(A, clkRate);
} else {
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
}
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(A), stream));
@@ -473,7 +535,11 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") {
REQUIRE(A != B);
// Run kernel with the new memory in the second stream
kernel500ms<<<32, blocks, 0, stream2>>>(B, clkRate);
if (IsGfx11()) {
kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate);
} else {
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
}
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipStreamSynchronize(stream2));
@@ -510,9 +576,15 @@ TEST_CASE("Unit_hipMemPoolApi_Default") {
int blocks = 2;
int clkRate;
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
if (IsGfx11()) {
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0));
kernel500ms_gfx11<<<32, blocks, 0, stream>>>(A, clkRate);
} else {
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
}
hipMemPoolAttr attr;
// Not a real free, since kernel isn't done
@@ -527,7 +599,11 @@ TEST_CASE("Unit_hipMemPoolApi_Default") {
HIP_CHECK(hipStreamSynchronize(stream));
// Second kernel launch with new memory
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
if (IsGfx11()) {
kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate);
} else {
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
}
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(B), stream));
catch/unit/memory/hipMemoryAllocateCoherent.cc
+22 -2
View File
@@ -41,6 +41,21 @@ __global__ void Kernel(float* hostRes, int clkRate) {
}
}
__global__ void Kernel_gfx11(float* hostRes, int clkRate) {
#if HT_AMD
int tid = threadIdx.x + blockIdx.x * blockDim.x;
hostRes[tid] = tid + 1;
__threadfence_system();
// expecting that the data is getting flushed to host here!
uint64_t start = wall_clock64()/clkRate, cur;
if (clkRate > 1) {
do { cur = wall_clock64()/clkRate-start;}while (cur < wait_sec);
} else {
do { cur = wall_clock64()/start;}while (cur < wait_sec);
}
#endif
}
TEST_CASE("Unit_hipHostMalloc_CoherentAccess") {
int blocks = 2;
float* hostRes;
@@ -49,9 +64,14 @@ TEST_CASE("Unit_hipHostMalloc_CoherentAccess") {
hostRes[0] = 0;
hostRes[1] = 0;
int clkRate;
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
if (IsGfx11()) {
HIPCHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0));
} else {
HIPCHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
}
std::cout << clkRate << std::endl;
hipLaunchKernelGGL(HIP_KERNEL_NAME(Kernel), dim3(1), dim3(blocks),
auto Kernel_used = IsGfx11() ? Kernel_gfx11 : Kernel;
hipLaunchKernelGGL(HIP_KERNEL_NAME(Kernel_used), dim3(1), dim3(blocks),
0, 0, hostRes, clkRate);
HIP_CHECK(hipGetLastError());
int eleCounter = 0;