# Converting result of round or floor as Int in Metal

**URL:** <https://discourse.julialang.org/t/converting-result-of-round-or-floor-as-int-in-metal/98545>\
**Category:** GPU\
**Created:** [May 9, 2023, 12:16pm UTC](https://discourse.julialang.org/t/converting-result-of-round-or-floor-as-int-in-metal/98545 "2023-05-09T12:16:37Z")\
**Posts on this page:** 11\
**Page:** 1

<div class="post-metadata">

**Author:** ![f.ij](https://avatars.discourse-cdn.com/v4/letter/f/8491ac/32.png) [@f.ij](https://discourse.julialang.org/u/f.ij)\
**Post date:** [May 9, 2023, 12:16pm UTC](https://discourse.julialang.org/t/converting-result-of-round-or-floor-as-int-in-metal/98545/1 "2023-05-09T12:16:37Z")

</div>

I’m trying to get into GPU programming and I’m trying to write a simple kernel that calculates the correlation function of an image. I’ve run into a problem where I cannot convert the result of the round or floor functions to an Int.

If I try to do the following as a very simple prototype

```julia
result = MtlArray(zeros(Int32, 10))

function test(result)
    i = thread_position_in_grid_1d()
    i1 = 1.f0
    i2 = 3.3f0
  
    result[i] = Int32(round(i1+i2))

    return
end

@metal threads = 10 groups = 1 test(result)

```

And run it, I get the following. Why is this not possible?

```julia
ERROR: InvalidIRError: compiling MethodInstance for test(::MtlDeviceVector{Int32, 1}) resulted in invalid LLVM IR
Reason: unsupported call to an unknown function (call to gpu_malloc)
Stacktrace:
 [1] malloc
   @ ~/.julia/packages/GPUCompiler/BxfIW/src/runtime.jl:89
 [2] macro expansion
   @ ~/.julia/packages/GPUCompiler/BxfIW/src/runtime.jl:184
 [3] macro expansion
   @ ./none:0
 [4] box
   @ ./none:0
 [5] box_float32
   @ ~/.julia/packages/GPUCompiler/BxfIW/src/runtime.jl:213
 [6] Int32
   @ ./float.jl:900
 [7] test
   @ ~/Library/Mobile Documents/com~apple~CloudDocs/Documents/PhD/JuliaProjects/InteractiveIsing.jl/Tests/MetalSampling.jl:53
Hint: catch this exception as `err` and call `code_typed(err; interactive = true)` to introspect the erronous code with Cthulhu.jl
Stacktrace:
  [1] check_ir(job::GPUCompiler.CompilerJob{GPUCompiler.MetalCompilerTarget, Metal.MetalCompilerParams}, args::LLVM.Module)
    @ GPUCompiler ~/.julia/packages/GPUCompiler/BxfIW/src/validation.jl:145
  [2] macro expansion
    @ ~/.julia/packages/GPUCompiler/BxfIW/src/driver.jl:407 [inlined]
  [3] macro expansion
    @ ~/.julia/packages/TimerOutputs/RsWnF/src/TimerOutput.jl:253 [inlined]
  [4] macro expansion
    @ ~/.julia/packages/GPUCompiler/BxfIW/src/driver.jl:406 [inlined]
  [5] emit_llvm(job::GPUCompiler.CompilerJob; libraries::Bool, deferred_codegen::Bool, optimize::Bool, cleanup::Bool, only_entry::Bool, validate::Bool, ctx::LLVM.ThreadSafeContext)
    @ GPUCompiler ~/.julia/packages/GPUCompiler/BxfIW/src/utils.jl:89
  [6] emit_llvm
    @ ~/.julia/packages/GPUCompiler/BxfIW/src/utils.jl:83 [inlined]
  [7] codegen(output::Symbol, job::GPUCompiler.CompilerJob; libraries::Bool, deferred_codegen::Bool, optimize::Bool, cleanup::Bool, strip::Bool, validate::Bool, only_entry::Bool, parent_job::Nothing, ctx::LLVM.ThreadSafeContext)
    @ GPUCompiler ~/.julia/packages/GPUCompiler/BxfIW/src/driver.jl:120
  [8] codegen
    @ ~/.julia/packages/GPUCompiler/BxfIW/src/driver.jl:94 [inlined]
  [9] compile(target::Symbol, job::GPUCompiler.CompilerJob; libraries::Bool, deferred_codegen::Bool, optimize::Bool, cleanup::Bool, strip::Bool, validate::Bool, only_entry::Bool, ctx::LLVM.ThreadSafeContext)
    @ GPUCompiler ~/.julia/packages/GPUCompiler/BxfIW/src/driver.jl:90
 [10] compile
    @ ~/.julia/packages/GPUCompiler/BxfIW/src/driver.jl:81 [inlined]
 [11] compile(job::GPUCompiler.CompilerJob, ctx::LLVM.ThreadSafeContext)
    @ Metal ~/.julia/packages/Metal/TtPHW/src/compiler/compilation.jl:59
 [12] #59
    @ ~/.julia/packages/Metal/TtPHW/src/compiler/compilation.jl:55 [inlined]
 [13] LLVM.ThreadSafeContext(f::Metal.var"#59#60"{GPUCompiler.CompilerJob{GPUCompiler.MetalCompilerTarget, Metal.MetalCompilerParams}})
    @ LLVM ~/.julia/packages/LLVM/TLGyi/src/executionengine/ts_module.jl:14
 [14] JuliaContext(f::Metal.var"#59#60"{GPUCompiler.CompilerJob{GPUCompiler.MetalCompilerTarget, Metal.MetalCompilerParams}})
    @ GPUCompiler ~/.julia/packages/GPUCompiler/BxfIW/src/driver.jl:35
 [15] compile
    @ ~/.julia/packages/Metal/TtPHW/src/compiler/compilation.jl:54 [inlined]
 [16] actual_compilation(cache::Dict{UInt64, Any}, key::UInt64, cfg::GPUCompiler.CompilerConfig{GPUCompiler.MetalCompilerTarget, Metal.MetalCompilerParams}, ft::Type, tt::Type, compiler::typeof(Metal.compile), linker::typeof(Metal.link))
    @ GPUCompiler ~/.julia/packages/GPUCompiler/BxfIW/src/cache.jl:53
 [17] cached_compilation(cache::Dict{UInt64, Any}, cfg::GPUCompiler.CompilerConfig{GPUCompiler.MetalCompilerTarget, Metal.MetalCompilerParams}, ft::Type, tt::Type, compiler::Function, linker::Function)
    @ GPUCompiler ~/.julia/packages/GPUCompiler/BxfIW/src/cache.jl:37
 [18] macro expansion
    @ ~/.julia/packages/Metal/TtPHW/src/compiler/execution.jl:161 [inlined]
 [19] macro expansion
    @ ./lock.jl:267 [inlined]
 [20] mtlfunction(f::typeof(test), tt::Type{Tuple{MtlDeviceVector{Int32, 1}}}; name::Nothing, kwargs::Base.Pairs{Symbol, Union{}, Tuple{}, NamedTuple{(), Tuple{}}})
    @ Metal ~/.julia/packages/Metal/TtPHW/src/compiler/execution.jl:157
 [21] mtlfunction(f::typeof(test), tt::Type{Tuple{MtlDeviceVector{Int32, 1}}})
    @ Metal ~/.julia/packages/Metal/TtPHW/src/compiler/execution.jl:155
 [22] top-level scope
    @ ~/.julia/packages/Metal/TtPHW/src/compiler/execution.jl:77

```

---

<div class="post-metadata">

**Author:** ![SteffenPL](https://sea2.discourse-cdn.com/julialang/user_avatar/discourse.julialang.org/steffenpl/32/206270_2.png) [@SteffenPL](https://discourse.julialang.org/u/SteffenPL)\
**Post date:** [May 9, 2023, 12:25pm UTC](https://discourse.julialang.org/t/converting-result-of-round-or-floor-as-int-in-metal/98545/2 "2023-05-09T12:25:38Z")

</div>

Does `round(Int32, i1+i2)` work?

---

<div class="post-metadata">

**Author:** ![f.ij](https://avatars.discourse-cdn.com/v4/letter/f/8491ac/32.png) [@f.ij](https://discourse.julialang.org/u/f.ij)\
**Post date:** [May 9, 2023, 12:27pm UTC](https://discourse.julialang.org/t/converting-result-of-round-or-floor-as-int-in-metal/98545/3 "2023-05-09T12:27:53Z")

</div>

Wow, yes this works! Why does the other code error, however? I assumed there were no type instabilities in the code. I assume round(x) returns a ::typeof(x) and Int32 obviously returns an Int32.

---

<div class="post-metadata">

**Author:** ![f.ij](https://avatars.discourse-cdn.com/v4/letter/f/8491ac/32.png) [@f.ij](https://discourse.julialang.org/u/f.ij)\
**Post date:** [May 9, 2023, 12:39pm UTC](https://discourse.julialang.org/t/converting-result-of-round-or-floor-as-int-in-metal/98545/4 "2023-05-09T12:39:37Z")

</div>

Okay, I seem to really be misunderstanding something. Your suggestion works, but if I now try: `round(Int32, sqrt(i1+i2))`

I get:

```julia
ERROR: InvalidIRError: compiling MethodInstance for test(::MtlDeviceVector{Int32, 1}) resulted in invalid LLVM IR
Reason: unsupported call to an unknown function (call to gpu_malloc)
Stacktrace:
 [1] malloc
   @ ~/.julia/packages/GPUCompiler/BxfIW/src/runtime.jl:89
 [2] macro expansion
   @ ~/.julia/packages/GPUCompiler/BxfIW/src/runtime.jl:184
 [3] macro expansion
   @ ./none:0
 [4] box
   @ ./none:0
 [5] box_float32
   @ ~/.julia/packages/GPUCompiler/BxfIW/src/runtime.jl:213
 [6] trunc
   @ ./float.jl:893
 [7] round
   @ ./float.jl:384
 [8] test
   @ ~/Library/Mobile Documents/com~apple~CloudDocs/Documents/PhD/JuliaProjects/InteractiveIsing.jl/Tests/MetalSampling.jl:57
Hint: catch this exception as `err` and call `code_typed(err; interactive = true)` to introspect the erronous code with Cthulhu.jl
Stacktrace:
  [1] check_ir(job::GPUCompiler.CompilerJob{GPUCompiler.MetalCompilerTarget, Metal.MetalCompilerParams}, args::LLVM.Module)
    @ GPUCompiler ~/.julia/packages/GPUCompiler/BxfIW/src/validation.jl:145
  [2] macro expansion
    @ ~/.julia/packages/GPUCompiler/BxfIW/src/driver.jl:407 [inlined]
  [3] macro expansion
    @ ~/.julia/packages/TimerOutputs/RsWnF/src/TimerOutput.jl:253 [inlined]
  [4] macro expansion
    @ ~/.julia/packages/GPUCompiler/BxfIW/src/driver.jl:406 [inlined]
  [5] emit_llvm(job::GPUCompiler.CompilerJob; libraries::Bool, deferred_codegen::Bool, optimize::Bool, cleanup::Bool, only_entry::Bool, validate::Bool, ctx::LLVM.ThreadSafeContext)
    @ GPUCompiler ~/.julia/packages/GPUCompiler/BxfIW/src/utils.jl:89
  [6] emit_llvm
    @ ~/.julia/packages/GPUCompiler/BxfIW/src/utils.jl:83 [inlined]
  [7] codegen(output::Symbol, job::GPUCompiler.CompilerJob; libraries::Bool, deferred_codegen::Bool, optimize::Bool, cleanup::Bool, strip::Bool, validate::Bool, only_entry::Bool, parent_job::Nothing, ctx::LLVM.ThreadSafeContext)
    @ GPUCompiler ~/.julia/packages/GPUCompiler/BxfIW/src/driver.jl:120
  [8] codegen
    @ ~/.julia/packages/GPUCompiler/BxfIW/src/driver.jl:94 [inlined]
  [9] compile(target::Symbol, job::GPUCompiler.CompilerJob; libraries::Bool, deferred_codegen::Bool, optimize::Bool, cleanup::Bool, strip::Bool, validate::Bool, only_entry::Bool, ctx::LLVM.ThreadSafeContext)
    @ GPUCompiler ~/.julia/packages/GPUCompiler/BxfIW/src/driver.jl:90
 [10] compile
    @ ~/.julia/packages/GPUCompiler/BxfIW/src/driver.jl:81 [inlined]
 [11] compile(job::GPUCompiler.CompilerJob, ctx::LLVM.ThreadSafeContext)
    @ Metal ~/.julia/packages/Metal/TtPHW/src/compiler/compilation.jl:59
 [12] #59
    @ ~/.julia/packages/Metal/TtPHW/src/compiler/compilation.jl:55 [inlined]
 [13] LLVM.ThreadSafeContext(f::Metal.var"#59#60"{GPUCompiler.CompilerJob{GPUCompiler.MetalCompilerTarget, Metal.MetalCompilerParams}})
    @ LLVM ~/.julia/packages/LLVM/TLGyi/src/executionengine/ts_module.jl:14
 [14] JuliaContext(f::Metal.var"#59#60"{GPUCompiler.CompilerJob{GPUCompiler.MetalCompilerTarget, Metal.MetalCompilerParams}})
    @ GPUCompiler ~/.julia/packages/GPUCompiler/BxfIW/src/driver.jl:35
 [15] compile
    @ ~/.julia/packages/Metal/TtPHW/src/compiler/compilation.jl:54 [inlined]
 [16] actual_compilation(cache::Dict{UInt64, Any}, key::UInt64, cfg::GPUCompiler.CompilerConfig{GPUCompiler.MetalCompilerTarget, Metal.MetalCompilerParams}, ft::Type, tt::Type, compiler::typeof(Metal.compile), linker::typeof(Metal.link))
    @ GPUCompiler ~/.julia/packages/GPUCompiler/BxfIW/src/cache.jl:53
 [17] cached_compilation(cache::Dict{UInt64, Any}, cfg::GPUCompiler.CompilerConfig{GPUCompiler.MetalCompilerTarget, Metal.MetalCompilerParams}, ft::Type, tt::Type, compiler::Function, linker::Function)
    @ GPUCompiler ~/.julia/packages/GPUCompiler/BxfIW/src/cache.jl:37
 [18] macro expansion
    @ ~/.julia/packages/Metal/TtPHW/src/compiler/execution.jl:161 [inlined]
 [19] macro expansion
    @ ./lock.jl:267 [inlined]
 [20] mtlfunction(f::typeof(test), tt::Type{Tuple{MtlDeviceVector{Int32, 1}}}; name::Nothing, kwargs::Base.Pairs{Symbol, Union{}, Tuple{}, NamedTuple{(), Tuple{}}})
    @ Metal ~/.julia/packages/Metal/TtPHW/src/compiler/execution.jl:157
 [21] mtlfunction(f::typeof(test), tt::Type{Tuple{MtlDeviceVector{Int32, 1}}})
    @ Metal ~/.julia/packages/Metal/TtPHW/src/compiler/execution.jl:155
 [22] top-level scope
    @ ~/.julia/packages/Metal/TtPHW/src/compiler/execution.jl:77

```

---

<div class="post-metadata">

**Author:** ![SteffenPL](https://sea2.discourse-cdn.com/julialang/user_avatar/discourse.julialang.org/steffenpl/32/206270_2.png) [@SteffenPL](https://discourse.julialang.org/u/SteffenPL)\
**Post date:** [May 9, 2023, 12:44pm UTC](https://discourse.julialang.org/t/converting-result-of-round-or-floor-as-int-in-metal/98545/5 "2023-05-09T12:44:39Z")

</div>

I think in all cases the problem is that even if the return value is correct, it is important the the function themselves do not use 64bit types. For example the `sqrt` is probably using a `Float64` if you provide `Integer` as inputs…

(Maybe `round( Float32( i1 + i2 ) )` works. But I’m also guessing.)

---

<div class="post-metadata">

**Author:** ![f.ij](https://avatars.discourse-cdn.com/v4/letter/f/8491ac/32.png) [@f.ij](https://discourse.julialang.org/u/f.ij)\
**Post date:** [May 9, 2023, 12:46pm UTC](https://discourse.julialang.org/t/converting-result-of-round-or-floor-as-int-in-metal/98545/6 "2023-05-09T12:46:54Z")

</div>

So you mean to say that there probably is some intermediate step in sqrt which uses Float64? I’m inputting two Float32’s into the sqrt though. Also, why does this problem not come up if I don’t use it as an input to the round function. If I just set `result[I] = sqrt(i1+i2)` it works fine.

Edit:

Of course if I do `result[i] = round(Int32, somefloat)`, I insert a `result = MtlArray(zeros(Int32,10))` into the test function (and for Floats I use the corresponding `MtlArray(zeros(Float32, 10))`). I made a typo before in my opening post which I fixed now.

---

<div class="post-metadata">

**Author:** ![SteffenPL](https://sea2.discourse-cdn.com/julialang/user_avatar/discourse.julialang.org/steffenpl/32/206270_2.png) [@SteffenPL](https://discourse.julialang.org/u/SteffenPL)\
**Post date:** [May 9, 2023, 1:03pm UTC](https://discourse.julialang.org/t/converting-result-of-round-or-floor-as-int-in-metal/98545/7 "2023-05-09T13:03:05Z")

</div>

Oh, maybe someone else here can help better. I tried to look into the source code of `round` and got confused. (I will look at it later since I want to know the solution as well 😉

---

<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:** [May 9, 2023, 7:21pm UTC](https://discourse.julialang.org/t/converting-result-of-round-or-floor-as-int-in-metal/98545/8 "2023-05-09T19:21:47Z")

</div>

> [@f.ij](#):
>
> Why does the other code error, however? I assumed there were no type instabilities in the code.

Not a type instability, but a checked conversion: `round` still returns a Float32, so the conversion to Int32 may throw (which is currently unsupported by Metal.jl). The typical workaround is to use `unsafe_trunc(Int32, ...)` or `...%Int32`.

---

<div class="post-metadata">

**Author:** ![cortner](https://sea2.discourse-cdn.com/julialang/user_avatar/discourse.julialang.org/cortner/32/204_2.png) [@cortner](https://discourse.julialang.org/u/cortner)\
**Post date:** [December 31, 2025, 5:30pm UTC](https://discourse.julialang.org/t/converting-result-of-round-or-floor-as-int-in-metal/98545/9 "2025-12-31T17:30:22Z")

</div>

apologies for reviving this thread, but I’ve run into exactly the same issue, which seems to still be around. I also don’t see an Issue on Metal.jl (there seems to be a similar but not the same [issue](https://github.com/JuliaGPU/KernelAbstractions.jl/issues/254) on KernelAbstraction.jl). I’d just like to understand whether this is actually a “bug” or an inherent limitation of Metal, or ???

And how “safe” is `unsafe_trunc` ? What are the potential pitfalls? Should I go through the extra work to use different `round` or `floor` on different backends?

Would it be worthwhile to have “official” workarounds in KernelAbstractions.jl that can be centrally updated as the backends improve?

---

<div class="post-metadata">

**Author:** ![GunnarFarneback](https://sea2.discourse-cdn.com/julialang/user_avatar/discourse.julialang.org/gunnarfarneback/32/1827_2.png) [@GunnarFarneback](https://discourse.julialang.org/u/GunnarFarneback)\
**Post date:** [December 31, 2025, 5:34pm UTC](https://discourse.julialang.org/t/converting-result-of-round-or-floor-as-int-in-metal/98545/10 "2025-12-31T17:34:03Z")

</div>

> [@cortner](#):
>
> And how “safe” is `unsafe_trunc` ?

Presumably exactly as unsafe as the docstring says:

> Return the nearest integral value of type T whose absolute value is less  
> than or equal to the absolute value of x. If the value is not representable  
> by T, an arbitrary value will be returned. See also trunc.

---

<div class="post-metadata">

**Author:** ![cortner](https://sea2.discourse-cdn.com/julialang/user_avatar/discourse.julialang.org/cortner/32/204_2.png) [@cortner](https://discourse.julialang.org/u/cortner)\
**Post date:** [December 31, 2025, 5:35pm UTC](https://discourse.julialang.org/t/converting-result-of-round-or-floor-as-int-in-metal/98545/11 "2025-12-31T17:35:27Z")

</div>

well, that didn’t sound so unsafe to me, so the only difference is that it doesn’t throw an exception I guess?

Would I be right in assuming that said potential for an exception is what prevents using `round, floor, ceil, trunc` and so forth?