HIP Green Context Implementation by manocharahul · Pull Request #3253 · ROCm/rocm-systems

manocharahul · 2026-02-12T21:17:14Z

Header File updated with structs and declaration
Added CUDA mappings to hipother
Added Basic Sanity test for each API

Motivation

Technical Details

JIRA ID

Test Plan

Test Result

Submission Checklist

Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

- Header File updated with structs and declaration - Added CUDA mappings to hipother - Added Basic Sanity test for each API

Copilot

Pull request overview

Implements the HIP “Green Context” API surface across HIP headers and the AMD/NVIDIA backends, wiring through dispatch/export/profiling infrastructure.

Changes:

Added new Green Context typedefs and public runtime API declarations.
Added AMD-side dispatch table plumbing, symbol exports, and API tracing/profiling IDs.
Added NVIDIA (hipnv) inline mappings to corresponding CUDA driver APIs plus an AMD stub implementation source.

Reviewed changes

Copilot reviewed 11 out of 11 changed files in this pull request and generated 7 comments.

Show a summary per file

File	Description
projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h	Adds hipnv typedefs and CUDA driver API mappings for Green Context calls
projects/hip/include/hip/hip_runtime_api.h	Exposes new Green Context types and function declarations in the public HIP runtime API
projects/clr/hipamd/src/hip_table_interface.cpp	Adds runtime entrypoints that forward Green Context calls through the dispatch table
projects/clr/hipamd/src/hip_hcc.map.in	Exports new Green Context symbols (version script)
projects/clr/hipamd/src/hip_greenctx.hpp	Introduces a GreenCtx class header and includes for the new feature area
projects/clr/hipamd/src/hip_greenctx.cpp	Adds AMD implementation stubs for the Green Context APIs
projects/clr/hipamd/src/hip_api_trace.cpp	Wires Green Context APIs into dispatch table initialization and ABI enforcement
projects/clr/hipamd/src/amdhip.def	Exports new Green Context functions on Windows (DEF file)
projects/clr/hipamd/src/CMakeLists.txt	Builds the new hip_greenctx.cpp source
projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h	Adds tracer/profiler IDs, args structs, and string formatting for new APIs
projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp	Extends dispatch table types and bumps runtime API table step version

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

Copilot · 2026-02-12T21:18:46Z

projects/clr/hipamd/src/hip_hcc.map.in

+hip_7.3 {
+global:
+    hipGreenCtxCreate;
+    hipGreenCtxDestroy;
+    hipGreenCtxStreamCreate;
+    hipStreamGetGreenCtx;
+    hipGreenCtxRecordEvent;
+    hipGreenCtxWaitEvent;
+    hipCtxFromGreenCtx;
+local:
+    *;
+} hip_7.2;


The version script defines hip_7.3 but declares it inherits from hip_7.2 (and closes with } hip_7.2;). In the shown file, the immediately preceding version is hip_7.1, so inheriting from hip_7.2 (and/or closing with hip_7.2) looks incorrect and can break symbol versioning at link time. Update the inherited version to the correct prior version (e.g., } hip_7.1;) or introduce a proper hip_7.2 block if that’s intended.

Copilot · 2026-02-12T21:18:47Z

projects/clr/hipamd/src/hip_greenctx.cpp

+  if (ctx != nullptr) {
+    *ctx = nullptr;
+  }
+
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGreenCtxDestroy(hipGreenCtx_t ctx) {
+  HIP_INIT_API(hipGreenCtxDestroy, ctx);
+
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGreenCtxStreamCreate(hipStream_t* stream, hipGreenCtx_t greenctx, unsigned int flags,
+                                   int priority) {
+  HIP_INIT_API(hipGreenCtxStreamCreate, stream, greenctx, flags, priority);
+
+  if (stream != nullptr) {
+    *stream = nullptr;
+  }
+
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipStreamGetGreenCtx(hipStream_t hStream, hipGreenCtx_t* greenCtx) {
+  HIP_INIT_API(hipStreamGetGreenCtx, hStream, greenCtx);
+
+  if (greenCtx != nullptr) {
+    *greenCtx = nullptr;
+  }
+
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGreenCtxRecordEvent(hipGreenCtx_t greenCtx, hipEvent_t event) {
+  HIP_INIT_API(hipGreenCtxRecordEvent, greenCtx, event);
+
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGreenCtxWaitEvent(hipGreenCtx_t greenCtx, hipEvent_t event) {
+  HIP_INIT_API(hipGreenCtxWaitEvent, greenCtx, event);
+
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipCtxFromGreenCtx(hipCtx_t* ctx, hipGreenCtx_t greenCtx) {
+  HIP_INIT_API(hipCtxFromGreenCtx, ctx, greenCtx);
+
+  if (ctx != nullptr) {
+    *ctx = nullptr;
+  }
+
+  HIP_RETURN(hipSuccess);


These AMD implementations are hard-coded to return hipSuccess while producing no functional result (e.g., hipGreenCtxCreate sets the output context to nullptr). This is observable incorrect behavior for callers because it reports success but does not create a usable context. If the feature is not supported yet, return an appropriate error (commonly hipErrorNotSupported) and avoid signaling success; if it is supported, implement creation/destruction/stream association semantics so that success implies a valid, usable handle.

Suggested change

if (ctx != nullptr) {

*ctx = nullptr;

}

HIP_RETURN(hipSuccess);

}

hipError_t hipGreenCtxDestroy(hipGreenCtx_t ctx) {

HIP_INIT_API(hipGreenCtxDestroy, ctx);

HIP_RETURN(hipSuccess);

}

hipError_t hipGreenCtxStreamCreate(hipStream_t* stream, hipGreenCtx_t greenctx, unsigned int flags,

int priority) {

HIP_INIT_API(hipGreenCtxStreamCreate, stream, greenctx, flags, priority);

if (stream != nullptr) {

*stream = nullptr;

}

HIP_RETURN(hipSuccess);

}

hipError_t hipStreamGetGreenCtx(hipStream_t hStream, hipGreenCtx_t* greenCtx) {

HIP_INIT_API(hipStreamGetGreenCtx, hStream, greenCtx);

if (greenCtx != nullptr) {

*greenCtx = nullptr;

}

HIP_RETURN(hipSuccess);

}

hipError_t hipGreenCtxRecordEvent(hipGreenCtx_t greenCtx, hipEvent_t event) {

HIP_INIT_API(hipGreenCtxRecordEvent, greenCtx, event);

HIP_RETURN(hipSuccess);

}

hipError_t hipGreenCtxWaitEvent(hipGreenCtx_t greenCtx, hipEvent_t event) {

HIP_INIT_API(hipGreenCtxWaitEvent, greenCtx, event);

HIP_RETURN(hipSuccess);

}

hipError_t hipCtxFromGreenCtx(hipCtx_t* ctx, hipGreenCtx_t greenCtx) {

HIP_INIT_API(hipCtxFromGreenCtx, ctx, greenCtx);

if (ctx != nullptr) {

*ctx = nullptr;

}

HIP_RETURN(hipSuccess);

HIP_RETURN(hipErrorNotSupported);

}

hipError_t hipGreenCtxDestroy(hipGreenCtx_t ctx) {

HIP_INIT_API(hipGreenCtxDestroy, ctx);

HIP_RETURN(hipErrorNotSupported);

}

hipError_t hipGreenCtxStreamCreate(hipStream_t* stream, hipGreenCtx_t greenctx, unsigned int flags,

int priority) {

HIP_INIT_API(hipGreenCtxStreamCreate, stream, greenctx, flags, priority);

HIP_RETURN(hipErrorNotSupported);

}

hipError_t hipStreamGetGreenCtx(hipStream_t hStream, hipGreenCtx_t* greenCtx) {

HIP_INIT_API(hipStreamGetGreenCtx, hStream, greenCtx);

HIP_RETURN(hipErrorNotSupported);

}

hipError_t hipGreenCtxRecordEvent(hipGreenCtx_t greenCtx, hipEvent_t event) {

HIP_INIT_API(hipGreenCtxRecordEvent, greenCtx, event);

HIP_RETURN(hipErrorNotSupported);

}

hipError_t hipGreenCtxWaitEvent(hipGreenCtx_t greenCtx, hipEvent_t event) {

HIP_INIT_API(hipGreenCtxWaitEvent, greenCtx, event);

HIP_RETURN(hipErrorNotSupported);

}

hipError_t hipCtxFromGreenCtx(hipCtx_t* ctx, hipGreenCtx_t greenCtx) {

HIP_INIT_API(hipCtxFromGreenCtx, ctx, greenCtx);

HIP_RETURN(hipErrorNotSupported);

Copilot · 2026-02-12T21:18:47Z

projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp

+  t_hipCtxFromGreenCtx hipCtxFromGreenCtx_fn;
+
+  // DO NOT EDIT ABOVE!
+  // HIP_RUNTIME_API_TABLE_STEP_VERSION == 24


The dispatch table section ends with a comment suggesting HIP_RUNTIME_API_TABLE_STEP_VERSION == 24, but the macro was bumped to 23 and this PR adds entries for step 23. Leaving a == 24 marker here is misleading for future edits and can cause confusion when reconciling ABI/versioning changes. Recommend removing the // HIP_RUNTIME_API_TABLE_STEP_VERSION == 24 line (or updating it only when the step version actually advances).

Suggested change

// HIP_RUNTIME_API_TABLE_STEP_VERSION == 24

Copilot · 2026-02-12T21:18:48Z

projects/clr/hipamd/src/hip_greenctx.hpp

+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 WARRANNTY OF ANY KIND, EXPRESS OR


Correct the spelling of 'WARRANNTY' to 'WARRANTY' in the license header.

Suggested change

THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR

THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR

Copilot · 2026-02-12T21:18:48Z

projects/clr/hipamd/src/hip_greenctx.hpp

+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE


Correct the spelling of 'FITNNESS' to 'FITNESS' in the license header.

Suggested change

FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE

FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE

Copilot · 2026-02-12T21:18:48Z

projects/clr/hipamd/src/hip_greenctx.hpp

+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN


Correct the spelling of 'INN' to 'IN' in the license header (appears multiple times).

Suggested change

LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,

OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN

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

Copilot · 2026-02-12T21:18:49Z

projects/clr/hipamd/src/hip_greenctx.hpp

+        GreenCtx();
+        ~GreenCtx();


The GreenCtx class declares a constructor/destructor but this PR doesn’t provide definitions in the added .cpp (which only contains the C API stubs). If this class is intended to be used, provide definitions (or default them inline) and document its role; if it’s not needed yet, consider removing it to avoid carrying an incomplete/unused abstraction.

Suggested change

GreenCtx();

~GreenCtx();

// Lightweight context object for future green threading / context management.

// Currently has no state; ctor/dtor are defaulted and do not perform any work.

GreenCtx() = default;

~GreenCtx() = default;

HIP Green Context Implementation

4b375ba

- Header File updated with structs and declaration - Added CUDA mappings to hipother - Added Basic Sanity test for each API

Copilot AI review requested due to automatic review settings February 12, 2026 21:17

manocharahul requested review from a team as code owners February 12, 2026 21:17

manocharahul marked this pull request as draft February 12, 2026 21:17

github-actions bot added project: clr project: hip project: hipother labels Feb 12, 2026

Copilot AI reviewed Feb 12, 2026

View reviewed changes

systems-assistant bot added the organization: ROCm label Feb 12, 2026

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

HIP Green Context Implementation#3253

HIP Green Context Implementation#3253
manocharahul wants to merge 1 commit intodevelopfrom
amd/dev/rmanocha/green-ctx-api

manocharahul commented Feb 12, 2026

Uh oh!

Copilot AI left a comment

Uh oh!

Copilot AI Feb 12, 2026

Uh oh!

Copilot AI Feb 12, 2026

Uh oh!

Copilot AI Feb 12, 2026

Uh oh!

Copilot AI Feb 12, 2026

Uh oh!

Copilot AI Feb 12, 2026

Uh oh!

Copilot AI Feb 12, 2026

Uh oh!

Copilot AI Feb 12, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

	THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
	THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR

	FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
	FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE

		LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
		OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN

-        GreenCtx();
-        ~GreenCtx();
+        // Lightweight context object for future green threading / context management.
+        // Currently has no state; ctor/dtor are defaulted and do not perform any work.
+        GreenCtx() = default;
+        ~GreenCtx() = default;

Conversation

manocharahul commented Feb 12, 2026

Motivation

Technical Details

JIRA ID

Test Plan

Test Result

Submission Checklist

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

Copilot AI Feb 12, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI Feb 12, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI Feb 12, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI Feb 12, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI Feb 12, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI Feb 12, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI Feb 12, 2026

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant