Matrix/Matrix multiplication is a computation that requires an evaluation as it relies on BLAS kernels. For that reason assign() is overloaded by special functions whenever an evaluation is required, but the overload that prevails remains the one for CUDADynamicMatrix.

This is a blocking feature for Blaze CUDA, it has got all my attention recently.

The issue here is that the workflow of smpAssign() is different than the one I expected. I might have to change the whole approach for Blaze CUDA, I've been thinking about making a separate cudaAssign() function. Klaus Iglberger suggested that I could overload DMatDMatMultExpr but it would only solve the problem for that specific computation, and I'd like it to be solved properly for all computations.

The problem however is that cudaAssign() would be external to the expressions and would not have access to the private type traits so I might need an additional type traits system if I take that direction. I'll give it a shot and see how that goes.

Performance being a main feature of Blaze, it has to be for Blaze CUDA too. Therefore benchmarking will be necessary at some point to make sure we reach that goal.

Adding type traits for better integration with the original Blaze. This would allow us to make sure the right assign functions are called given the operands' types.

This part will require a lot of attention, type traits are rather easy to implement but are structural elements to the library.

Error management macro is being worked on, most of the work will consist in making sure all CUDA runtime errors are handled.

We definitely need them. These should be easy to get from the original Blaze.

This might be done once the code gets a nice refactoring. The structure might be subject to change, so this issue will be on hold for now.

blaze::CUDAReduce doesn't work for large sizes. I've been unable to find the source of the bug for days now and I'm running out of ideas.

Above a certain threshold the CUDA reduce kernel (the __global__ function) will start outputting inaccurate values. I've been trying to pinpoint the issue, to add synchronization directives but nothing seems to help.

Referring to STEllAR-GROUP/blaze#1

Most of the work is done now, cudaAssign() needs overloads for every expression to support partial evaluation properly, following the same implementation pattern as in DMatDMatAddExpr.h:

• External to the original expression templates
• Implement the same functionalities as their CPU counterparts
• Follow the same enable condition as their CPU counterparts
• Call cudaAssign() instead of assign()

cuBLAS will be used as much as possible to implement them.