18

There is this answer on another question about the use of cudaMalloc((void**)&device_array, num_bytes), which uses void** as output argument instead of passing a void* as return value like the standard malloc.

It criticizes NVIDIA's API and states :

Casting, as in (void**)&device_array, is invalid C and results in undefined behavior.

and has been upvoted several times (8 as of now), so I assume there is some truth in it.

I don't understand what's wrong with casting there.

  • What is invalid C here?
  • In what case would this lead to undefined behavior?

All I know is that it compiles without warning and runs with the intended behavior for me. But I am not knowledgeable with C up to standard specification level.

Community
  • 1
  • 1
bct
  • 285
  • 2
  • 11
  • Same problem as if you cast `int*` to `float*`. `void**` is not special in the same way `void*` is. – user253751 Jun 07 '16 at 05:36
  • Then, is it incorrect to write this? `int val = 0x4229B26C; printf("the float represented in memory by 0x%X has a decimal value of %.03f\n", val, *((float*)&val));` – bct Jun 07 '16 at 06:01
  • Take a look at [this SO post](http://stackoverflow.com/questions/15818906/does-this-pointer-casting-break-strict-aliasing-rule) – LPs Jun 07 '16 at 06:12
  • @bct: yes the code in your comment is incorrect. But the code in the question is very different, as there is a cast, but not an access to the memory pointed to by the address, – atturri Jun 07 '16 at 06:35
  • @atturri The access is inside `cudaMalloc`. – user253751 Jun 07 '16 at 06:55

1 Answers1

11

The problem is that void* has a special meaning in C, with special rules (1). It is the only pointer type to/from which you can safely convert any other pointer type. However, these special rules do not apply recursively to void**.

Meaning that code like int* ptr = malloc(x); is perfectly fine, but

int* ptr; 
cudaMalloc(&ptr, x); // bad

is not fine! A pointer conversion from int** to void** is not well-defined. In theory this could cause undefined behavior and misalignment (2).

In addition, there might also be problems with pointer aliasing. The compiler is free to assume that the contents of a void* is never accessed through a int** and could therefore optimize the code in unexpected ways, leading to undefined behavior for violation of the strict aliasing rule (6.5).

Which means you will have to write code like this in order to safely use the function:

void* vptr; 
int*  iptr;

cudaMalloc(&vptr, x);
iptr = vptr;

(1) C11 6.3.2.3/1:

A pointer to void may be converted to or from a pointer to any object type. A pointer to any object type may be converted to a pointer to void and back again; the result shall compare equal to the original pointer.

(2) C11 6.3.2.3/7:

A pointer to an object type may be converted to a pointer to a different object type. If the resulting pointer is not correctly aligned for the referenced type, the behavior is undefined.

Lundin
  • 195,001
  • 40
  • 254
  • 396
  • I don't really see where the API is horrible thing originated. It seems that all the functions that take the device_array take ``void*``, so there is no need to use ``int*`` anywhere. Probably the example code of the other question is the only horrible thing. – atturri Jun 07 '16 at 06:42
  • @atturri `void*` is not a meaningful type by its own. Somewhere in the program, there will be a conversion to the intended type. But if that conversion is always hidden behind an abstraction layer, the API might be fine. I don't really know anything about Cuda so I can't tell. I'll remove my comment about that, since it's subjective without any context given. – Lundin Jun 07 '16 at 06:47
  • @Lundin Nice answer. About atturi's comment, in the case of cuda memory allocation void* should never be converted to something else, at least on host code. Actually it points to an address in graphics memory that is not accessible by the CPU... So after thinking about it, I'd argue that the problem is just that the type of device_array should be void* and not int*? Then everything would be fine?.... – bct Jun 07 '16 at 07:03
  • Hmm, also, `void *` and `int *` might have different representations on memory; which would also mean that Grzegorz's answer is wrong. – Antti Haapala -- Слава Україні Jun 07 '16 at 07:09
  • @bct Yeah but the problem is they need 2 return values. If the contents of the `void*` is never accessed by the caller or converted to any other type, it should be ok. But in that case the API is still questionable, const correctness would mean they should have used `const void**`. – Lundin Jun 07 '16 at 07:41
  • also, following the explanation, would cudaMalloc(&(void*)device_array, num_bytes) be legal then? (as all the casts go through 'void*' this way) – bct Jun 07 '16 at 07:51
  • @Lundin: If they use ``const void**`` then the application will have to keep a ``const void*`` variable, which means that they will not be able to pass it to some (I'm making this up) API function like ``cudaWriteTo(void* pDest, [...])``. If the user has a ``void*`` variable, they should know they should never dereference it, neither for reading nor for writing. – atturri Jun 07 '16 at 09:41
  • While most of this answer is good and correct, I find the first paragraph confusing and irrelevant to what's actually wrong with the bad cast. It's accessing the object via the wrong type that's the problem. Yes `void *` is special in that it can represent any pointer and implicitly converts, but you never access an object of type `void` via `void *` because objects of type `void` don't exist. On the other hand, objects of type `void *` do exist, and they're distinct from objects of type `int *`. – R.. GitHub STOP HELPING ICE Jun 07 '16 at 12:08
  • I agree with R that this answer is at least confusing, if not incorrect. The use of the second reference to suggest that double pointers may present a problem is simply incorrect. A `void**` pointer is a pointer to a pointer. An `int**` pointer is a pointer to a pointer. On a given platform, pointer sizes are the same, and therefore a pointer to a pointer can never be misaligned, even if cast from one type to another. – Robert Crovella Jun 07 '16 at 13:13
  • The underlying single pointer can be misaligned (in the general case) but this is true generically of C (as the second reference suggests) and not uniquely arising when we "recursively" take the address of a pointer. – Robert Crovella Jun 07 '16 at 13:13
  • 1
    Considering that NVIDIA's official example code consistently demos `cudaMalloc(&non_void_ptr, size)`, I will assume that is is well defined for [nvcc](https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/). It's not like it will be portable code for other GPU compilers anyway. – Matt Eding Sep 14 '21 at 17:54