Some basic example with Kernel

[kolkov]

Hi!
Please provide some basic example how to work with this library properly.
How to create basic Kernel and use it into computation with GPU.

Thanks!

[chewxy]

Yeah I'll get round to doing that.

[kolkov]

It is better to give two examples. One is quite simple, to understand the basics, and one more complex, to understand where to move on.
Like in Nvidia docs.
Thanks!

[cobnst]

@chewxy @kolkov well how to found the example?0.0

[Icaro-Lima]

I'm here 2 years later ... Any examples?

[malkhamis]

maybe this test file can serve as an example 🤔 ?

[Icaro-Lima]

Thank you! This is exactly what I need!

[neurlang]

I got a working kernel here

tl;dr: The order of operations which seems to be working is:
initCUDA()

device, err := cu.GetDevice(0)
ctx, err := device.MakeContext(cu.SchedAuto)
err = ctx.Lock()
// then do malloc etc for permanent memory
mod, err := cu.LoadData(kernel.PTXreduceCUDA)
fn, err := mod.Function("reduce")
stream, err := cu.MakeStream(cu.DefaultStream)

reduceCUDA()

err := cu.SetCurrentContext(ctx)
// single use memory malloc/ permanent memory clear (memset 0)
err = cu.MemcpyHtoD(d_input_nums, unsafe.Pointer(&input_numbers[0]), inputNumsSize) // copy some data to kernel
err = d_fn.LaunchAndSync(x[1][0], x[1][1], x[1][2], x[0][0], x[0][1], x[0][2], 0, d_stream, args)
err = cu.MemcpyDtoH(unsafe.Pointer(&result[0]), d_result, resultSize) // copy solution data back from device to host

destroyCUDA()

cu.MemFree(memory) // free all memories
ctx.Unlock()
ctx.Destroy()

the kernel itself is in the kernel/ subdir, compiled to .ptx as this:

#!/bin/bash
nvcc -ptx reduceCUDA.cu -o reduceCUDA.ptx

the kernel does have a sort of main function:

extern "C" __global__ void reduce(uint8_t *d_set, uint32_t *d_nums, uint32_t *alphabet, uint32_t* out) {
// the code
}

Hope this serves as a solid example to the next person.

[chewxy]

hey @neurlang do you want to put this in a PR as an example?

[neurlang]

pushing hashtron kernel to the main cu repo doesn't feel right. it would shift the maintenance burden to you, i still need to tweak the kernel sometimes, etc. I think I will keep it in my repo instead. Anyone can look there anyway.

Speaking of this repo (gorgonia/cu), I would prefer to put here more testcases which stress the code paths related to my kernel, to make sure that hashtron behavior is not broken here accidentally by you or by someone, now that it works pretty stable.

On the other hand, now that we (or simply me) know how to make a CU kernel, it would be a good idea to write it to the gorgonia.github.io site as a tutorial to make a CU kernel.

I'll open a separate issue on there instead.