Kernel-supporting jll #389

wsmoses · 2024-12-17T00:31:10Z

No description provided.

github-actions · 2024-12-17T01:06:35Z

src/XLA.jl

@@ -131,6 +131,7 @@ function __init__()
        end
    end

+    @ccall MLIR.API.mlir_c.RegisterCustomCallTarget("enzymexla_gpu"::Cstring, cglobal((:EnzymeGPUCustomCall, MLIR.API.mlir_c))::Ptr{Cvoid}, "CUDA"::Cstring)::Cvoid


[JuliaFormatter] _{reported by reviewdog 🐶}

Suggested change

@ccall MLIR.API.mlir_c.RegisterCustomCallTarget("enzymexla_gpu"::Cstring, cglobal((:EnzymeGPUCustomCall, MLIR.API.mlir_c))::Ptr{Cvoid}, "CUDA"::Cstring)::Cvoid

@ccall MLIR.API.mlir_c.RegisterCustomCallTarget(

"enzymexla_gpu"::Cstring,

cglobal((:EnzymeGPUCustomCall, MLIR.API.mlir_c))::Ptr{Cvoid},

"CUDA"::Cstring,

)::Cvoid

github-actions · 2024-12-17T01:14:21Z

ext/ReactantCUDAExt.jl

+Reactant.@reactant_overlay @noinline function (func::LLVMFunc{F,tt})(args...; convert=Val(false), blocks::CuDim=1, threads::CuDim=1,
                cooperative::Bool=false, shmem::Integer=0, call_kwargs...) where{F, tt}


[JuliaFormatter] _{reported by reviewdog 🐶}

Suggested change

Reactant.@reactant_overlay @noinline function (func::LLVMFunc{F,tt})(args...; convert=Val(false), blocks::CuDim=1, threads::CuDim=1,

cooperative::Bool=false, shmem::Integer=0, call_kwargs...) where{F, tt}

Reactant.@reactant_overlay @noinline function (func::LLVMFunc{F,tt})(

args...;

convert=Val(false),

blocks::CuDim=1,

threads::CuDim=1,

cooperative::Bool=false,

shmem::Integer=0,

call_kwargs...,

) where {F,tt}

github-actions · 2024-12-17T01:35:08Z

test/cuda.jl

    @show @code_hlo optimize=false square!(A)
+    @show @code_hlo optimize=:before_kernel square!(A)


[JuliaFormatter] _{reported by reviewdog 🐶}

Suggested change

@show @code_hlo optimize=false square!(A)

@show @code_hlo optimize=:before_kernel square!(A)

@show @code_hlo optimize = false square!(A)

@show @code_hlo optimize = :before_kernel square!(A)

github-actions · 2024-12-17T03:45:56Z

ext/ReactantCUDAExt.jl

+    for idx in (blockdim.x, blockdim.y, blockdim.z, threaddim.x, threaddim.y, threaddim.z, shmem)
+        push!(operands, Reactant.TracedUtils.promote_to(Reactant.TracedRNumber{Int}, idx).mlir_data)


[JuliaFormatter] _{reported by reviewdog 🐶}

Suggested change

for idx in (blockdim.x, blockdim.y, blockdim.z, threaddim.x, threaddim.y, threaddim.z, shmem)

push!(operands, Reactant.TracedUtils.promote_to(Reactant.TracedRNumber{Int}, idx).mlir_data)

for idx in

(blockdim.x, blockdim.y, blockdim.z, threaddim.x, threaddim.y, threaddim.z, shmem)

push!(

operands,

Reactant.TracedUtils.promote_to(Reactant.TracedRNumber{Int}, idx).mlir_data,

)

github-actions · 2024-12-17T03:45:56Z

ext/ReactantCUDAExt.jl

+        push!(operands, Reactant.TracedUtils.promote_to(Reactant.TracedRNumber{Int}, idx).mlir_data)
+    end
+    for arg in mlir_args
+	push!(operands, arg)


[JuliaFormatter] _{reported by reviewdog 🐶}

Suggested change

push!(operands, arg)

push!(operands, arg)

github-actions · 2024-12-17T03:45:56Z

ext/ReactantCUDAExt.jl

+	MLIR.IR.NamedAttribute("fn", MLIR.IR.FlatSymbolRefAttribute(Base.String(fname))),
+	MLIR.IR.NamedAttribute("output_operand_aliases", MLIR.IR.Attribute(output_operand_aliases))


[JuliaFormatter] _{reported by reviewdog 🐶}

Suggested change

MLIR.IR.NamedAttribute("fn", MLIR.IR.FlatSymbolRefAttribute(Base.String(fname))),

MLIR.IR.NamedAttribute("output_operand_aliases", MLIR.IR.Attribute(output_operand_aliases))

MLIR.IR.NamedAttribute("fn", MLIR.IR.FlatSymbolRefAttribute(Base.String(fname))),

MLIR.IR.NamedAttribute(

"output_operand_aliases", MLIR.IR.Attribute(output_operand_aliases)

),

github-actions · 2024-12-17T03:45:56Z

ext/ReactantCUDAExt.jl


-    call = MLIR.Dialects.stablehlo.custom_call(mlir_args; result_0=restys, call_target_name="reactant_gpu_call", output_operand_aliases, backend_config=MLIR.IR.Attribute(fname))
-    # call = MLIR.Dialects.stablehlo.custom_call(mlir_args; result_0=restys, call_target_name="reactant_gpu_call", output_operand_aliases, backend_config=MLIR.IR.Attribute(func.mod))
    for (i, res) in enumerate(rarrays)
       res.mlir_data = transpose_val(MLIR.IR.result(call, i))


[JuliaFormatter] _{reported by reviewdog 🐶}

Suggested change

res.mlir_data = transpose_val(MLIR.IR.result(call, i))

res.mlir_data = transpose_val(MLIR.IR.result(call, i))

github-actions · 2024-12-17T03:45:57Z

ext/ReactantCUDAExt.jl

@@ -379,7 +394,7 @@ function compiler_cache(ctx::MLIR.IR.Context)
    return cache
 end

-Reactant.@reactant_override @noinline function CUDA.cufunction(f::F, tt::TT=Tuple{}; kwargs...) where {F,TT}
+Reactant.@reactant_overlay @noinline function CUDA.cufunction(f::F, tt::TT=Tuple{}; kwargs...) where {F,TT}


[JuliaFormatter] _{reported by reviewdog 🐶}

Suggested change

Reactant.@reactant_overlay @noinline function CUDA.cufunction(f::F, tt::TT=Tuple{}; kwargs...) where {F,TT}

Reactant.@reactant_overlay @noinline function CUDA.cufunction(

f::F, tt::TT=Tuple{}; kwargs...

) where {F,TT}

github-actions · 2024-12-17T03:45:57Z

ext/ReactantCUDAExt.jl

@@ -379,7 +394,7 @@ function compiler_cache(ctx::MLIR.IR.Context)
    return cache
 end

-Reactant.@reactant_override @noinline function CUDA.cufunction(f::F, tt::TT=Tuple{}; kwargs...) where {F,TT}
+Reactant.@reactant_overlay @noinline function CUDA.cufunction(f::F, tt::TT=Tuple{}; kwargs...) where {F,TT}
    res = Base.@lock CUDA.cufunction_lock begin
        # compile the function
 	cache = compiler_cache(MLIR.IR.context())


[JuliaFormatter] _{reported by reviewdog 🐶}

Suggested change

cache = compiler_cache(MLIR.IR.context())

cache = compiler_cache(MLIR.IR.context())

github-actions · 2024-12-17T04:07:12Z

src/Compiler.jl

@@ -304,7 +306,28 @@ function compile_mlir!(mod, f, args; optimize::Union{Bool,Symbol}=true)

    optimize isa Bool && (optimize = ifelse(optimize, :all, :none))

+    toolkit = ""
+    if isdefined(Reactant_jll, :ptxas_path)
+	 toolkit = Reactant_jll.ptxas_path[1:end-length("/bin/ptxas")]


[JuliaFormatter] _{reported by reviewdog 🐶}

Suggested change

toolkit = Reactant_jll.ptxas_path[1:end-length("/bin/ptxas")]

toolkit = Reactant_jll.ptxas_path[1:(end - length("/bin/ptxas"))]

github-actions · 2024-12-17T04:07:12Z

src/Compiler.jl

+                    "remove-unnecessary-enzyme-ops",
+                    "enzyme-simplify-math",
+                    opt_passes,
+		    kern


[JuliaFormatter] _{reported by reviewdog 🐶}

Suggested change

kern

kern,

github-actions · 2024-12-17T04:07:12Z

src/Compiler.jl

@@ -340,6 +363,7 @@ function compile_mlir!(mod, f, args; optimize::Union{Bool,Symbol}=true)
                    "remove-unnecessary-enzyme-ops",
                    "enzyme-simplify-math",
                    opt_passes,
+		    kern


[JuliaFormatter] _{reported by reviewdog 🐶}

Suggested change

kern

kern,

github-actions · 2024-12-17T04:07:12Z

src/Compiler.jl

@@ -348,7 +372,7 @@ function compile_mlir!(mod, f, args; optimize::Union{Bool,Symbol}=true)
        run_pass_pipeline!(mod, join([opt_passes, "enzyme-batch", opt_passes], ","))
        run_pass_pipeline!(mod, "enzyme,arith-raise{stablehlo=true}"; enable_verifier=false)
        run_pass_pipeline!(
-            mod, "canonicalize,remove-unnecessary-enzyme-ops,enzyme-simplify-math"
+            mod, "canonicalize,remove-unnecessary-enzyme-ops,enzyme-simplify-math,"*kern


[JuliaFormatter] _{reported by reviewdog 🐶}

Suggested change

mod, "canonicalize,remove-unnecessary-enzyme-ops,enzyme-simplify-math,"*kern

mod, "canonicalize,remove-unnecessary-enzyme-ops,enzyme-simplify-math," * kern

wsmoses added 8 commits December 16, 2024 18:30

Kernel-supporting jll

22db5df

fix rulescc

b6d0615

adapt to hedron dep

f6b2238

init target

6121847

fixup

d62d01e

additional fixups

c3192d9

fixup

e86fac6

parse

6cea3df

github-actions bot reviewed Dec 17, 2024

View reviewed changes

overlay

e3f4df2

github-actions bot reviewed Dec 17, 2024

View reviewed changes

wsmoses added 2 commits December 16, 2024 19:28

fix

019a690

registry utils

7a4a403

github-actions bot reviewed Dec 17, 2024

View reviewed changes

wsmoses added 6 commits December 16, 2024 19:40

callname

8fef9ff

reg

fe43909

fix

be19519

fix bld

b0f679f

cleanup

1cd21ff

no pip

8b5d131

github-actions bot reviewed Dec 17, 2024

View reviewed changes

fix

3072670

github-actions bot reviewed Dec 17, 2024

View reviewed changes

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Kernel-supporting jll #389

Kernel-supporting jll #389

wsmoses commented Dec 17, 2024

github-actions bot Dec 17, 2024

github-actions bot Dec 17, 2024

github-actions bot Dec 17, 2024

github-actions bot Dec 17, 2024

github-actions bot Dec 17, 2024

github-actions bot Dec 17, 2024

github-actions bot Dec 17, 2024

github-actions bot Dec 17, 2024

github-actions bot Dec 17, 2024

github-actions bot Dec 17, 2024

github-actions bot Dec 17, 2024

github-actions bot Dec 17, 2024

github-actions bot Dec 17, 2024

		Reactant.@reactant_overlay @noinline function (func::LLVMFunc{F,tt})(args...; convert=Val(false), blocks::CuDim=1, threads::CuDim=1,
		cooperative::Bool=false, shmem::Integer=0, call_kwargs...) where{F, tt}

		@show @code_hlo optimize=false square!(A)
		@show @code_hlo optimize=:before_kernel square!(A)

		for idx in (blockdim.x, blockdim.y, blockdim.z, threaddim.x, threaddim.y, threaddim.z, shmem)
		push!(operands, Reactant.TracedUtils.promote_to(Reactant.TracedRNumber{Int}, idx).mlir_data)

		MLIR.IR.NamedAttribute("fn", MLIR.IR.FlatSymbolRefAttribute(Base.String(fname))),
		MLIR.IR.NamedAttribute("output_operand_aliases", MLIR.IR.Attribute(output_operand_aliases))

	res.mlir_data = transpose_val(MLIR.IR.result(call, i))
	res.mlir_data = transpose_val(MLIR.IR.result(call, i))

	cache = compiler_cache(MLIR.IR.context())
	cache = compiler_cache(MLIR.IR.context())

	toolkit = Reactant_jll.ptxas_path[1:end-length("/bin/ptxas")]
	toolkit = Reactant_jll.ptxas_path[1:(end - length("/bin/ptxas"))]

	mod, "canonicalize,remove-unnecessary-enzyme-ops,enzyme-simplify-math,"*kern
	mod, "canonicalize,remove-unnecessary-enzyme-ops,enzyme-simplify-math," * kern

Kernel-supporting jll #389

Are you sure you want to change the base?

Kernel-supporting jll #389

Conversation

wsmoses commented Dec 17, 2024

github-actions bot Dec 17, 2024

Choose a reason for hiding this comment

github-actions bot Dec 17, 2024

Choose a reason for hiding this comment

github-actions bot Dec 17, 2024

Choose a reason for hiding this comment

github-actions bot Dec 17, 2024

Choose a reason for hiding this comment

github-actions bot Dec 17, 2024

Choose a reason for hiding this comment

github-actions bot Dec 17, 2024

Choose a reason for hiding this comment

github-actions bot Dec 17, 2024

Choose a reason for hiding this comment

github-actions bot Dec 17, 2024

Choose a reason for hiding this comment

github-actions bot Dec 17, 2024

Choose a reason for hiding this comment

github-actions bot Dec 17, 2024

Choose a reason for hiding this comment

github-actions bot Dec 17, 2024

Choose a reason for hiding this comment

github-actions bot Dec 17, 2024

Choose a reason for hiding this comment

github-actions bot Dec 17, 2024

Choose a reason for hiding this comment