Odd behavior with __constant__ memory, empty ctors, structs, and arrays

This code compiles correctly

struct MyVec4 
	__device__ MyVec4() {}
	float x, y, z, w;
struct MyMat4 { MyVec4 col0, col1, col2, col3; };
__constant__ MyMat4 myMatrix;
__constant__ MyVec4 myArrayOfVectors[4];

But altering the MyMat4 definition slightly causes the error “dynamic initialization is not supported for device, constant and shared variables”

struct MyMat4 { MyVec4 cols[4]; };

Removing the ctor from MyVec4 works around the issue; but this is not practical in the real use case, the empty ctor is necessary to denote that the class does allow default no-initialization construction.

Am I overlooking something, or are there oddities in detecting whether dynamic initialization is necessary?