What happens if you call cudaMalloc with size zero?

Accelerated Computing CUDA CUDA Programming and Performance
1

In the following code, I have a template MyClass where the template parameter Obj is itself a class that must have a device method named Func.

template <typename Obj>
class MyClass {
  public:
  __device__ void CallFunc() {
     obj.Func();
  }
  
  Obj obj;
};

MyClass<SomeClass>* buffer;
cudaMalloc(&buffer, sizeof(MyClass<SomeClass>) * 100);

As you can see above, if I try to cudaMalloc an array of MyClass (with template parameter SomeClass), it’s possible that SomeClass could have size zero because SomeClass is only required to have a method named Func and doesn’t need to have any member data. So what happens if I cudaMalloc a buffer of size 0? Is the behavior undefined? What happens if I’m in device code and I use buffer->CallFunc()? Would it work if SomeClass is size zero or is the behavior undefined?

2

According to my testing, in host code, if you call cudaMalloc with a byte size of zero (the second parameter), cudaMalloc returns a NULL pointer but not a runtime error. A NULL pointer by definition is not usable.

Note that even if you have a non-zero size, simply allocating space for an object does not make that object usable in any way. This last statement is a C++ statement, it is not unique or specific to CUDA. You must construct the object somehow. And constructing an object using a NULL pointer is non-sensical.

3

The observed behavior of cudaMalloc() is consistent with the behavior of malloc() prescribed by the ISO-C standard (which the ISO-C++ standard incorporates by reference since the function is in cstdlib):

If the size of the space requested is zero, the behavior is implementation-defined:
either a null pointer is returned, or the behavior is as if the size were some
nonzero value, except that the returned pointer shall not be used to access an object.

Returning NULL for a zero-sized allocation request is therefore not indicative of an error condition, meaning returning cudaSuccess is the logical and correct design choice.

4

sizeof(MyClass) cannot be 0

The size of a C++ class/struct object is necessarily at least 1, regardless of whether it has any members. See this StackOverflow answer.

5

From the ISO-C++11 standard, chapter 9 Classes, paragraph 4:

Complete objects and member subobjects of class type shall have nonzero size.107
[…]
107) Base class subobjects are not so constrained.

6

Just to clarify the footnote @njuffa quoted:

It doesn’t mean that if you have class B inheriting class A, then instances of A can have size 0. It means that instances of B don’t need to have extra size for being A’s, if A itself has no members. So if we have struct A {} ; and struct B : public A { int x; } then sizeof(A) will be non-zero (e.g. 1) and sizeof(B) will be sizeof(int), not sizeof(A)+sizeof(int).