19#ifndef TMVA_DNN_ARCHITECTURES_CUDA_CUDAMATRIX
20#define TMVA_DNN_ARCHITECTURES_CUDA_CUDAMATRIX
23#include "RConfigure.h"
24#ifdef R__HAS_STD_STRING_VIEW
25#undef R__HAS_STD_STRING_VIEW
26#define R__HAS_STD_EXPERIMENTAL_STRING_VIEW
30#include "cuda_runtime.h"
32#include "curand_kernel.h"
37#define CUDACHECK(ans) {cudaError((ans), __FILE__, __LINE__); }
45inline void cudaError(cudaError_t code,
const char *
file,
int line,
bool abort=
true);
60template<
typename AFloat>
104template<
typename AFloat>
203 if (code != cudaSuccess)
205 fprintf(stderr,
"CUDA Error: %s %s %d\n", cudaGetErrorString(code),
file,
line);
206 if (abort) exit(code);
211template<
typename AFloat>
213 : fDevicePointer(devicePointer)
219template<
typename AFloat>
223 cudaMemcpy(& buffer, fDevicePointer,
sizeof(AFloat),
224 cudaMemcpyDeviceToHost);
229template<
typename AFloat>
233 cudaMemcpyDeviceToDevice);
237template<
typename AFloat>
240 AFloat buffer = value;
241 cudaMemcpy(fDevicePointer, & buffer,
sizeof(AFloat),
242 cudaMemcpyHostToDevice);
246template<
typename AFloat>
250 cudaMemcpy(& buffer, fDevicePointer,
sizeof(AFloat),
251 cudaMemcpyDeviceToHost);
253 cudaMemcpy(fDevicePointer, & buffer,
sizeof(AFloat),
254 cudaMemcpyHostToDevice);
258template<
typename AFloat>
262 cudaMemcpy(& buffer, fDevicePointer,
sizeof(AFloat),
263 cudaMemcpyDeviceToHost);
265 cudaMemcpy(fDevicePointer, & buffer,
sizeof(AFloat),
266 cudaMemcpyHostToDevice);
270template<
typename AFloat>
273 return fElementBuffer.GetComputeStream();
277template<
typename AFloat>
280 return fElementBuffer.SetComputeStream(stream);
284template<
typename AFloat>
288 cudaEventCreateWithFlags(&event, cudaEventDisableTiming);
289 cudaEventRecord(event,
A.GetComputeStream());
290 cudaStreamWaitEvent(fElementBuffer.GetComputeStream(), event, 0);
291 cudaEventDestroy(event);
295template<
typename AFloat>
298 AFloat buffer = value;
299 cudaMemcpy(fDeviceReturn, & buffer,
sizeof(AFloat), cudaMemcpyHostToDevice);
303template<
typename AFloat>
307 cudaMemcpy(& buffer, fDeviceReturn,
sizeof(AFloat), cudaMemcpyDeviceToHost);
312template<
typename AFloat>
315 AFloat * elementPointer = fElementBuffer;
316 elementPointer += j * fNRows + i;
void operator-=(AFloat value)
TCudaDeviceReference(AFloat *devicePointer)
void operator=(const TCudaDeviceReference &other)
void operator+=(AFloat value)
TCudaDeviceBuffer< AFloat > fElementBuffer
TCudaMatrix & operator=(const TCudaMatrix &)=default
static curandState_t * fCurandStates
TCudaMatrix(const TMatrixT< AFloat > &)
static AFloat GetDeviceReturn()
Transfer the value in the device return buffer to the host.
void SetComputeStream(cudaStream_t stream)
static AFloat * fDeviceReturn
Buffer for kernel return values.
cudaStream_t GetComputeStream() const
size_t GetNoElements() const
void InitializeCuda()
Initializes all shared devices resource and makes sure that a sufficient number of curand states are ...
static Bool_t gInitializeCurand
TCudaDeviceReference< AFloat > operator()(size_t i, size_t j) const
Access to elements of device matrices provided through TCudaDeviceReference class.
static AFloat * GetDeviceReturnPointer()
Return device pointer to the device return buffer.
const cublasHandle_t & GetCublasHandle() const
static AFloat * fOnes
Vector used for summations of columns.
static void ResetDeviceReturn(AFloat value=0.0)
Set the return buffer on the device to the specified value.
const AFloat * GetDataPointer() const
TCudaMatrix(TCudaDeviceBuffer< AFloat > buffer, size_t m, size_t n)
static size_t fNCurandStates
TCudaMatrix(const TCudaMatrix &)=default
TCudaDeviceBuffer< AFloat > GetDeviceBuffer() const
void Synchronize(const TCudaMatrix &) const
Blocking synchronization with the associated compute stream, if it's not the default stream.
static AFloat * GetOnes()
static cublasHandle_t fCublasHandle
static size_t fInstances
Current number of matrix instances.
TCudaMatrix(size_t i, size_t j)
TCudaMatrix & operator=(TCudaMatrix &&)=default
void InitializeCurandStates()
AFloat * GetDataPointer()
static size_t fNOnes
Current length of the one vector.
TCudaMatrix(TCudaMatrix &&)=default
static curandState_t * GetCurandStatesPointer()
void Print(Option_t *name="") const
Print the matrix as a table of elements.
void cudaError(cudaError_t code, const char *file, int line, bool abort=true)
Function to check cuda return code.
create variable transformations