Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Change request: Give A2D matrices a trivial constructor #125

Open
A-CGray opened this issue Dec 7, 2024 · 1 comment
Open

Change request: Give A2D matrices a trivial constructor #125

A-CGray opened this issue Dec 7, 2024 · 1 comment

Comments

@A-CGray
Copy link
Contributor

A-CGray commented Dec 7, 2024

I was recently trying to created an array of A2D matrices in shared memory inside a GPU kernel, e.g:

__shared__ A2D::Mat<double, numNodes, numStates> localNodeStates[elemPerBlock];

However, when I did this I got the following:

../element/FEKernels.h(738): warning #20054-D: dynamic initialization is not supported for a function-scope static __shared__ variable within a __device__/__global__ function

I think the issue is that A2D matrices do not have a trivial constructor, the default constructor sets all entries to zero, so when you try to create one of the matrices in shared memory, CUDA doesn't know which thread should call it's constructor.

Cuda does support dynamic allocation of shared memory that would get around this issue, but IMO it's a pretty nasty approach as it requires passing in the correct allocation sizes when you call the kernel.

I verified that you can successfully create A2D matrices in shared memory if you replace the current default constructor with default:

// Current constructor
// A2D_FUNCTION Mat() {
//   for (int i = 0; i < M * N; i++) {
//     A[i] = 0.0;
//   }
// }

// New constructor
A2D_FUNCTION Mat() = default;

This however has the downside that we can no longer rely on matrices being zeroed out by default, which might break things elsewhere in A2D.

@A-CGray
Copy link
Contributor Author

A-CGray commented Dec 7, 2024

@gjkennedy @sean-engelstad this is the shared memory issue I was talking about yesterday

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

No branches or pull requests

1 participant