# How to write device code?

**URL:** <https://discourse.julialang.org/t/how-to-write-device-code/11527>\
**Category:** GPU\
**Tags:** question\
**Created:** [June 8, 2018, 7:24am UTC](https://discourse.julialang.org/t/how-to-write-device-code/11527 "2018-06-08T07:24:38Z")\
**Posts on this page:** 5\
**Page:** 1

<div class="post-metadata">

**Author:** ![floswald](https://sea2.discourse-cdn.com/julialang/user_avatar/discourse.julialang.org/floswald/32/195_2.png) [@floswald](https://discourse.julialang.org/u/floswald)\
**Post date:** [June 8, 2018, 7:24am UTC](https://discourse.julialang.org/t/how-to-write-device-code/11527/1 "2018-06-08T07:24:38Z")

</div>

hi there

newcomer question here. suppose i am looking at some cuda code that has a function like

```julia
__device__ int binaryVal(const REAL x, const int n, const REAL* X)
{
}

```

where I read the ` __device__ ` thing as saying that thus function will only be called by the gpu device itself (not called directly from the cpu).

I understand that the `@cuda` macro has the effect of exposing a function on the gpu to the main julia context, as in

```julia
__global__ void vfStep(const parameters param, const REAL* K, const REAL* Z,
		       const REAL* P, const REAL* V0, REAL* V, REAL* G) 
{
}

```

so, question: is there a _julian_ difference between ` __global__ ` and ` __device__ `? if so, how do i specify it? if not, I just say `@cuda` in front of all code needed on the GPU and that’s it?

thanks!

---

<div class="post-metadata">

**Author:** ![maleadt](https://sea2.discourse-cdn.com/julialang/user_avatar/discourse.julialang.org/maleadt/32/10097_2.png) [@maleadt](https://discourse.julialang.org/u/maleadt)\
**Post date:** [June 8, 2018, 8:01am UTC](https://discourse.julialang.org/t/how-to-write-device-code/11527/2 "2018-06-08T08:01:17Z")

</div>

> [@floswald](#):
>
> so, question: is there a _julian_ difference between ` __global__ ` and ` __device__ ` ? if so, how do i specify it? if not, I just say `@cuda` in front of all code needed on the GPU and that’s it?

You don’t need to specify which functions are device, global or host functions, just add `@cuda` in front of a call to what you would have declared ` __global__ ` in CUDA C.  
And every other function you call from there is implicitly a ` __device__ ` function.

Note that the actual implementation might be device specific, eg. by calling GPU incompatible code or using device intrinsics, and for that reason we will probably add such an annotation in the future.

---

<div class="post-metadata">

**Author:** ![floswald](https://sea2.discourse-cdn.com/julialang/user_avatar/discourse.julialang.org/floswald/32/195_2.png) [@floswald](https://discourse.julialang.org/u/floswald)\
**Post date:** [June 8, 2018, 9:33am UTC](https://discourse.julialang.org/t/how-to-write-device-code/11527/3 "2018-06-08T09:33:19Z")

</div>

I see - very helpful! so I like looking at [CUDAnative.jl/examples/pairwise.jl at 2271487e4fd493c6a66a9b9bb48110fa6f6b6a1e · JuliaGPU/CUDAnative.jl · GitHub](https://github.com/JuliaGPU/CUDAnative.jl/blob/2271487e4fd493c6a66a9b9bb48110fa6f6b6a1e/examples/pairwise.jl#L85-L99)

- I specify a low-level implementation function, using appropriate data types, and using `CUDAnative.cospi` (for example) as std lib functions:

```julia-auto
function haversine_gpu(lat1::Float32, lon1::Float32, lat2::Float32, lon2::Float32, radius::Float32)

```

- I can specify another function, again using `CuArrays`, to provide an iteration over that function

```julia-auto
function pairwise_dist_kernel(lat::CuDeviceVector{Float32}, lon::CuDeviceVector{Float32},
                              rowresult::CuDeviceMatrix{Float32}, n)

```

- (notice to myself: not a single `@cuda` used so far!)
- then there is an interface function, setting up blocks and thread sizes, casting standard `Array{Float32}` as `CuArray`, and eventually emitting a call to above `pairwise_dist_kernel` as a `@cuda`
- right?

### couple more things

- `Float32`. you recommend that for most (all?) gpus? which characteristic of my GPU would tell me more about that?
- we **always** provide an output arrray to a `@cuda` function (i.e we always return `nothing`)?
- thanks for this great example and the excellent work! 🙂

---

<div class="post-metadata">

**Author:** ![maleadt](https://sea2.discourse-cdn.com/julialang/user_avatar/discourse.julialang.org/maleadt/32/10097_2.png) [@maleadt](https://discourse.julialang.org/u/maleadt)\
**Post date:** [June 8, 2018, 10:08am UTC](https://discourse.julialang.org/t/how-to-write-device-code/11527/4 "2018-06-08T10:08:00Z")

</div>

Sounds mostly right; some notes:

> I specify a low-level implementation function, using appropriate data types, and using CUDAnative.cospi (for example) as std lib functions:

Types aren’t necessary, so you can code generically. Also note that, with Julia, you can eg. pass `CUDAnative.cospi` as an argument to that function.

> I can specify another function, again using CuArrays, to provide an iteration over that function

These are `CuDeviceArray`s, a type by CUDAnative, not to be confused by CuArray which implements the _host_ GPU array type (as implemented by either CuArrays.jl or CUDAdrv.jl).

Depending on which type of iteration you need, you could also just broadcast or reduce a scalar function like `haversine_gpu` over some `CuArray`s, avoiding the need for any `@cuda` or GPU-specific code like `pairwise_dist_kernel`. At this point, there’s not many of those iteration functions implemented in CuArrays though, so you’ll often end up implementing your own kernels (if you keep them generic, please contribute them to CuArrays.jl 🙂).

> then there is an interface function, setting up blocks and thread sizes, casting standard Array{Float32} as CuArray, and eventually emitting a call to above pairwise\_dist\_kernel as a @cuda

It would probably make sense for your outer interface to take a CuArray and work with that directly, making it possible to compose operations without paying the cost of transferring to and from device memory.

> Float32. you recommend that for most (all?) gpus? which characteristic of my GPU would tell me more about that?

Generally yes, Float64 is _way_ slower (think 50-100x). Some high-end GPUs feature slightly better FP64 performance, but it’s still slow and better avoided.

> we always provide an output arrray to a @cuda function (i.e we always return nothing)?

Kernels can’t return values. Nowadays (Julia 0.7 + latest CUDAnative) you don’t need to `return nothing` explicitly, as kernels are wrapped automatically.

---

<div class="post-metadata">

**Author:** ![duanestorti](https://sea2.discourse-cdn.com/julialang/user_avatar/discourse.julialang.org/duanestorti/32/37374_2.png) [@duanestorti](https://discourse.julialang.org/u/duanestorti)\
**Post date:** [September 4, 2020, 10:32pm UTC](https://discourse.julialang.org/t/how-to-write-device-code/11527/5 "2020-09-04T22:32:48Z")

</div>

Coming from a CUDA and numba background, your reply about implicit recognition of device functions seems REALLY important. It would be great if this tidbit of critical information could be readily found in the CUDA.jl docs. Perhaps as addition to the FAQ?