# Writing a Metal Kernel
**URL:** https://discourse.julialang.org/t/writing-a-metal-kernel/100891
**Category:** GPU
**Created:** [June 27, 2023, 1:44pm UTC](https://discourse.julialang.org/t/writing-a-metal-kernel/100891 "2023-06-27T13:44:05Z")
**Posts on this page:** 10
**Page:** 1
### Author:  [@matthieu](https://discourse.julialang.org/u/matthieu)
#### Post date: [June 27, 2023, 1:44pm UTC](https://discourse.julialang.org/t/writing-a-metal-kernel/100891/1 "2023-06-27T13:44:05Z")
I have the following kernel in CUDA
```julia
function inv!(scale, nthreads)
nblocks = cld(length(scale), nthreads)
@cuda threads=nthreads blocks=nblocks inv_kernel!(scale)
end
function inv_kernel!(scale)
index = (blockIdx().x - 1) * blockDim().x + threadIdx().x
stride = blockDim().x * gridDim().x
@inbounds for i = index:stride:length(scale)
scale[i] = 1 / sqrt(scale[i]))
end
end
```
How can I write an equivalent kernel in Metal?
---
### Author:  [@maleadt](https://discourse.julialang.org/u/maleadt)
#### Post date: [June 28, 2023, 5:55am UTC](https://discourse.julialang.org/t/writing-a-metal-kernel/100891/2 "2023-06-28T05:55:43Z")
Did you look at the examples? [Metal.jl/examples/vadd.jl at main · JuliaGPU/Metal.jl · GitHub](https://github.com/JuliaGPU/Metal.jl/blob/main/examples/vadd.jl)
---
### Author:  [@matthieu](https://discourse.julialang.org/u/matthieu)
#### Post date: [June 28, 2023, 12:11pm UTC](https://discourse.julialang.org/t/writing-a-metal-kernel/100891/3 "2023-06-28T12:11:12Z")
Thanks but these example do not seem to cover the keyword groups (the equivalent to blocks in CUDA?). That is, how should I modify your linked example to allow for parallelism?
---
### Author:  [@matthieu](https://discourse.julialang.org/u/matthieu)
#### Post date: [June 28, 2023, 12:14pm UTC](https://discourse.julialang.org/t/writing-a-metal-kernel/100891/4 "2023-06-28T12:14:31Z")
Here is what I have written. Is that correct / efficient?
```julia
function inv!(scale::MtlVector, nthreads::Integer)
nblocks = cld(length(scale), nthreads)
@metal threads=nthreads groups=nblocks inv_kernel!(scale)
end
function inv_kernel!(scale)
i = thread_position_in_grid_1d()
if i <= length(scale)
scale[i] = 1 / sqrt(scale[i])
end
return nothing
end
```
---
### Author:  [@Ronis\_BR](https://discourse.julialang.org/u/Ronis_BR)
#### Post date: [June 28, 2023, 12:47pm UTC](https://discourse.julialang.org/t/writing-a-metal-kernel/100891/5 "2023-06-28T12:47:17Z")
Hi @matthieu !
I think so, I would have done something very similar. However, is there any problem for using broadcasting: `scale .= 1 ./ sqrt.(scale)`?
```julia
julia> vec = MtlVector([10.0f0^(n - 1) for n in 1:20])
20-element MtlVector{Float32}:
1.0
10.0
100.0
1000.0
10000.0
100000.0
1.0f6
1.0f7
1.0f8
1.0f9
1.0f10
1.0f11
1.0f12
1.0f13
1.0f14
1.0f15
1.0f16
1.0f17
1.0f18
1.0f19
julia> vec .= 1 ./ sqrt.(vec)
20-element MtlVector{Float32}:
1.0
0.31622776
0.1
0.03162278
0.01
0.0031622779
0.001
0.0003162278
0.0001
3.1622774f-5
1.0f-5
3.1622778f-6
1.0f-6
3.1622776f-7
1.0f-7
3.1622776f-8
1.0f-8
3.1622776f-9
1.0f-9
3.1622777f-10
```
---
### Author:  [@matthieu](https://discourse.julialang.org/u/matthieu)
#### Post date: [June 28, 2023, 12:49pm UTC](https://discourse.julialang.org/t/writing-a-metal-kernel/100891/6 "2023-06-28T12:49:12Z")
Right, this is just a simple example — the real kernel I need to write involve indices / multiple arrays.
---
### Author:  [@Ronis\_BR](https://discourse.julialang.org/u/Ronis_BR)
#### Post date: [June 28, 2023, 12:49pm UTC](https://discourse.julialang.org/t/writing-a-metal-kernel/100891/7 "2023-06-28T12:49:46Z")
Ah ok! So you are in the right track 🙂
---
### Author:  [@Simo\_S](https://discourse.julialang.org/u/Simo_S)
#### Post date: [August 31, 2024, 12:45pm UTC](https://discourse.julialang.org/t/writing-a-metal-kernel/100891/8 "2024-08-31T12:45:09Z")
A question related to this. Indeed in the CUDA example above there is the “stride loop” which decouples the number of threads from the array length:
```julia
index = (blockIdx().x - 1) * blockDim().x + threadIdx().x
stride = blockDim().x * gridDim().x. # Total number of threads
```
what would be the correct way to do this in Metal.jl? Is it as simple as
```julia
index = thread_position_in_grid_1d()
stride = threads_per_grid_1d()
```
or do I need to factor in the thread groups?
---
### Author:  [@maleadt](https://discourse.julialang.org/u/maleadt)
#### Post date: [September 1, 2024, 7:21am UTC](https://discourse.julialang.org/t/writing-a-metal-kernel/100891/9 "2024-09-01T07:21:27Z")
Those `per_grid`/`in_grid` indexing functions should be sufficient, e.g., see how we implement Broadcast in Metal.jl: [Metal.jl/src/broadcast.jl at adac3bd000be7dc6b64725a8e8a7a47a154c4dfe · JuliaGPU/Metal.jl · GitHub](https://github.com/JuliaGPU/Metal.jl/blob/adac3bd000be7dc6b64725a8e8a7a47a154c4dfe/src/broadcast.jl#L138-L147)
---
### Author:  [@Simo\_S](https://discourse.julialang.org/u/Simo_S)
#### Post date: [September 1, 2024, 3:35pm UTC](https://discourse.julialang.org/t/writing-a-metal-kernel/100891/10 "2024-09-01T15:35:54Z")
Thanks!