diff --git a/projects/clr/CHANGELOG.md b/projects/clr/CHANGELOG.md index f35ca829a0..ba1fb40610 100644 --- a/projects/clr/CHANGELOG.md +++ b/projects/clr/CHANGELOG.md @@ -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 diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp b/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp index 051ee774cc..499f131518 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp @@ -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 // ******************************************************************************************* // // diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h b/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h index d9b3941439..390d5fe273 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h @@ -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("; diff --git a/projects/clr/hipamd/src/amdhip.def b/projects/clr/hipamd/src/amdhip.def index 89db8a4ef5..008d684a40 100644 --- a/projects/clr/hipamd/src/amdhip.def +++ b/projects/clr/hipamd/src/amdhip.def @@ -516,3 +516,4 @@ hipLibraryLoadFromFile hipLibraryUnload hipLibraryGetKernel hipLibraryGetKernelCount +hipStreamCopyAttributes diff --git a/projects/clr/hipamd/src/hip_api_trace.cpp b/projects/clr/hipamd/src/hip_api_trace.cpp index c2db14f15f..5e782edd82 100644 --- a/projects/clr/hipamd/src/hip_api_trace.cpp +++ b/projects/clr/hipamd/src/hip_api_trace.cpp @@ -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(, , 8) // // HIP_ENFORCE_ABI_VERSIONING(
, 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 diff --git a/projects/clr/hipamd/src/hip_hcc.map.in b/projects/clr/hipamd/src/hip_hcc.map.in index c9fecec49c..c7974006da 100644 --- a/projects/clr/hipamd/src/hip_hcc.map.in +++ b/projects/clr/hipamd/src/hip_hcc.map.in @@ -638,6 +638,7 @@ global: hipLibraryUnload; hipLibraryGetKernel; hipLibraryGetKernelCount; + hipStreamCopyAttributes; local: *; } hip_7.1; \ No newline at end of file diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp index f2081b97f8..ec178d4bea 100644 --- a/projects/clr/hipamd/src/hip_stream.cpp +++ b/projects/clr/hipamd/src/hip_stream.cpp @@ -878,7 +878,8 @@ hipError_t hipStreamSetAttribute(hipStream_t stream, hipStreamAttrID attr, HIP_RETURN(hipErrorStreamCaptureUnsupported); } - hip::Stream* s = reinterpret_cast(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(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 diff --git a/projects/clr/hipamd/src/hip_table_interface.cpp b/projects/clr/hipamd/src/hip_table_interface.cpp index 90d27455f6..0143f9deef 100644 --- a/projects/clr/hipamd/src/hip_table_interface.cpp +++ b/projects/clr/hipamd/src/hip_table_interface.cpp @@ -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); } diff --git a/projects/hip-tests/catch/unit/stream/CMakeLists.txt b/projects/hip-tests/catch/unit/stream/CMakeLists.txt index bdaa5050c9..4ef76f368b 100644 --- a/projects/hip-tests/catch/unit/stream/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/stream/CMakeLists.txt @@ -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 diff --git a/projects/hip-tests/catch/unit/stream/hipStreamCopyAttributes.cc b/projects/hip-tests/catch/unit/stream/hipStreamCopyAttributes.cc new file mode 100644 index 0000000000..f3d9672b5d --- /dev/null +++ b/projects/hip-tests/catch/unit/stream/hipStreamCopyAttributes.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 + +/** + * 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)); +} diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h index d354e0b8e4..88fcf6cea9 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -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 /** * @} diff --git a/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h b/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h index 9cc9975a62..f775ea8868 100644 --- a/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h +++ b/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h @@ -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)); } diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/enum_string.hpp b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/enum_string.hpp index b998fe029c..4c9ed13717 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/enum_string.hpp +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/enum_string.hpp @@ -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) diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/api_args.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/api_args.h index 1ecca43ad6..012b709941 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/api_args.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/api_args.h @@ -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 diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/runtime_api_id.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/runtime_api_id.h index 7781669996..b0f6314e99 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/runtime_api_id.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/runtime_api_id.h @@ -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; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/abi.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/abi.cpp index d80304a87e..ea0b0931ad 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/abi.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/abi.cpp @@ -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 diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.def.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.def.cpp index 03590a1421..f14a55609b 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.def.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.def.cpp @@ -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 diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/stream.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/stream.cpp index 1d04dd1e41..55f9641b1c 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/stream.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/stream.cpp @@ -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 : std::true_type \ + { \ + static constexpr auto reason = REASON; \ + }; + +template +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 void update_table(Tp* _orig, std::integral_constant) { - using table_type = typename hip_table_lookup::type; - using info_type = hip_api_info; - using function_args_type = decltype(info_type::get_args_type()); + using table_type = typename hip_table_lookup::type; + using info_type = hip_api_info; + using explicit_disable_type = explicit_disable_update; + 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) constexpr auto num_args = function_args_type::size(); - if constexpr(common::mpl::is_one_of::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::value) { constexpr auto stream_idx = common::mpl::index_of::value;