Why does cudaMalloc() use pointer to pointer?

Question

For example, cudaMalloc((void**)&device_array, num_bytes);

This question has been asked before, and the reply was "because cudaMalloc returns an error code", but I don't get it - what has a double pointer got to do with returning an error code? Why can't a simple pointer do the job?

If I write

cudaError_t catch_status;
catch_status = cudaMalloc((void**)&device_array, num_bytes);

the error code will be put in catch_status, and returning a simple pointer to the allocated GPU memory should suffice, shouldn't it?

Answer 1

In C, data can be passed to functions by value or via simulated pass-by-reference (i.e. by a pointer to the data). By value is a one-way methodology, by pointer allows for two-way data flow between the function and its calling environment.

When a data item is passed to a function via the function parameter list, and the function is expected to modify the original data item so that the modified value shows up in the calling environment, the correct C method for this is to pass the data item by pointer. In C, when we pass by pointer, we take the address of the item to be modified, creating a pointer (perhaps a pointer to a pointer in this case) and hand the address to the function. This allows the function to modify the original item (via the pointer) in the calling environment.

Normally malloc returns a pointer, and we can use assignment in the calling environment to assign this returned value to the desired pointer. In the case of cudaMalloc, the CUDA designers chose to use the returned value to carry an error status rather than a pointer. Therefore the setting of the pointer in the calling environment must occur via one of the parameters passed to the function, by reference (i.e. by pointer). Since it is a pointer value that we want to set, we must take the address of the pointer (creating a pointer to a pointer) and pass that address to the cudaMalloc function.

Answer 2

Adding to Robert's answer, but to first reiterate, it is a C API, which means it does not support references, which would allow you to modify the value of a pointer (not just what is pointed to) inside the function. The answer by Robert Crovella explained this. Also note that it needs to be void because C also does not support function overloading.

Further, when using a C API within a C++ program (but you have not stated this), it is common to wrap such a function in a template. For example,

template<typename T>
cudaError_t cudaAlloc(T*& d_p, size_t elements)
{
    return cudaMalloc((void**)&d_p, elements * sizeof(T));
}

There are two differences with how you would call the above cudaAlloc function:

1. Pass the device pointer directly, without using the address-of operator (&) when calling it, and without casting to a void type.
2. The second argument elements is now the number of elements rather than the number of bytes. The sizeof operator facilitates this. This is arguably more intuitive to specify elements and not worry about bytes.

For example:

float *d = nullptr;  // floats, 4 bytes per elements
size_t N = 100;      // 100 elements

cudaError_t err = cudaAlloc(d,N);      // modifies d, input is not bytes

if (err != cudaSuccess)
    std::cerr << "Unable to allocate device memory" << std::endl;

Answer 3

I guess the signature of cudaMalloc function could be better explained by an example. It is basically assigning a buffer through a pointer to that buffer (a pointer to pointer), like the following method:

int cudaMalloc(void **memory, size_t size)
{
    int errorCode = 0;

    *memory = new char[size];

    return errorCode;
}

As you can see, the method takes a memory pointer to pointer, on which it saves the new allocated memory. It then returns the error code (in this case as an integer, but it is actually an enum).

The cudaMalloc function could be designed as it follows also:

void * cudaMalloc(size_t size, int * errorCode = nullptr)
{
    if(errorCode)
        errorCode = 0;

    char *memory = new char[size];

    return memory;
}

In this second case, the error code is set through a pointer implicit set to null (for the case people do not bother with the error code at all). Then the allocated memory is returned.

The first method can be used as is the actual cudaMalloc right now:

float *p;
int errorCode;
errorCode = cudaMalloc((void**)&p, sizeof(float));

While the second one can be used as follows:

float *p;
int errorCode;
p = (float *) cudaMalloc(sizeof(float), &errorCode);

These two methods are functionally equivalent, while they have different signatures, and the people from cuda decided to go for the first method, returning the error code and assigning the memory through a pointer, while most people say that the second method would have been a better choice.