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

rocrtst: Move kernel object loading outside of loops

Browse Source 
Negative queue validation tests were doing many redundant from-file
kernel object loads in a loop. This was creating many simulataneous open
file handles within many dynamically allocated CodeObject objects. While
the CodeObject class implements RAII on the file handles to cleanup on
destruction, clear_code_object() only gets called on the destruction of
the TestBase-derived test objects (these being a suite abstraction).

Due to this we were hitting file open() EMFILE errors (too many open
files) in gfx94x CPX mode. Move LoadKernelFromObjFile outside of the
test loops and clear_code_object() for each test on each agent.

Signed-off-by: Graham Sider <Graham.Sider@amd.com>
Change-Id: I6f9d23fd122720c49a58c22698f097906d2fc97c


[ROCm/ROCR-Runtime commit: 7a4c9273d7]
This commit is contained in:
Graham Sider
2023-04-06 16:09:11 -04:00
committed by David Yat Sin
parent 11541cc283
commit e2fc46c189
1 changed files with 31 additions and 25 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocr-runtime/rocrtst/suites/negative/queue_validation.cc
+31 -25
View File
@@ -175,6 +175,10 @@ void QueueValidation::QueueValidationForInvalidDimension(hsa_agent_t cpuAgent,
hsa_agent_t gpuAgent) {
hsa_status_t err;
// Create the executable, get symbol by name and load the code object
err = rocrtst::LoadKernelFromObjFile(this, &gpuAgent);
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
// get queue size
uint32_t queue_max = 0;
err = hsa_agent_get_info(gpuAgent,
@@ -201,11 +205,6 @@ void QueueValidation::QueueValidationForInvalidDimension(hsa_agent_t cpuAgent,
CallbackQueueErrorHandling, &user_data[ii], 0, 0, &queue[ii]);
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
// Create the executable, get symbol by name and load the code object
err = rocrtst::LoadKernelFromObjFile(this, &gpuAgent);
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
// setting the dimesion more than 3
aql().setup = 4;
aql().kernel_object = kernel_object();
@@ -249,6 +248,8 @@ void QueueValidation::QueueValidationForInvalidDimension(hsa_agent_t cpuAgent,
ASSERT_EQ(user_data[ii].cb_triggered, true);
if (queue[ii]) { hsa_queue_destroy(queue[ii]); }
}
clear_code_object();
}
@@ -256,6 +257,10 @@ void QueueValidation::QueueValidationInvalidGroupMemory(hsa_agent_t cpuAgent,
hsa_agent_t gpuAgent) {
hsa_status_t err;
// Create the executable, get symbol by name and load the code object
err = rocrtst::LoadKernelFromObjFile(this, &gpuAgent);
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
// Fill up the kernel packet except header
err = rocrtst::InitializeAQLPacket(this, &aql());
ASSERT_EQ(HSA_STATUS_SUCCESS, err);
@@ -287,11 +292,6 @@ void QueueValidation::QueueValidationInvalidGroupMemory(hsa_agent_t cpuAgent,
CallbackQueueErrorHandling, &user_data[ii], 0, 0, &queue[ii]);
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
// Create the executable, get symbol by name and load the code object
err = rocrtst::LoadKernelFromObjFile(this, &gpuAgent);
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
aql().kernel_object = kernel_object();
// Request a large group memory segment size
aql().group_segment_size = (uint32_t)-1;
@@ -336,12 +336,19 @@ void QueueValidation::QueueValidationInvalidGroupMemory(hsa_agent_t cpuAgent,
ASSERT_EQ(user_data[ii].cb_triggered, true);
if (queue[ii]) { hsa_queue_destroy(queue[ii]); }
}
clear_code_object();
}
void QueueValidation::QueueValidationForInvalidKernelObject(hsa_agent_t cpuAgent,
hsa_agent_t gpuAgent) {
hsa_status_t err;
// Create the executable, get symbol by name and load the code object
err = rocrtst::LoadKernelFromObjFile(this, &gpuAgent);
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
// Fill up the kernel packet except header
err = rocrtst::InitializeAQLPacket(this, &aql());
ASSERT_EQ(HSA_STATUS_SUCCESS, err);
@@ -372,11 +379,6 @@ void QueueValidation::QueueValidationForInvalidKernelObject(hsa_agent_t cpuAgent
CallbackQueueErrorHandling, &user_data[ii], 0, 0, &queue[ii]);
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
// Create the executable, get symbol by name and load the code object
err = rocrtst::LoadKernelFromObjFile(this, &gpuAgent);
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
// setting the null code object
aql().kernel_object = 0;
@@ -420,12 +422,18 @@ void QueueValidation::QueueValidationForInvalidKernelObject(hsa_agent_t cpuAgent
ASSERT_EQ(user_data[ii].cb_triggered, true);
if (queue[ii]) { hsa_queue_destroy(queue[ii]); }
}
clear_code_object();
}
void QueueValidation::QueueValidationForInvalidPacket(hsa_agent_t cpuAgent,
hsa_agent_t gpuAgent) {
hsa_status_t err;
// Create the executable, get symbol by name and load the code object
err = rocrtst::LoadKernelFromObjFile(this, &gpuAgent);
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
// Fill up the kernel packet except header
err = rocrtst::InitializeAQLPacket(this, &aql());
ASSERT_EQ(HSA_STATUS_SUCCESS, err);
@@ -456,11 +464,6 @@ void QueueValidation::QueueValidationForInvalidPacket(hsa_agent_t cpuAgent,
CallbackQueueErrorHandling, &user_data[ii], 0, 0, &queue[ii]);
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
// Create the executable, get symbol by name and load the code object
err = rocrtst::LoadKernelFromObjFile(this, &gpuAgent);
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
const uint32_t queue_mask = queue[ii]->size - 1;
// Load index for writing header later to command queue at same index
@@ -499,12 +502,18 @@ void QueueValidation::QueueValidationForInvalidPacket(hsa_agent_t cpuAgent,
ASSERT_EQ(user_data[ii].cb_triggered, true);
if (queue[ii]) { hsa_queue_destroy(queue[ii]); }
}
clear_code_object();
}
void QueueValidation::QueueValidationForInvalidWorkGroupSize(hsa_agent_t cpuAgent,
hsa_agent_t gpuAgent) {
hsa_status_t err;
// Create the executable, get symbol by name and load the code object
err = rocrtst::LoadKernelFromObjFile(this, &gpuAgent);
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
// Fill up the kernel packet except header
err = rocrtst::InitializeAQLPacket(this, &aql());
ASSERT_EQ(HSA_STATUS_SUCCESS, err);
@@ -537,11 +546,6 @@ void QueueValidation::QueueValidationForInvalidWorkGroupSize(hsa_agent_t cpuAgen
CallbackQueueErrorHandling, &user_data[ii][jj - 1], 0, 0, &queue[ii]);
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
// Create the executable, get symbol by name and load the code object
err = rocrtst::LoadKernelFromObjFile(this, &gpuAgent);
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
aql().setup |= jj << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;
aql().workgroup_size_x = (jj == 1) ? (uint16_t)-1 : 1;
aql().workgroup_size_y = (jj == 2) ? (uint16_t)-1 : 1;
@@ -591,6 +595,8 @@ void QueueValidation::QueueValidationForInvalidWorkGroupSize(hsa_agent_t cpuAgen
ASSERT_EQ(user_data[ii][jj].cb_triggered, true);
}
}
clear_code_object();
}