CUDA GPU

Functions
__global__ void	cudakernel_vectorAdd (const LL_float32 *A, const LL_float32 *B, LL_float32 *C, int n)
__global__ void	kernel1 (LL_float32 *data, int n)
cudaError_t	launch_kernel1 (int nblocks, int thrPerBlk, int shared, cudaStream_t stream, LL_float32 *data, int n)
__global__ void	kernel2 (LL_float32 *data, int n)
void	cudaDevicePtrDestroyer (Ptr_t p)
void	cudaHostPtrDestroyer (Ptr_t p)
int	array_matmul_cublas_f32 (const LL_float32 *A, const LL_float32 *B, LL_float32 *C, int m, int k, int n)
int	array_matmul_cublas_f64 (const LL_float64 *A, const LL_float64 *B, LL_float64 *C, int m, int k, int n)
Sexpr_t	mop3_cpp_vector_add (Lamb &lamb, Sexpr_t sexpr, Sexpr_t env_exec)
Sexpr_t	mop3_cudaLaunch_kernel1 (Lamb &lamb, Sexpr_t sexpr, Sexpr_t env_exec)
Sexpr_t	mop3_cudakernel_vectorAdd (Lamb &lamb, Sexpr_t sexpr, Sexpr_t env_exec)
Sexpr_t	mop3_cudaMalloc (Lamb &lamb, Sexpr_t sexpr, Sexpr_t env_exec)
Sexpr_t	mop3_cudaMallocHost (Lamb &lamb, Sexpr_t sexpr, Sexpr_t env_exec)
Sexpr_t	mop3_cudaFree (Lamb &lamb, Sexpr_t sexpr, Sexpr_t env_exec)
Sexpr_t	mop3_cudaMemcpy (Lamb &lamb, Sexpr_t sexpr, Sexpr_t env_exec)
Sexpr_t	mop3_cudaMemset (Lamb &lamb, Sexpr_t sexpr, Sexpr_t env_exec)
void	cudaStreamDestroyer (Ptr_t p)
Sexpr_t	mop3_cudaStreamCreate (Lamb &lamb, Sexpr_t sexpr, Sexpr_t env_exec)
Sexpr_t	mop3_cudaStreamSynchronize (Lamb &lamb, Sexpr_t sexpr, Sexpr_t env_exec)
Sexpr_t	mop3_cudaDeviceSynchronize (Lamb &lamb, Sexpr_t sexpr, Sexpr_t env_exec)
Sexpr_t	mop3_cudaGetDeviceCount (Lamb &lamb, Sexpr_t sexpr, Sexpr_t env_exec)
Sexpr_t	mop3_cudaGetDeviceProp (Lamb &lamb, Sexpr_t sexpr, Sexpr_t env_exec)
Sexpr_t	Cuda_install_mop3 (Lamb &lamb, Sexpr_t sexpr, Sexpr_t env_exec)

Detailed Description

LambLisp CUDA GPU builtins.

Function Documentation

cudaDevicePtrDestroyer()

void cudaDevicePtrDestroyer

(

Ptr_t

p

)

GC destructor for device memory.

cudaHostPtrDestroyer()

void cudaHostPtrDestroyer

(

Ptr_t

p

)

GC destructor for pinned host memory.

array_matmul_cublas_f32()

int array_matmul_cublas_f32	(	const LL_float32 *	A,
		const LL_float32 *	B,
		LL_float32 *	C,
		int	m,
		int	k,
		int	n )

Row-major C(m x n) = A(m x k) * B(k x n) via cuBLAS, on HOST buffers. Our arrays are host-resident (shape . data) pairs, so we round-trip through the device: cudaMalloc + H2D + gemm + D2H + cudaFree.

cuBLAS is COLUMN-major but our buffers are ROW-major. Rather than copy/transpose, use the identity (A*B)^T = B^T * A^T. A row-major (m x k) buffer, reinterpreted column-major, already IS A^T (k x m). So computing, in cuBLAS column-major terms, Ct(n x m) = B^T(n x k) * A^T(k x m) — i.e. sgemm(N,N, n, m, k, dB, n, dA, k, dC, n) — writes exactly our desired row-major C(m x n) into dC with no transposes. Returns 0 on success, nonzero on any CUDA/cuBLAS error.

array_matmul_cublas_f64()

int array_matmul_cublas_f64	(	const LL_float64 *	A,
		const LL_float64 *	B,
		LL_float64 *	C,
		int	m,
		int	k,
		int	n )

Double-precision counterpart of array_matmul_cublas_f32 (uses cublasDgemm).

mop3_cpp_vector_add()

Sexpr_t mop3_cpp_vector_add	(	Lamb &	lamb,
		Sexpr_t	sexpr,
		Sexpr_t	env_exec )

CPU-side vector add: C[i] = A[i] + B[i] for i in [0,n). Args: (A B C n) – A, B, C are realvectors; n is element count. Available with or without CUDA; used in three-way correctness tests.

mop3_cudaLaunch_kernel1()

Sexpr_t mop3_cudaLaunch_kernel1	(	Lamb &	lamb,
		Sexpr_t	sexpr,
		Sexpr_t	env_exec )

Launch kernel1 (element *= 2) on a device buffer; args: (stream data [n]).

mop3_cudakernel_vectorAdd()

Sexpr_t mop3_cudakernel_vectorAdd	(	Lamb &	lamb,
		Sexpr_t	sexpr,
		Sexpr_t	env_exec )

GPU vector add C[i] = A[i]+B[i]; args: (A B C n) as device/pinned pointers or bytevectors.

mop3_cudaMalloc()

Sexpr_t mop3_cudaMalloc	(	Lamb &	lamb,
		Sexpr_t	sexpr,
		Sexpr_t	env_exec )

Allocate n bytes of device memory; returns a T_CPP_HEAP object with cudaFree destructor.

mop3_cudaMallocHost()

Sexpr_t mop3_cudaMallocHost	(	Lamb &	lamb,
		Sexpr_t	sexpr,
		Sexpr_t	env_exec )

Allocate n bytes of page-locked host memory; returns a T_CPP_HEAP object with cudaFreeHost destructor.

mop3_cudaFree()

Sexpr_t mop3_cudaFree	(	Lamb &	lamb,
		Sexpr_t	sexpr,
		Sexpr_t	env_exec )

Free one or more device memory objects returned by cudaMalloc.

mop3_cudaMemcpy()

Sexpr_t mop3_cudaMemcpy	(	Lamb &	lamb,
		Sexpr_t	sexpr,
		Sexpr_t	env_exec )

Combines cudaMemcpy() and cudaMemcpyAsync(). If the stream argument is provided then the asynchronous copy will be used. count is in bytes. dst and src may be device pointers (cudaMalloc), pinned host pointers (cudaMallocHost), or bytevectors.

mop3_cudaMemset()

Sexpr_t mop3_cudaMemset	(	Lamb &	lamb,
		Sexpr_t	sexpr,
		Sexpr_t	env_exec )

Fill count bytes of device memory with a byte value: (cudaMemset ptr val count).

mop3_cudaStreamCreate()

Sexpr_t mop3_cudaStreamCreate	(	Lamb &	lamb,
		Sexpr_t	sexpr,
		Sexpr_t	env_exec )

Create a CUDA stream; returns a T_CPP_HEAP object with cudaStreamDestroy destructor.

mop3_cudaStreamSynchronize()

Sexpr_t mop3_cudaStreamSynchronize	(	Lamb &	lamb,
		Sexpr_t	sexpr,
		Sexpr_t	env_exec )

Block until all operations in the given CUDA stream have completed.

mop3_cudaDeviceSynchronize()

Sexpr_t mop3_cudaDeviceSynchronize	(	Lamb &	lamb,
		Sexpr_t	sexpr,
		Sexpr_t	env_exec )

Block until all CUDA work on the current device has completed.

mop3_cudaGetDeviceCount()

Sexpr_t mop3_cudaGetDeviceCount	(	Lamb &	lamb,
		Sexpr_t	sexpr,
		Sexpr_t	env_exec )

Return the number of CUDA-capable devices available on the host.

mop3_cudaGetDeviceProp()

Sexpr_t mop3_cudaGetDeviceProp	(	Lamb &	lamb,
		Sexpr_t	sexpr,
		Sexpr_t	env_exec )

Return an association list of device properties for the given device index (default 0).