SWDEV-545950 - Add hipStreamCopyAttributes API Implementation (#914)
* SWDEV-545950 - Add hipStreamCopyAttributes API Implementation * Add unit test for hipStreamCopyAttributes API * Add ChangeLog and nvidia mapping for the API * Update rocprofiler-sdk with new HIP API details * [rocprofiler-sdk] handle hipStreamCopyAttributes in stream tracing service - this new HIP function has multiple stream arguments and needs to be skipped because it does not have an explicit create/destroy/set functionality * Update HIP_RUNTIME_API_TABLE_STEP_VERSION in clr and rocprofiler-sdk * Resolve merge conflicts --------- Co-authored-by: Jonathan R. Madsen <jonathanrmadsen@gmail.com>
Этот коммит содержится в:
Satyanvesh Dittakavi
коммит произвёл
GitHub
GitHub
родитель
7f79d0febc
Коммит
46e683d41a
@@ -2,6 +2,13 @@
|
||||
|
||||
Full documentation for HIP is available at [rocm.docs.amd.com](https://rocm.docs.amd.com/projects/HIP/en/latest/index.html)
|
||||
|
||||
## HIP 7.2 for ROCm 7.2
|
||||
|
||||
### Added
|
||||
|
||||
* New HIP APIs
|
||||
- `hipStreamCopyAttributes` Copies attributes from source stream to destination stream
|
||||
|
||||
## HIP 7.1 for ROCm 7.1
|
||||
|
||||
### Added
|
||||
|
||||
@@ -63,7 +63,7 @@
|
||||
#define HIP_API_TABLE_STEP_VERSION 0
|
||||
#define HIP_COMPILER_API_TABLE_STEP_VERSION 0
|
||||
#define HIP_TOOLS_API_TABLE_STEP_VERSION 0
|
||||
#define HIP_RUNTIME_API_TABLE_STEP_VERSION 15
|
||||
#define HIP_RUNTIME_API_TABLE_STEP_VERSION 16
|
||||
|
||||
// HIP API interface
|
||||
// HIP compiler dispatch functions
|
||||
@@ -713,6 +713,7 @@ typedef hipError_t (*t_hipStreamAddCallback)(hipStream_t stream, hipStreamCallba
|
||||
typedef hipError_t (*t_hipStreamAttachMemAsync)(hipStream_t stream, void* dev_ptr, size_t length,
|
||||
unsigned int flags);
|
||||
typedef hipError_t (*t_hipStreamBeginCapture)(hipStream_t stream, hipStreamCaptureMode mode);
|
||||
typedef hipError_t (*t_hipStreamCopyAttributes)(hipStream_t dst, hipStream_t src);
|
||||
typedef hipError_t (*t_hipStreamCreate)(hipStream_t* stream);
|
||||
typedef hipError_t (*t_hipStreamCreateWithFlags)(hipStream_t* stream, unsigned int flags);
|
||||
typedef hipError_t (*t_hipStreamCreateWithPriority)(hipStream_t* stream, unsigned int flags,
|
||||
@@ -1679,8 +1680,11 @@ struct HipDispatchTable {
|
||||
t_hipLibraryGetKernel hipLibraryGetKernel_fn;
|
||||
t_hipLibraryGetKernelCount hipLibraryGetKernelCount_fn;
|
||||
|
||||
// HIP_RUNTIME_API_TABLE_STEP_VERSION = 16
|
||||
t_hipStreamCopyAttributes hipStreamCopyAttributes_fn;
|
||||
|
||||
// DO NOT EDIT ABOVE!
|
||||
// HIP_RUNTIME_API_TABLE_STEP_VERSION == 15
|
||||
// HIP_RUNTIME_API_TABLE_STEP_VERSION == 17
|
||||
|
||||
// ******************************************************************************************* //
|
||||
//
|
||||
|
||||
@@ -462,7 +462,8 @@ enum hip_api_id_t {
|
||||
HIP_API_ID_hipLibraryGetKernel = 442,
|
||||
HIP_API_ID_hipLibraryGetKernelCount = 443,
|
||||
HIP_API_ID_hipMemGetHandleForAddressRange = 444,
|
||||
HIP_API_ID_LAST = 444,
|
||||
HIP_API_ID_hipStreamCopyAttributes = 445,
|
||||
HIP_API_ID_LAST = 445,
|
||||
|
||||
HIP_API_ID_hipChooseDevice = HIP_API_ID_CONCAT(HIP_API_ID_,hipChooseDevice),
|
||||
HIP_API_ID_hipGetDeviceProperties = HIP_API_ID_CONCAT(HIP_API_ID_,hipGetDeviceProperties),
|
||||
@@ -885,6 +886,7 @@ static inline const char* hip_api_name(const uint32_t id) {
|
||||
case HIP_API_ID_hipStreamBatchMemOp: return "hipStreamBatchMemOp";
|
||||
case HIP_API_ID_hipStreamBeginCapture: return "hipStreamBeginCapture";
|
||||
case HIP_API_ID_hipStreamBeginCaptureToGraph: return "hipStreamBeginCaptureToGraph";
|
||||
case HIP_API_ID_hipStreamCopyAttributes: return "hipStreamCopyAttributes";
|
||||
case HIP_API_ID_hipStreamCreate: return "hipStreamCreate";
|
||||
case HIP_API_ID_hipStreamCreateWithFlags: return "hipStreamCreateWithFlags";
|
||||
case HIP_API_ID_hipStreamCreateWithPriority: return "hipStreamCreateWithPriority";
|
||||
@@ -1323,6 +1325,7 @@ static inline uint32_t hipApiIdByName(const char* name) {
|
||||
if (strcmp("hipStreamBatchMemOp", name) == 0) return HIP_API_ID_hipStreamBatchMemOp;
|
||||
if (strcmp("hipStreamBeginCapture", name) == 0) return HIP_API_ID_hipStreamBeginCapture;
|
||||
if (strcmp("hipStreamBeginCaptureToGraph", name) == 0) return HIP_API_ID_hipStreamBeginCaptureToGraph;
|
||||
if (strcmp("hipStreamCopyAttributes", name) == 0) return HIP_API_ID_hipStreamCopyAttributes;
|
||||
if (strcmp("hipStreamCreate", name) == 0) return HIP_API_ID_hipStreamCreate;
|
||||
if (strcmp("hipStreamCreateWithFlags", name) == 0) return HIP_API_ID_hipStreamCreateWithFlags;
|
||||
if (strcmp("hipStreamCreateWithPriority", name) == 0) return HIP_API_ID_hipStreamCreateWithPriority;
|
||||
@@ -3731,6 +3734,10 @@ typedef struct hip_api_data_s {
|
||||
size_t numDependencies;
|
||||
hipStreamCaptureMode mode;
|
||||
} hipStreamBeginCaptureToGraph;
|
||||
struct {
|
||||
hipStream_t dst;
|
||||
hipStream_t src;
|
||||
} hipStreamCopyAttributes;
|
||||
struct {
|
||||
hipStream_t* stream;
|
||||
hipStream_t stream__val;
|
||||
@@ -6397,6 +6404,11 @@ typedef struct hip_api_data_s {
|
||||
cb_data.args.hipStreamBeginCaptureToGraph.numDependencies = (size_t)numDependencies; \
|
||||
cb_data.args.hipStreamBeginCaptureToGraph.mode = (hipStreamCaptureMode)mode; \
|
||||
};
|
||||
// hipStreamCopyAttributes[('hipStream_t', 'dst'), ('hipStream_t', 'src')]
|
||||
#define INIT_hipStreamCopyAttributes_CB_ARGS_DATA(cb_data) { \
|
||||
cb_data.args.hipStreamCopyAttributes.dst = (hipStream_t)dst; \
|
||||
cb_data.args.hipStreamCopyAttributes.src = (hipStream_t)src; \
|
||||
};
|
||||
// hipStreamCreate[('hipStream_t*', 'stream')]
|
||||
#define INIT_hipStreamCreate_CB_ARGS_DATA(cb_data) { \
|
||||
cb_data.args.hipStreamCreate.stream = (hipStream_t*)stream; \
|
||||
@@ -8227,6 +8239,9 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
|
||||
if (data->args.hipStreamBeginCaptureToGraph.dependencies) data->args.hipStreamBeginCaptureToGraph.dependencies__val = *(data->args.hipStreamBeginCaptureToGraph.dependencies);
|
||||
if (data->args.hipStreamBeginCaptureToGraph.dependencyData) data->args.hipStreamBeginCaptureToGraph.dependencyData__val = *(data->args.hipStreamBeginCaptureToGraph.dependencyData);
|
||||
break;
|
||||
// hipStreamCopyAttributes[('hipStream_t', 'dst'), ('hipStream_t', 'src')]
|
||||
case HIP_API_ID_hipStreamCopyAttributes:
|
||||
break;
|
||||
// hipStreamCreate[('hipStream_t*', 'stream')]
|
||||
case HIP_API_ID_hipStreamCreate:
|
||||
if (data->args.hipStreamCreate.stream) data->args.hipStreamCreate.stream__val = *(data->args.hipStreamCreate.stream);
|
||||
@@ -11571,6 +11586,12 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
|
||||
oss << ", numDependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamBeginCaptureToGraph.numDependencies);
|
||||
oss << ", mode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamBeginCaptureToGraph.mode);
|
||||
oss << ")";
|
||||
break;
|
||||
case HIP_API_ID_hipStreamCopyAttributes:
|
||||
oss << "hipStreamCopyAttributes(";
|
||||
oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamCopyAttributes.dst);
|
||||
oss << ", src="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamCopyAttributes.src);
|
||||
oss << ")";
|
||||
break;
|
||||
case HIP_API_ID_hipStreamCreate:
|
||||
oss << "hipStreamCreate(";
|
||||
|
||||
@@ -516,3 +516,4 @@ hipLibraryLoadFromFile
|
||||
hipLibraryUnload
|
||||
hipLibraryGetKernel
|
||||
hipLibraryGetKernelCount
|
||||
hipStreamCopyAttributes
|
||||
|
||||
@@ -594,6 +594,7 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback
|
||||
unsigned int flags);
|
||||
hipError_t hipStreamAttachMemAsync(hipStream_t stream, void* dev_ptr, size_t length,
|
||||
unsigned int flags);
|
||||
hipError_t hipStreamCopyAttributes(hipStream_t dst, hipStream_t src);
|
||||
hipError_t hipStreamBeginCapture(hipStream_t stream, hipStreamCaptureMode mode);
|
||||
hipError_t hipStreamCreate(hipStream_t* stream);
|
||||
hipError_t hipStreamCreateWithFlags(hipStream_t* stream, unsigned int flags);
|
||||
@@ -1264,6 +1265,7 @@ void UpdateDispatchTable(HipDispatchTable* ptrDispatchTable) {
|
||||
ptrDispatchTable->hipStreamAddCallback_fn = hip::hipStreamAddCallback;
|
||||
ptrDispatchTable->hipStreamAttachMemAsync_fn = hip::hipStreamAttachMemAsync;
|
||||
ptrDispatchTable->hipStreamBeginCapture_fn = hip::hipStreamBeginCapture;
|
||||
ptrDispatchTable->hipStreamCopyAttributes_fn = hip::hipStreamCopyAttributes;
|
||||
ptrDispatchTable->hipStreamCreate_fn = hip::hipStreamCreate;
|
||||
ptrDispatchTable->hipStreamCreateWithFlags_fn = hip::hipStreamCreateWithFlags;
|
||||
ptrDispatchTable->hipStreamCreateWithPriority_fn = hip::hipStreamCreateWithPriority;
|
||||
@@ -2084,15 +2086,17 @@ HIP_ENFORCE_ABI(HipDispatchTable, hipLibraryLoadFromFile_fn, 497);
|
||||
HIP_ENFORCE_ABI(HipDispatchTable, hipLibraryUnload_fn, 498);
|
||||
HIP_ENFORCE_ABI(HipDispatchTable, hipLibraryGetKernel_fn, 499);
|
||||
HIP_ENFORCE_ABI(HipDispatchTable, hipLibraryGetKernelCount_fn, 500);
|
||||
// HIP_RUNTIME_API_TABLE_STEP_VERSION == 16
|
||||
HIP_ENFORCE_ABI(HipDispatchTable, hipStreamCopyAttributes_fn, 501);
|
||||
// if HIP_ENFORCE_ABI entries are added for each new function pointer in the table, the number below
|
||||
// will be +1 of the number in the last HIP_ENFORCE_ABI line. E.g.:
|
||||
//
|
||||
// HIP_ENFORCE_ABI(<table>, <functor>, 8)
|
||||
//
|
||||
// HIP_ENFORCE_ABI_VERSIONING(<table>, 9) <- 8 + 1 = 9
|
||||
HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 501)
|
||||
HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 502)
|
||||
|
||||
static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 15,
|
||||
static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 16,
|
||||
"If you get this error, add new HIP_ENFORCE_ABI(...) code for the new function "
|
||||
"pointers and then update this check so it is true");
|
||||
#endif
|
||||
|
||||
@@ -638,6 +638,7 @@ global:
|
||||
hipLibraryUnload;
|
||||
hipLibraryGetKernel;
|
||||
hipLibraryGetKernelCount;
|
||||
hipStreamCopyAttributes;
|
||||
local:
|
||||
*;
|
||||
} hip_7.1;
|
||||
@@ -878,7 +878,8 @@ hipError_t hipStreamSetAttribute(hipStream_t stream, hipStreamAttrID attr,
|
||||
HIP_RETURN(hipErrorStreamCaptureUnsupported);
|
||||
}
|
||||
|
||||
hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
|
||||
constexpr bool wait = false;
|
||||
hip::Stream* s = hip::getStream(stream, wait);
|
||||
|
||||
switch (attr) {
|
||||
case hipStreamAttributeSynchronizationPolicy: {
|
||||
@@ -912,7 +913,8 @@ hipError_t hipStreamGetAttribute(hipStream_t stream, hipStreamAttrID attr,
|
||||
|
||||
getStreamPerThread(stream);
|
||||
|
||||
hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
|
||||
constexpr bool wait = false;
|
||||
hip::Stream* s = hip::getStream(stream, wait);
|
||||
|
||||
switch (attr) {
|
||||
case hipStreamAttributeSynchronizationPolicy: {
|
||||
@@ -930,4 +932,18 @@ hipError_t hipStreamGetAttribute(hipStream_t stream, hipStreamAttrID attr,
|
||||
|
||||
HIP_RETURN(hipSuccess);
|
||||
}
|
||||
|
||||
hipError_t hipStreamCopyAttributes(hipStream_t dst, hipStream_t src) {
|
||||
HIP_INIT_API(hipStreamCopyAttributes, dst, src);
|
||||
|
||||
getStreamPerThread(src);
|
||||
getStreamPerThread(dst);
|
||||
|
||||
constexpr bool wait = false;
|
||||
hip::Stream* src_stream = hip::getStream(src, wait);
|
||||
hip::Stream* dst_stream = hip::getStream(dst, wait);
|
||||
// Currently, SyncPolicy is the only stream attribute we can set during runtime
|
||||
dst_stream->SetSyncPolicy(src_stream->GetSyncPolicy());
|
||||
HIP_RETURN(hipSuccess);
|
||||
}
|
||||
} // namespace hip
|
||||
|
||||
@@ -1396,6 +1396,9 @@ hipError_t hipStreamAttachMemAsync(hipStream_t stream, void* dev_ptr, size_t len
|
||||
hipError_t hipStreamBeginCapture(hipStream_t stream, hipStreamCaptureMode mode) {
|
||||
return hip::GetHipDispatchTable()->hipStreamBeginCapture_fn(stream, mode);
|
||||
}
|
||||
hipError_t hipStreamCopyAttributes(hipStream_t dst, hipStream_t src) {
|
||||
return hip::GetHipDispatchTable()->hipStreamCopyAttributes_fn(dst, src);
|
||||
}
|
||||
hipError_t hipStreamCreate(hipStream_t* stream) {
|
||||
return hip::GetHipDispatchTable()->hipStreamCreate_fn(stream);
|
||||
}
|
||||
|
||||
@@ -21,7 +21,8 @@ set(TEST_SRC
|
||||
hipStreamLegacy_Ext.cc
|
||||
hipStreamLegacy_compiler_options.cc
|
||||
hipStreamGetId.cc
|
||||
hipStreamSetGetAttributes.cc)
|
||||
hipStreamSetGetAttributes.cc
|
||||
hipStreamCopyAttributes.cc)
|
||||
|
||||
if(HIP_PLATFORM MATCHES "amd")
|
||||
set(TEST_SRC ${TEST_SRC} hipStreamGetCUMask.cc hipStreamWithCUMask.cc
|
||||
|
||||
@@ -0,0 +1,97 @@
|
||||
/*
|
||||
Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in 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:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
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 THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/**
|
||||
* @addtogroup hipStreamCopyAttributes hipStreamCopyAttributes
|
||||
* @{
|
||||
* @ingroup StreamTest
|
||||
* `hipStreamCopyAttributes (hipStream_t dst, hipStream_t src)` -
|
||||
* copies attributes from one stream to other
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Test that creates two streams and copies attributes from one to another
|
||||
* ------------------------
|
||||
* - catch\unit\stream\hipStreamCopyAttributes.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 7.2
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipStreamCopyAttributes_Basic") {
|
||||
hipStream_t stream1, stream2, stream3, stream4;
|
||||
hipStreamAttrValue val1, val2;
|
||||
HIP_CHECK(hipStreamCreate(&stream1));
|
||||
HIP_CHECK(hipStreamCreate(&stream2));
|
||||
|
||||
SECTION("Two Non Default Streams") {
|
||||
val1.syncPolicy = hipSyncPolicySpin;
|
||||
|
||||
// Set the sync policy attribute of stream1 to hipSyncPolicySpin
|
||||
HIP_CHECK(hipStreamSetAttribute(stream1, hipStreamAttributeSynchronizationPolicy, &val1));
|
||||
|
||||
// Copy attributes from stream1 to stream2
|
||||
HIP_CHECK(hipStreamCopyAttributes(stream2, stream1));
|
||||
|
||||
// Query stream2 to verify the copied sync policy
|
||||
HIP_CHECK(hipStreamGetAttribute(stream2, hipStreamAttributeSynchronizationPolicy, &val2));
|
||||
|
||||
REQUIRE(val2.syncPolicy == hipSyncPolicySpin);
|
||||
}
|
||||
|
||||
SECTION("Copy attributes from Null Stream to Legacy Stream") {
|
||||
stream3 = nullptr;
|
||||
stream4 = hipStreamLegacy;
|
||||
val1.syncPolicy = hipSyncPolicyYield;
|
||||
|
||||
// Set the sync policy attribute of stream1 to hipSyncPolicySpin
|
||||
HIP_CHECK(hipStreamSetAttribute(stream3, hipStreamAttributeSynchronizationPolicy, &val1));
|
||||
|
||||
// Copy attributes from null stream to legacy stream
|
||||
HIP_CHECK(hipStreamCopyAttributes(stream4, stream3));
|
||||
|
||||
// Query stream2 to verify the copied sync policy
|
||||
HIP_CHECK(hipStreamGetAttribute(stream4, hipStreamAttributeSynchronizationPolicy, &val2));
|
||||
|
||||
REQUIRE(val2.syncPolicy == hipSyncPolicyYield);
|
||||
}
|
||||
|
||||
SECTION("Copy attributes from streamperthread to another stream") {
|
||||
stream3 = hipStreamPerThread;
|
||||
val1.syncPolicy = hipSyncPolicyBlockingSync;
|
||||
|
||||
// Set the sync policy attribute of stream1 to hipSyncPolicySpin
|
||||
HIP_CHECK(hipStreamSetAttribute(stream3, hipStreamAttributeSynchronizationPolicy, &val1));
|
||||
|
||||
// Copy attributes from streamperthread to non default stream
|
||||
HIP_CHECK(hipStreamCopyAttributes(stream2, stream3));
|
||||
|
||||
// Query stream2 to verify the copied sync policy
|
||||
HIP_CHECK(hipStreamGetAttribute(stream2, hipStreamAttributeSynchronizationPolicy, &val2));
|
||||
|
||||
REQUIRE(val2.syncPolicy == hipSyncPolicyBlockingSync);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamDestroy(stream1));
|
||||
HIP_CHECK(hipStreamDestroy(stream2));
|
||||
}
|
||||
@@ -3048,6 +3048,14 @@ hipError_t hipStreamSetAttribute(hipStream_t stream, hipStreamAttrID attr,
|
||||
hipError_t hipStreamGetAttribute(hipStream_t stream, hipStreamAttrID attr,
|
||||
hipStreamAttrValue* value_out);
|
||||
|
||||
/**
|
||||
*@brief Copies attributes from source stream to destination stream.
|
||||
* @param[in] dst - Destination stream
|
||||
* @param[in] src - Source stream
|
||||
* @returns #hipSuccess, #hipErrorInvalidValue
|
||||
*/
|
||||
hipError_t hipStreamCopyAttributes(hipStream_t dst, hipStream_t src);
|
||||
|
||||
// end doxygen Stream
|
||||
/**
|
||||
* @}
|
||||
|
||||
@@ -3346,6 +3346,10 @@ inline static hipError_t hipStreamGetAttribute(hipStream_t stream, hipStreamAttr
|
||||
return hipCUDAErrorTohipError(cudaStreamGetAttribute(stream, attr, value_out));
|
||||
}
|
||||
|
||||
inline static hipError_t hipStreamCopyAttributes(hipStream_t dst, hipStream_t src) {
|
||||
return hipCUDAErrorTohipError(cudaStreamCopyAttributes(dst, src));
|
||||
}
|
||||
|
||||
inline static hipError_t hipDriverGetVersion(int* driverVersion) {
|
||||
return hipCUDAErrorTohipError(cudaDriverGetVersion(driverVersion));
|
||||
}
|
||||
|
||||
@@ -987,6 +987,9 @@ ROCPROFILER_ENUM_LABEL(ROCPROFILER_HIP_RUNTIME_API_ID_hipLibraryUnload)
|
||||
ROCPROFILER_ENUM_LABEL(ROCPROFILER_HIP_RUNTIME_API_ID_hipLibraryGetKernel)
|
||||
ROCPROFILER_ENUM_LABEL(ROCPROFILER_HIP_RUNTIME_API_ID_hipLibraryGetKernelCount)
|
||||
#endif
|
||||
#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 16
|
||||
ROCPROFILER_ENUM_LABEL(ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamCopyAttributes)
|
||||
#endif
|
||||
#if HIP_RUNTIME_API_TABLE_STEP_VERSION == 0
|
||||
static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 442);
|
||||
#elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 1
|
||||
@@ -1019,6 +1022,8 @@ static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 477);
|
||||
static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 496);
|
||||
#elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 15
|
||||
static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 501);
|
||||
#elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 16
|
||||
static_assert(ROCPROFILER_HIP_RUNTIME_API_ID_LAST == 502);
|
||||
#else
|
||||
# if !defined(ROCPROFILER_UNSAFE_NO_VERSION_CHECK) && \
|
||||
(defined(ROCPROFILER_CI) && ROCPROFILER_CI > 0)
|
||||
|
||||
@@ -3324,6 +3324,13 @@ typedef union rocprofiler_hip_api_args_t
|
||||
hipLibrary_t library;
|
||||
} hipLibraryGetKernelCount;
|
||||
#endif
|
||||
#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 16
|
||||
struct
|
||||
{
|
||||
hipStream_t dst;
|
||||
hipStream_t src;
|
||||
} hipStreamCopyAttributes;
|
||||
#endif
|
||||
} rocprofiler_hip_api_args_t;
|
||||
|
||||
ROCPROFILER_EXTERN_C_FINI
|
||||
|
||||
@@ -558,6 +558,9 @@ typedef enum rocprofiler_hip_runtime_api_id_t // NOLINT(performance-enum-size)
|
||||
ROCPROFILER_HIP_RUNTIME_API_ID_hipLibraryUnload,
|
||||
ROCPROFILER_HIP_RUNTIME_API_ID_hipLibraryGetKernel,
|
||||
ROCPROFILER_HIP_RUNTIME_API_ID_hipLibraryGetKernelCount,
|
||||
#endif
|
||||
#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 16
|
||||
ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamCopyAttributes,
|
||||
#endif
|
||||
ROCPROFILER_HIP_RUNTIME_API_ID_LAST,
|
||||
} rocprofiler_hip_runtime_api_id_t;
|
||||
|
||||
@@ -601,6 +601,10 @@ ROCP_SDK_ENFORCE_ABI(::HipDispatchTable, hipLibraryGetKernel_fn, 499);
|
||||
ROCP_SDK_ENFORCE_ABI(::HipDispatchTable, hipLibraryGetKernelCount_fn, 500);
|
||||
#endif
|
||||
|
||||
#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 16
|
||||
ROCP_SDK_ENFORCE_ABI(::HipDispatchTable, hipStreamCopyAttributes_fn, 501);
|
||||
#endif
|
||||
|
||||
#if HIP_RUNTIME_API_TABLE_STEP_VERSION == 0
|
||||
ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 442)
|
||||
#elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 1
|
||||
@@ -633,6 +637,8 @@ ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 477)
|
||||
ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 496)
|
||||
#elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 15
|
||||
ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 501)
|
||||
#elif HIP_RUNTIME_API_TABLE_STEP_VERSION == 16
|
||||
ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 502)
|
||||
#else
|
||||
INTERNAL_CI_ROCP_SDK_ENFORCE_ABI_VERSIONING(::HipDispatchTable, 0)
|
||||
#endif
|
||||
|
||||
@@ -627,6 +627,10 @@ HIP_API_INFO_DEFINITION_V(ROCPROFILER_HIP_TABLE_ID_Runtime, ROCPROFILER_HIP_RUNT
|
||||
HIP_API_INFO_DEFINITION_V(ROCPROFILER_HIP_TABLE_ID_Runtime, ROCPROFILER_HIP_RUNTIME_API_ID_hipLibraryGetKernel, hipLibraryGetKernel, hipLibraryGetKernel_fn, pKernel, library, name);
|
||||
HIP_API_INFO_DEFINITION_V(ROCPROFILER_HIP_TABLE_ID_Runtime, ROCPROFILER_HIP_RUNTIME_API_ID_hipLibraryGetKernelCount, hipLibraryGetKernelCount, hipLibraryGetKernelCount_fn, count, library);
|
||||
#endif
|
||||
|
||||
#if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 16
|
||||
HIP_API_INFO_DEFINITION_V(ROCPROFILER_HIP_TABLE_ID_Runtime, ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamCopyAttributes, hipStreamCopyAttributes, hipStreamCopyAttributes_fn, dst, src);
|
||||
#endif
|
||||
// clang-format on
|
||||
|
||||
#else
|
||||
|
||||
@@ -416,13 +416,38 @@ enable_stream_stack()
|
||||
return false;
|
||||
}
|
||||
|
||||
#define HIP_RUNTIME_API_TABLE_VERSION \
|
||||
ROCPROFILER_SDK_COMPUTE_VERSION( \
|
||||
HIP_RUNTIME_API_TABLE_MAJOR_VERSION, 0, HIP_RUNTIME_API_TABLE_STEP_VERSION)
|
||||
|
||||
#define HIP_STREAM_EXPLICIT_DISABLE(TABLE, OPERATION, REASON) \
|
||||
template <> \
|
||||
struct explicit_disable_update<TABLE, OPERATION> : std::true_type \
|
||||
{ \
|
||||
static constexpr auto reason = REASON; \
|
||||
};
|
||||
|
||||
template <size_t TableIdx, size_t OpIdx>
|
||||
struct explicit_disable_update : std::false_type
|
||||
{};
|
||||
|
||||
#if HIP_RUNTIME_API_TABLE_VERSION >= ROCPROFILER_SDK_COMPUTE_VERSION(0, 0, 16)
|
||||
HIP_STREAM_EXPLICIT_DISABLE(ROCPROFILER_HIP_TABLE_ID_Runtime,
|
||||
ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamCopyAttributes,
|
||||
"has multiple HIP stream arguments")
|
||||
#endif
|
||||
|
||||
#undef HIP_RUNTIME_API_TABLE_VERSION
|
||||
#undef HIP_STREAM_EXPLICIT_DISABLE
|
||||
|
||||
template <size_t TableIdx, typename Tp, size_t OpIdx>
|
||||
void
|
||||
update_table(Tp* _orig, std::integral_constant<size_t, OpIdx>)
|
||||
{
|
||||
using table_type = typename hip_table_lookup<TableIdx>::type;
|
||||
using info_type = hip_api_info<TableIdx, OpIdx>;
|
||||
using function_args_type = decltype(info_type::get_args_type());
|
||||
using table_type = typename hip_table_lookup<TableIdx>::type;
|
||||
using info_type = hip_api_info<TableIdx, OpIdx>;
|
||||
using explicit_disable_type = explicit_disable_update<TableIdx, OpIdx>;
|
||||
using function_args_type = decltype(info_type::get_args_type());
|
||||
|
||||
static_assert(info_type::table_idx == ROCPROFILER_HIP_TABLE_ID_Runtime,
|
||||
"This function should only be instantiated for HIP runtime API");
|
||||
@@ -440,7 +465,15 @@ update_table(Tp* _orig, std::integral_constant<size_t, OpIdx>)
|
||||
|
||||
constexpr auto num_args = function_args_type::size();
|
||||
|
||||
if constexpr(common::mpl::is_one_of<hipStream_t, function_args_type>::value)
|
||||
if constexpr(explicit_disable_type::value)
|
||||
{
|
||||
ROCP_INFO << fmt::format(
|
||||
"[hip stream] {} is explicitly disabled from stream tracing: {}",
|
||||
info_type::name,
|
||||
explicit_disable_type::reason);
|
||||
return;
|
||||
}
|
||||
else if constexpr(common::mpl::is_one_of<hipStream_t, function_args_type>::value)
|
||||
{
|
||||
constexpr auto stream_idx =
|
||||
common::mpl::index_of<hipStream_t, function_args_type>::value;
|
||||
|
||||
Ссылка в новой задаче
Block a user