AI/rocm-systems
2
0
Rozštěpit 0
Zdrojový kód Úkoly Pull requesty Akce Balíčky Projekty Vydání Wiki Aktivita

EXSWHTEC-382 - Implement tests for Launch API functions #454

Procházet zdrojové kódy 
Change-Id: I0720758144e89adaa43bcbcc6262dbb16cd4e2be
Tento commit je obsažen v:
Mirza Halilčević
2023-12-28 19:32:13 +01:00
odevzdal Rakesh Roy
rodič e0c3f64e78
revize 66e2885107
8 změnil soubory, kde provedl 249 přidání a 23 odebrání
Stáhněte soubor záplaty Stáhněte rozdílový soubor Rozbalit všechny soubory Sbalit všechny soubory
catch/unit/executionControl/hipLaunchCooperativeKernel.cc
+9 -9
Zobrazit soubor
@@ -61,19 +61,19 @@ TEST_CASE("Unit_hipLaunchCooperativeKernel_Positive_Parameters") {
}
SECTION("blockDim.x == maxBlockDimX") {
const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimX);
const unsigned int x = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimX, 0);
HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
dim3{x, 1, 1}, nullptr, 0, nullptr));
}
SECTION("blockDim.y == maxBlockDimY") {
const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimY);
const unsigned int y = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimY, 0);
HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
dim3{y, 1, 1}, nullptr, 0, nullptr));
}
SECTION("blockDim.z == maxBlockDimZ") {
const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimZ);
const unsigned int z = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimZ, 0);
HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
dim3{z, 1, 1}, nullptr, 0, nullptr));
}
@@ -128,28 +128,28 @@ TEST_CASE("Unit_hipLaunchCooperativeKernel_Negative_Parameters") {
}
SECTION("blockDim.x > maxBlockDimX") {
const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimX) + 1u;
const unsigned int x = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimX, 0) + 1u;
HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
dim3{x, 1, 1}, nullptr, 0, nullptr),
hipErrorInvalidConfiguration);
}
SECTION("blockDim.y > maxBlockDimY") {
const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimY) + 1u;
const unsigned int y = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimY, 0) + 1u;
HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
dim3{1, y, 1}, nullptr, 0, nullptr),
hipErrorInvalidConfiguration);
}
SECTION("blockDim.z > maxBlockDimZ") {
const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimZ) + 1u;
const unsigned int z = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimZ, 0) + 1u;
HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
dim3{1, 1, z}, nullptr, 0, nullptr),
hipErrorInvalidConfiguration);
}
SECTION("blockDim.x * blockDim.y * blockDim.z > maxThreadsPerBlock") {
const unsigned int max = GetDeviceAttribute(0, hipDeviceAttributeMaxThreadsPerBlock);
const unsigned int max = GetDeviceAttribute(hipDeviceAttributeMaxThreadsPerBlock, 0);
const unsigned int dim = std::ceil(std::cbrt(max));
HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
dim3{dim, dim, dim}, nullptr, 0, nullptr),
@@ -163,7 +163,7 @@ TEST_CASE("Unit_hipLaunchCooperativeKernel_Negative_Parameters") {
HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&max_blocks,
reinterpret_cast<void*>(kernel), 1, 0));
const unsigned int multiproc_count =
GetDeviceAttribute(0, hipDeviceAttributeMultiprocessorCount);
GetDeviceAttribute(hipDeviceAttributeMultiprocessorCount, 0);
const unsigned int dim = std::ceil(std::cbrt(max_blocks * multiproc_count));
HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast<void*>(kernel), dim3{dim, dim, dim},
dim3{1, 1, 1}, nullptr, 0, nullptr),
@@ -171,7 +171,7 @@ TEST_CASE("Unit_hipLaunchCooperativeKernel_Negative_Parameters") {
}
SECTION("sharedMemBytes > maxSharedMemoryPerBlock") {
const unsigned int max = GetDeviceAttribute(0, hipDeviceAttributeMaxSharedMemoryPerBlock) + 1u;
const unsigned int max = GetDeviceAttribute(hipDeviceAttributeMaxSharedMemoryPerBlock, 0) + 1u;
HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
dim3{1, 1, 1}, nullptr, max, nullptr),
hipErrorCooperativeLaunchTooLarge);