From fbb4e4c2c339cfbc0c853fc971ce316938f7f430 Mon Sep 17 00:00:00 2001 From: Tony Date: Thu, 7 May 2020 19:15:42 -0400 Subject: [PATCH] Make HSA_QUEUE_TYPE_COOPERATIVE a queue type value - Correct defintion of HSA_QUEUE_TYPE_COOPERATIVE to be a queue type and not a bit mask. - Correct implementation of hsa_queue_type_t to treat is as an enumeration type and not a bit mask. In particular HSA_QUEUE_TYPE_COOPERATIVE is a distinct queue type that uses the multi producer protocol, and is not a bit set value. Change-Id: I9415be8853671e5511e16e306caf16020e8c84af [ROCm/ROCR-Runtime commit: bccb25fc33c3ca635f34e96a9b74a8e2be3686ed] --- .../core/runtime/amd_aql_queue.cpp | 2 +- .../core/runtime/amd_gpu_agent.cpp | 4 +- .../runtime/hsa-runtime/core/runtime/hsa.cpp | 9 +--- .../runtime/hsa-runtime/inc/hsa.h | 41 ++++++++++--------- 4 files changed, 26 insertions(+), 30 deletions(-) diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp index 4ce4f9093b..6ee90cbb29 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp @@ -309,7 +309,7 @@ AqlQueue::~AqlQueue() { } void AqlQueue::Destroy() { - if (amd_queue_.hsa_queue.type & HSA_QUEUE_TYPE_COOPERATIVE) { + if (amd_queue_.hsa_queue.type == HSA_QUEUE_TYPE_COOPERATIVE) { agent_->GWSRelease(); return; } diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp index b40049a394..c6350024f7 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp @@ -598,7 +598,7 @@ void GpuAgent::InitGWS() { if (status != HSAKMT_STATUS_SUCCESS) throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES, "GWS allocation failed."); - queue->amd_queue_.hsa_queue.type = HSA_QUEUE_TYPE_COOPERATIVE | HSA_QUEUE_TYPE_MULTI; + queue->amd_queue_.hsa_queue.type = HSA_QUEUE_TYPE_COOPERATIVE; gws_queue_.ref_ct_ = 0; return queue.release(); }); @@ -933,7 +933,7 @@ hsa_status_t GpuAgent::QueueCreate(size_t size, hsa_queue_type32_t queue_type, uint32_t group_segment_size, core::Queue** queue) { // Handle GWS queues. - if (queue_type & HSA_QUEUE_TYPE_COOPERATIVE) { + if (queue_type == HSA_QUEUE_TYPE_COOPERATIVE) { ScopedAcquire lock(&gws_queue_.lock_); auto ret = (*gws_queue_.queue_).get(); if (ret != nullptr) { diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp index 7de29f965a..1d08c25643 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp @@ -700,13 +700,8 @@ hsa_status_t hsa_queue_create( agent->GetInfo(HSA_AGENT_INFO_QUEUE_TYPE, &agent_queue_type); assert(HSA_STATUS_SUCCESS == status); - if (agent_queue_type == HSA_QUEUE_TYPE_SINGLE && - ((type & HSA_QUEUE_TYPE_SINGLE) != HSA_QUEUE_TYPE_SINGLE)) { - return HSA_STATUS_ERROR_INVALID_QUEUE_CREATION; - } - - if ((type & HSA_QUEUE_TYPE_COOPERATIVE) && - ((type & HSA_QUEUE_TYPE_SINGLE) != HSA_QUEUE_TYPE_MULTI)) { + if ((agent_queue_type == HSA_QUEUE_TYPE_SINGLE) && + (type != HSA_QUEUE_TYPE_SINGLE)) { return HSA_STATUS_ERROR_INVALID_QUEUE_CREATION; } diff --git a/projects/rocr-runtime/runtime/hsa-runtime/inc/hsa.h b/projects/rocr-runtime/runtime/hsa-runtime/inc/hsa.h index 54dc78460b..0e5936852e 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/inc/hsa.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/inc/hsa.h @@ -2,24 +2,24 @@ // // The University of Illinois/NCSA // Open Source License (NCSA) -// +// // Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved. -// +// // Developed by: -// +// // AMD Research and AMD HSA Software Development -// +// // Advanced Micro Devices, Inc. -// +// // www.amd.com -// +// // Permission is hereby granted, free of charge, to any person obtaining a copy // of this software and associated documentation files (the "Software"), to // deal with 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: -// +// // - Redistributions of source code must retain the above copyright notice, // this list of conditions and the following disclaimers. // - Redistributions in binary form must reproduce the above copyright @@ -29,7 +29,7 @@ // nor the names of its contributors may be used to endorse or promote // products derived from this Software without specific prior written // permission. -// +// // 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 @@ -632,7 +632,7 @@ hsa_status_t HSA_API hsa_system_major_extension_supported( uint16_t version_major, uint16_t *version_minor, bool* result); - + /** * @deprecated @@ -711,7 +711,7 @@ hsa_status_t HSA_API hsa_system_get_major_extension_table( uint16_t extension, uint16_t version_major, size_t table_length, - void *table); + void *table); /** * @brief Struct containing an opaque handle to an agent, a device that participates in @@ -1283,7 +1283,7 @@ hsa_status_t HSA_API hsa_agent_major_extension_supported( uint16_t version_major, uint16_t *version_minor, bool* result); - + /** @} */ @@ -2184,24 +2184,25 @@ typedef struct hsa_region_s { */ typedef enum { /** - * Queue supports multiple producers. + * Queue supports multiple producers. Use of multiproducer queue mechanics is + * required. */ HSA_QUEUE_TYPE_MULTI = 0, /** * Queue only supports a single producer. In some scenarios, the application * may want to limit the submission of AQL packets to a single agent. Queues * that support a single producer may be more efficient than queues supporting - * multiple producers. + * multiple producers. Use of multiproducer queue mechanics is not supported. */ HSA_QUEUE_TYPE_SINGLE = 1, /** - * Queue supports cooperative dispatches able to use GWS synchronization. - * Queues of this type must also be of type HSA_QUEUE_TYPE_MULTI and - * may be limited in number. The runtime may return the same queue to serve - * multiple hsa_queue_create calls when this type is given. Callers must - * inspect the returned queue to discover queue size. Queues of this type - * are reference counted and require a matching number of hsa_queue_destroy - * calls to release. Use of multiproducer queue mechanics is required. See + * Queue supports multiple producers and cooperative dispatches. Cooperative + * dispatches are able to use GWS synchronization. Queues of this type may be + * limited in number. The runtime may return the same queue to serve multiple + * ::hsa_queue_create calls when this type is given. Callers must inspect the + * returned queue to discover queue size. Queues of this type are reference + * counted and require a matching number of ::hsa_queue_destroy calls to + * release. Use of multiproducer queue mechanics is required. See * ::HSA_AMD_AGENT_INFO_COOPERATIVE_QUEUES to query agent support for this * type. */