From e2fc46c189c9f77914ecaa3ffbad7d3946211890 Mon Sep 17 00:00:00 2001 From: Graham Sider Date: Thu, 6 Apr 2023 16:09:11 -0400 Subject: [PATCH] rocrtst: Move kernel object loading outside of loops 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 Change-Id: I6f9d23fd122720c49a58c22698f097906d2fc97c [ROCm/ROCR-Runtime commit: 7a4c9273d7d8b20e11e2e10c4b45d05c8ae9c9d1] --- .../suites/negative/queue_validation.cc | 56 ++++++++++--------- 1 file changed, 31 insertions(+), 25 deletions(-) diff --git a/projects/rocr-runtime/rocrtst/suites/negative/queue_validation.cc b/projects/rocr-runtime/rocrtst/suites/negative/queue_validation.cc index 260c377c52..d03a9b4a4c 100755 --- a/projects/rocr-runtime/rocrtst/suites/negative/queue_validation.cc +++ b/projects/rocr-runtime/rocrtst/suites/negative/queue_validation.cc @@ -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(); }