I have a minor API suggestion. Currently to register a host pointer for mapping it has to be convertible to void * (i.e. non-const). This allows the information at the pointer to be both read and written to. However, if you know you only want to read the information from a pointer, it seems that we should be able to register a constant host pointer.
Take for example:
const std::vector<int> a = {0,1,2,3,4,5};
int * d_a;
cudaMalloc((void**)d_a),6*sizeof(int));
cudaHostRegister(a.data(),6*sizeof(int),cudaHostRegisterDefault);
cudaMemcpyAsync(d_a, a, 6*sizeof(int), cudaMemcpyHostToDevice);
cudaUnregister(a.data());
//...
//do work on d_a
//...
int * h_a_transformed = new int[6];
cudaHostRegister(h_a_transformed,6*sizeof(int),cudaHostRegisterDefault);
cudaMemcpyAsync(h_a_transformed, d_a, 6*sizeof(int), cudaMemcpyDeviceToHost);
cudaUnregister(h_a_transformed);
cudaDeviceSynchronize();
This will fail at line 2 because a.data() returns a const int * which cannot be converted to void*. Now I don’t have to use Host registered memory or cudaMemcpyAsync (particularly for this example). I could just pass a.data() to a regular cudaMemcpy. But it would be nice to be able to pin memory labeled as const int * for performance …
I’m not sure if there is a reason why this cannot be done or if this is an overlooked corner case. If the latter, it would be nice if an overload was provided in the runtime API (and of course the corresponding cuMemHostRegister for the driver API and unregister functions for each).
If you’d like to see a change to CUDA, the best approach is to file a bug at developer.nvidia.com, and mark the title of the bug with the RFE keyword (Request For Enhancement) which will let folks know that you’re requesting a change, rather than reporting a functional defect.
Some thoughts occur to me as I look at what you have shown, but it’s unclear what your focus is, whether its just the const-qualified character or the desire to register these types of things. Your code appears to be an idea sketch, so I’ll leave it at that.
I apologize that my question is unclear and that my example made it worse: the request for enhancement is that I would like to be able to register/un-register const-qualified host pointers. The example above is meant to show a use-case where the user would receive a const T* from an API, tries to register it, and failing because cudaHostRegister only takes in a void*. The solution without breaking const-ness is to have overloads of the register functions that take in a const void *.
Note: for my own work, I can work around this easily enough because I’m the one designing the relevant API! However, I thought for the general use case, being able to un/register const T* pointers, without breaking their const-ness, could be useful if possible.
One side-point you brought up before was:
That’s actually a really good question. I’ve tested it a bunch times and it seems to “work”, but I don’t know if I am just getting lucky or if there is a guarantee of correctness whereby unregistering the host pointer after an async memcpy call will still guarantee that the memcpy will complete correctly. I would think/hope it would, but it would be nice to get confirmation as the documentation doesn’t really specify.