Dpdl - GPU compute

			www.dpdl.io

by SEE Solutions ©

Dpdl enables to execute compute intensive operations and algorithms directly on GPGPUs or GPUs.

This can be performed via 'embedded code sections' implemented directly in OCL (OpenCL) or Wgsl (Web GPU Shading Language) code, or by other foreign programming models available.

The following Dpdl language plug-ins that enable to perform GPU compute are currently available:

Dpdl language plug-in opencl
Dpdl language plug-in wgsl
Dpdl language plug-in triton

OpenCL

This dedicated Dpdl language plug-in for 'OpenCL' can execute OCL (OpenCL Compute Language) compute code, which can be executed directly on GPGPUs, GPU's and even on multiple ordinary CPUs.

Exmample:

import('native')


println("testing simple OpenGl algorithm on GPU...")

println("loading native libraries 'libc' and 'libdpdlopencl'...")

object libc = native.loadLib("c")

raise(libc, "Error in loading 'libc'")

object libgpu = native.loadLib("dpdlopencl")

raise(libgpu, "Error in loading 'libdpdlopencl'")

println("getting dpdlopencl version...")

string ver = libgpu.DPDLNATIVE_OPENCL_getVersion()

println("version: " + ver)

# we allocate a data buffer with 10000 entries, each 8 bytes of memory
long DATA_SIZE = 10000L
long byte_cnt = 8L

println("allocating input buffer...")

object input_arr1 = libc.malloc(DATA_SIZE * byte_cnt)
input_arr1.setMemory(0L, DATA_SIZE * byte_cnt, 0x00B)

object input_arr2 = libc.malloc(DATA_SIZE * byte_cnt)
input_arr2.setMemory(0L, DATA_SIZE * byte_cnt, 0x00B)

object output_arr = libc.malloc(DATA_SIZE * byte_cnt)
input_arr.setMemory(0L, DATA_SIZE * byte_cnt, 0x00B)

println("populating some random input data...")

long cnt = 0L
double d_val
int tmp
for(cnt < DATA_SIZE)
	tmp = randInt(16, 1000)
	d_val = to_double(tmp)

	input_arr1.setDouble(cnt, d_val)
	input_arr2.setDouble(cnt, d_val+3000.0d)

	cnt = cnt+1L
endfor

# indexes of input & output vectors
int idx_vec_a = 0
int idx_vec_b = 1
int idx_vec_c = 2

int status

status = libgpu.DPDLNATIVE_OPENCL_addVec(idx_vec_a, "f64", DATA_SIZE, input_arr1)
status = libgpu.DPDLNATIVE_OPENCL_addVec(idx_vec_b, "f64", DATA_SIZE, input_arr2)
status = libgpu.DPDLNATIVE_OPENCL_addVec(idx_vec_c, "f64", DATA_SIZE, output_arr)

println("executing kernel...")

>>ocl
	#if defined(cl_khr_fp64)
	#  pragma OPENCL EXTENSION cl_khr_fp64: enable
	#elif defined(cl_amd_fp64)
	#  pragma OPENCL EXTENSION cl_amd_fp64: enable
	#else
	#  error double precision is not supported
	#endif

	kernel void dpdl_kernel(global const double *a, global const double *b, global double *c) {
		size_t gid = get_global_id(0);
		c[gid] = a[gid] + b[gid];
	};
<<

int exit_code = dpdl_exit_code()

println("embedded ocl exit code: " + exit_code)

int wait_copy = libgpu.DPDLNATIVE_OPENCL_waitReadBuffer(idx_vec_c, output_arr, DATA_SIZE)

println("OUTPUT: ")

double d_val_out
long c = 0L
for(c < DATA_SIZE)
	d_val_out = output_arr.getDouble(c)

	println("output_arr[" + c + "]=" + d_val_out)

	c = c+1L
endfor

println("finished")

Index

WGSL WebGPU Shading Language

This dedicated dpdl language plug-in for 'Wgsl' can execute WGSL (WebGPU Shading Language) compute code, which is executed on GPUs and complies to the WebGPU specification, directly embedded within Dpdl code.

The WGSL language can be used to programmatically execute workloads of graphic compute pipelines as well as compute algorithms directly on GPUs and GPGPU's

The 'Wgsl' dpdl language plug-in supports all major GPU platforms such as Nvidia, Intel, AMD, and others. The WGSL code embedded with this plug-in can run on a wide variety of laptops, workstations, mobile devices and any hardware with Vulkan, Metal, or DirectX support.

Example:

import('native')


println("testing Wgsl algorithm on GPU...")

println("loading native libraries 'c' and 'dpdlgpu'...")

object libc = native.loadLib("c")

object libgpu = native.loadLib("dpdlgpu")

raise(libgpu, "Error in loading 'dpdlgpu' lib")

println("getting dpdlgpu lib version...")

string ver = libgpu.DPDLNATIVE_GPU_getVersion()

println("version: " + ver)

# we allocate a data buffer with 10000 entries
long DATA_SIZE = 10000L
long byte_cnt = 4L

# indexes of the tensors & kernel created
int idx_tensor_in = 1
int idx_tensor_out = 1
int idx_kernel = 1

println("allocating input buffer...")

object input_arr = libc.malloc(DATA_SIZE * byte_cnt)
input_arr.setMemory(0L, DATA_SIZE * byte_cnt, 0x00B)

object output_arr = libc.malloc(DATA_SIZE * byte_cnt)
input_arr.setMemory(0L, DATA_SIZE * byte_cnt, 0x00B)

println("populating input data...")

long cnt = 0L
float f_val
int tmp
for(cnt < DATA_SIZE)
	tmp = randInt(16, 1000)
	f_val = to_float(tmp)

	input_arr.setFloat(cnt, f_val)

	cnt = cnt+1L
endfor

int status_in = libgpu.DPDLNATIVE_GPU_createTensorInputS1x1(idx_tensor_in, "f32", DATA_SIZE, input_arr)
int status_out =  libgpu.DPDLNATIVE_GPU_createTensorOutputS1x1(idx_tensor_out, "f32", DATA_SIZE)


dpdl_stack_var_put("precision", "f32")
dpdl_stack_var_put("workgroup_size", "256, 1, 1")

dpdl_stack_push("dpdl:applyvars")

>>wgsl
const SCALE_FACTOR: f32 = 0.7978845608028654; // sqrt(2.0 / Math.PI)

@group(0) @binding(0) var<storage, read_write> inp: array<{{precision}}>;
@group(0) @binding(1) var<storage, read_write> out: array<{{precision}}>;

@compute @workgroup_size({{workgroup_size}})
fn main(
    @builtin(global_invocation_id) GlobalInvocationID: vec3<u32>) {
    let i: u32 = GlobalInvocationID.x;
    if (i < arrayLength(&inp)) {
        let x: f32 = inp[i];
        out[i] = select(0.5 * x * (1.0 + tanh(SCALE_FACTOR
                 * (x + .044715 * x * x * x))), x, x > 10.0);
    }
}
<<
int exit_code = dpdl_exit_code()

println("Wgsl exit code: " + exit_code)

int status_kernel = libgpu.DPDLNATIVE_GPU_createKernelS1x1(idx_kernel, idx_tensor_in, idx_tensor_out, DATA_SIZE)

int status_dispatch = libgpu.DPDLNATIVE_GPU_dispatchKernel(idx_kernel)

int wait_copy = libgpu.DPDLNATIVE_GPU_waitCopyCPU(idx_kernel, idx_tensor_out, output_arr, DATA_SIZE)

println("OUTPUT: ")

float f_val_in, f_val_out
long c = 0L
for(c < DATA_SIZE)
	f_val_in = input_arr.getFloat(c)
	f_val_out = output_arr.getFloat(c)

	println(f_val_in + "=" + f_val_out)

	c = c+1L
endfor

println("finished")

Index

Triton

This dedicated Dpdl language plug-in for 'Triton' language allows to execute DNN compute kernels for AI on GPU hardware.

Triton is a python-based domain specific language (https://triton-lang.org) that provides useful abstractions and simplifies the development of compute kernels. It lowers the compute kernel definition to LLVM-IR dialects (MLIR) and PTX formats that run directly on GPUs.

The plug-in can execute Triton code as it is, namely in Python code, but it has also the capability to convert the same python code to JVM compatible code that makes use of the Java toolkit HAT (Heterogeneous Accelerator Toolkit), which is then compiled and executed on GPU hardware via CUDA or OpenCL backends.

Both approaches perform the lowering to the appropriate LLVM dialect for the GPUs where it's running

Example: simple example of an addition of two tensors

println("executing a compute kernel with Triton on GPU...")

object tensors = getObj("Tensors")

int tns_size = 98432

object tns_x = tensors.random(tns_size, 1)

object tns_y = tensors.random(tns_size, 1)

dpdl_stack_push(tns_size, tns_x, tns_y)

>>triton(my_compute)
	import torch

	import triton
	import triton.language as tl


	DEVICE = triton.runtime.driver.active.get_active_torch_device()

	@triton.jit
	def add(x_ptr, y_ptr, output_ptr, n_elements, BLOCK_SIZE: tl.constexpr):

		pid = tl.program_id(axis=0)

		block_start = pid * BLOCK_SIZE
		offsets = block_start + tl.arange(0, BLOCK_SIZE)

		mask = offsets < n_elements

		x = tl.load(x_ptr + offsets, mask=mask)
		y = tl.load(y_ptr + offsets, mask=mask)

		output = x + y

		tl.store(output_ptr + offsets, output, mask=mask)


	def dpdl_kernel(size, x: torch.Tensor, y: torch.Tensor):

	    output = torch.empty_like(x)

	    assert x.device == DEVICE and y.device == DEVICE and output.device == DEVICE

	    n_elements = output.numel()

	    grid = lambda meta: (triton.cdiv(n_elements, meta['BLOCK_SIZE']), )

	    add[grid](x, y, output, n_elements, BLOCK_SIZE=1024)

	    return output


<<
int exit_code = dpdl_exit_code()

println("exit_code: " + exit_code)

raise(exit_code, "Error in executing triton kernel")

object output = dpdl_stack_obj_get(my_compute)

println("output: ")

println(output)

Triton on JVM

The Triton DNN compute kernel based on python can also be converted to JVM compatible code that makes use of the HAT (Heterogeneous Accelerator Toolkit) toolkit (jdk.incubator.code). This is achieved through a novel reflection methoand code model lowering. The converted java code is than compiled and executed.

The dpdl plug-in option to active this is option: dpdlplugin:-convert HAT

Example

println("executing a compute kernel with Triton on GPU...")

object tensors = getObj("Tensors")

int tns_size = 98432

object tns_x = tensors.random(tns_size, 1)

object tns_y = tensors.random(tns_size, 1)

dpdl_stack_push("dpdlplugin:-convert HAT", tns_size, tns_x, tns_y)

>>triton(my_compute)
	import torch

	import triton
	import triton.language as tl


	DEVICE = triton.runtime.driver.active.get_active_torch_device()

	@triton.jit
	def add(x_ptr, y_ptr, output_ptr, n_elements, BLOCK_SIZE: tl.constexpr):

		pid = tl.program_id(axis=0)

		block_start = pid * BLOCK_SIZE
		offsets = block_start + tl.arange(0, BLOCK_SIZE)

		mask = offsets < n_elements

		x = tl.load(x_ptr + offsets, mask=mask)
		y = tl.load(y_ptr + offsets, mask=mask)

		output = x + y

		tl.store(output_ptr + offsets, output, mask=mask)


	def dpdl_kernel(size, x: torch.Tensor, y: torch.Tensor):

	    output = torch.empty_like(x)

	    assert x.device == DEVICE and y.device == DEVICE and output.device == DEVICE

	    n_elements = output.numel()

	    grid = lambda meta: (triton.cdiv(n_elements, meta['BLOCK_SIZE']), )

	    add[grid](x, y, output, n_elements, BLOCK_SIZE=1024)

	    return output


<<
int exit_code = dpdl_exit_code()

println("exit_code: " + exit_code)

raise(exit_code, "Error in executing Triton kernel")

object output = dpdl_stack_obj_get(my_compute)

println("output: ")

println(output)

Index

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Dpdl - GPU compute

Index

OpenCL

WGSL WebGPU Shading Language

Triton

Triton on JVM

Example

FilesExpand file tree

Dpdl_GPU_compute.md

Latest commit

History

Dpdl_GPU_compute.md

File metadata and controls

Dpdl - GPU compute

Index

OpenCL

WGSL WebGPU Shading Language

Triton

Triton on JVM

Example