[Issue]: hipGraphKernelNodeSetParams fails to copy input hipKernelNodeParams causing the runtime to later attempt to free() memory owned by the client

[IMbackK]

Problem Description

When hipGraphKernelNodeSetParams is called the hipKernelNodeParams should be copied by the runtime, but this dose not happen. Instead the hipKernelNodeParams's address is taken as is, this causes a failure later in hipGraphDestroy() as the runtime will attempt to free() the parameters pointer as passed into hipGraphKernelNodeSetParams, which could be on stack or deep inside an allocation etc and is not owned by the runtime in the first place.

Operating System

Any

CPU

Epyc 7552

GPU

MI100

ROCm Version

6.3.1

ROCm Component

No response

Steps to Reproduce

reproducer:

#include <iostream>
#include <vector>
#include <hip/hip_runtime.h>

#define HIP_CHECK(fn) { hipError_t err = fn; if(err != hipSuccess){fprintf(stderr, "Error: %s: %s at %d\n", hipGetErrorName(err), hipGetErrorString(err), __LINE__); exit(1);} }

__global__ void dummyKernel(float* out, float* in,  size_t size)
{
	int idx = blockIdx.x*blockDim.x+threadIdx.x;
	if(idx < size)
		out[idx]=2*in[idx];
}

int main()
{
	hipGraph_t graph;
	hipGraphExec_t instance;
	hipStream_t stream;

	HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking));

	constexpr size_t count = 500000;
	float* inDeviceBuffer;
	float* outDeviceBufferA;
	float* outDeviceBufferB;
	float* inHostBuffer;
	HIP_CHECK(hipMalloc(&inDeviceBuffer, sizeof(*inDeviceBuffer)*count));
	HIP_CHECK(hipMalloc(&outDeviceBufferA, sizeof(*inDeviceBuffer)*count));
	HIP_CHECK(hipMalloc(&outDeviceBufferB, sizeof(*inDeviceBuffer)*count));
	HIP_CHECK(hipHostMalloc(&inHostBuffer, sizeof(*inDeviceBuffer)*count));
	for(size_t i = 0; i < count; ++i)
		inHostBuffer[i] = i;
	HIP_CHECK(hipMemcpy(inDeviceBuffer, inHostBuffer, count*sizeof(*inDeviceBuffer), hipMemcpyHostToDevice));

	std::cout<<"outDeviceBufferA: "<<outDeviceBufferA<<'\n';
	std::cout<<"outDeviceBufferB: "<<outDeviceBufferB<<'\n';

	HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeRelaxed));
	for(size_t i = 0; i < 500; i++)
		dummyKernel<<<count/1024+1, 1024, 0, stream>>>(inDeviceBuffer, outDeviceBufferA, count);
	HIP_CHECK(hipStreamEndCapture(stream, &graph));
	HIP_CHECK(hipGraphInstantiate(&instance, graph, NULL, NULL, 0));
	HIP_CHECK(hipGraphLaunch(instance, stream));
	HIP_CHECK(hipStreamSynchronize(stream));

	size_t num_nodes;
	HIP_CHECK(hipGraphGetNodes(graph, nullptr, &num_nodes));
	std::cout<<"Graph has "<<num_nodes<<" nodes\n";

	std::vector<hipGraphNode_t> nodes(num_nodes);
	std::vector<hipKernelNodeParams> params(num_nodes);

	HIP_CHECK(hipGraphGetNodes(graph, nodes.data(), &num_nodes));

	for (size_t i = 0; i < num_nodes; i++)
	{
		hipGraphNodeType node_type;
		HIP_CHECK(hipGraphNodeGetType(nodes[i], &node_type));
		if(node_type == hipGraphNodeTypeKernel)
			HIP_CHECK(hipGraphKernelNodeGetParams(nodes[i], &params[i]));
	}

	constexpr int intervention_node = 250;
	hipGraphNodeType node_type;
	HIP_CHECK(hipGraphNodeGetType(nodes[intervention_node], &node_type));
	if(node_type == hipGraphNodeTypeKernel)
	{
		std::cout<<"node "<<intervention_node<<" is a kernel node\n";
	}
	else
	{
		std::cout<<"node "<<intervention_node<<" is not a kernel node!!\n";
		return 1;
	}

	void** updated_output_buffer_location_ptr = (void**)&outDeviceBufferB;
	std::cout<<"Original output buffer: "<<*(void**)params[intervention_node].kernelParams[1]<<'\n';
	params[intervention_node].kernelParams[1] = updated_output_buffer_location_ptr;
	HIP_CHECK(hipGraphKernelNodeSetParams(nodes[intervention_node], &params[intervention_node]));
	std::cout<<"Replaced output buffer: "<<*(void**)params[intervention_node].kernelParams[1]<<'\n';

	hipGraphNode_t errorNode;
	hipGraphExecUpdateResult result_info;
	HIP_CHECK(hipGraphExecUpdate(instance, graph, &errorNode, &result_info));
	HIP_CHECK(hipGraphLaunch(instance, stream));
	HIP_CHECK(hipStreamSynchronize(stream));

	HIP_CHECK(hipGraphExecDestroy(instance));
	HIP_CHECK(hipGraphDestroy(graph));
	return 0;
}

result:

% ./graphRepo
outDeviceBufferA: 0x78158c800000
outDeviceBufferB: 0x78158c400000
Graph has 500 nodes
node 250 is a kernel node
Original output buffer: 0x78158c800000
Replaced output buffer: 0x78158c400000
free(): invalid pointer

(Optional for Linux users) Output of /opt/rocm/bin/rocminfo --support

No response

Additional Information

No response

[ppanchad-amd]

Hi @IMbackK. Internal ticket has been created to investigate this issue. Thanks!

[schung-amd]

Hi @IMbackK, thanks for pointing this out. I think this is only partially related to hipGraphKernelNodeSetParams; commenting out that line still results in an invalid free, and as far as I can tell this should be making a copy (see: 1,2). Rather, the issue occurs because hipGraphKernelNodeGetParams gives a pointer to the node params:

clr/hipamd/src/hip_graph_internal.hpp

Line 1193 in c3f49c8

void GetParams(hipKernelNodeParams* params) { *params = kernelParams_; }

so params[intervention_node].kernelParams[1] = updated_output_buffer_location_ptr; modifies the node params directly.

From the CUDA docs, "The kernelParams or extra array returned in pNodeParams, as well as the argument values it points to, are owned by the node. This memory remains valid until the node is destroyed or its parameters are modified, and should not be modified directly." This sounds like it's in line with our implementation, but from your linked PR I assume this is working differently on CUDA? If so I'll reach out to the internal team and see if this is something we've overlooked. It could be that CUDA is performing a deeper copy than we are in their GraphKernelNodeSetParams, which would mask this issue.

As an aside, this also means you don't need to allocate new memory for the workaround; although the CUDA specs advise against doing so, you can directly assign the value with *(void**)params[intervention_node].kernelParams[1] = *(void**)updated_output_buffer_location_ptr; which will not cause an invalid free.

[IMbackK]

From the CUDA docs, "The kernelParams or extra array returned in pNodeParams, as well as the argument values it points to, are owned by the node. This memory remains valid until the node is destroyed or its parameters are modified, and should not be modified directly." This sounds like it's in line with our implementation, but from your linked PR I assume this is working differently on CUDA? If so I'll reach out to the internal team and see if this is something we've overlooked. It could be that CUDA is performing a deeper copy than we are in their GraphKernelNodeSetParams, which would mask this issue.

Indeed the reproducer works fine on cuda, because GraphKernelNodeSetParams performs a deep copy.
That said I dont see how the workflow in the reproducer could work without there being a memory leak somewhere.

I am also now puzzled how this interface is supposed to be used, given the recommendation not to modify the parameters in place. The only way i can see this working without a leak anywhere and without modifying the parameters in place would be:

1. for the user to make a deep copy by hand of the parameters recived via GraphKernelNodeGetParams
2. the runtime to destroy the original parameters when they are replaced by the call to GraphKernelNodeSetParams
3. GraphKernelNodeSetParams doing a deep copy of the supplied parameters
4. the user also destroying his own copy.

Maybe i am missing something.

As an aside, this also means you don't need to allocate new memory for the workaround; although the CUDA specs advise against doing so, you can directly assign the value with *(void**)params[intervention_node].kernelParams[1] = *(void**)updated_output_buffer_location_ptr; which will not cause an invalid free.

Yes that would indeed be a better workaround given the above parameters.

[schung-amd]

I am also now puzzled how this interface is supposed to be used, given the recommendation not to modify the parameters in place

Agreed, from a reading of the CUDA docs there doesn't seem to be a recommended way to modify the node parameters without constructing the params object from scratch as you describe. https://developer.nvidia.com/blog/constructing-cuda-graphs-with-dynamic-parameters/ reinforces this, suggesting that graphs be built by hand to begin with if dynamic node parameters are desired. In that case the user would have their own copies of the params objects which can be modified freely.

I'll check in with the internal team regarding the copy depth of GraphKernelNodeSetParams and also check if direct assignment is permissible in this context. It may be the case that direct assignment is simply dangerous rather than non-functional, in which case you could avoid the deep copy on both ROCm and CUDA.

[schung-amd]

@IMbackK Sorry for the delay. I've spoken with the internal team and we don't plan to change the copy depth here, as we follow the public CUDA docs and this is not precisely specified there. Let me know if there's something in the CUDA docs I've missed that specifies how deep the copy should be, and I can forward that to the team.

Regarding the safety of directly modifying the node parameter values against the CUDA spec recommendation, it is safe to do so today but we can't guarantee that this will be safe in the future, so we recommend avoiding doing this if possible. I understand making the parameters object by hand is not ideal, but there doesn't seem to be a better option at this point in time.

I'll close this for now, but feel free to continue discussion here if you'd like additional guidance. If there's something to be done on our end we can reopen if necessary.

[IMbackK]

hi @schung-amd thank you for your efforts in this regard, i agree that the current implementation is fine in accordance to the cuda docs.

I do think that NVIDIA made a pretty poor decision here. In the interest of improving on their work you could consider adding a function to perform a deep copy of a hipKernelNodeParams struct, let say hipCopyKernelNodeParamsExt and hipFreeKernelNodeParamsExt.

This way the workflow could be improved without sacrificing any compatibility with cuda code. The workflow would then be:

the user gets the struct via hipGraphKernelNodeGetParams.
the user makes a local copy with hipCopyKernelNodeParamsExt and modifies it to suit.
the user updates the node with his modified params via hipGraphKernelNodeSetParams
the user deletes his extra copy via hipFreeKernelNodeParamsExt