text:
function preprocess_q2σs(file::String; batch_size=50)
    f = jldopen(file, "r")
    points = f["points"]

    # 4096 points 1 channel, 10000 data in a data fragment
    xs = Float32.(points)

    # σs, 10000 data in data fragment
    ys = f["σs"]

    close(f)

    return DataLoader((xs, ys), batchsize=batch_size, shuffle=true)
end

function preprocess_q2ρ(file::String; batch_size=50, dim=100)
    f = jldopen(file, "r")
    points = f["points"]

    # 2*4096 points, 10000 data in a data fragment
    xs = Float32.(points)

    # 𝛒s, 100x100, 10000 data in data fragment
    ys = Float32.(reinterpret(reshape, Float64, f["𝛒s"][1:dim, 1:dim, :]))

    close(f)

    return DataLoader((xs, ys), batchsize=batch_size, shuffle=true)
end

function preprocess_q2args(file::String; batch_size=50)
    f = jldopen(file, "r")
    points = f["points"]

    # 4096 points 1 channel, 10000 data in a data fragment
    xs = Float32.(points)

    # r, θ, n̄, 10000 data in data fragment
    ys = f["args"]

    close(f)

    return DataLoader((xs, ys), batchsize=batch_size, shuffle=true)
end


text:
module AlgebraicMultigrid

import IterativeSolvers: gauss_seidel!
using LinearAlgebra
using SparseArrays, Printf
using Base.Threads

using LinearAlgebra: rmul!

# const mul! = A_mul_B!

const MT = false
const AMG = AlgebraicMultigrid

include("utils.jl")
export approximate_spectral_radius

include("strength.jl")
export Classical, SymmetricStrength

include("splitting.jl")
export RS

include("gallery.jl")
export poisson

include("smoother.jl")
export GaussSeidel, SymmetricSweep, ForwardSweep, BackwardSweep,
        JacobiProlongation

include("multilevel.jl")
export solve

include("classical.jl")
export ruge_stuben

include("aggregate.jl")
export StandardAggregation

include("aggregation.jl")
export fit_candidates, smoothed_aggregation

include("preconditioner.jl")
export aspreconditioner

end # module

ext: jl


text:
s = BenchmarkGroup()
suite["centrality"] = s

srand(1)
for G in DGLIST
    for n in [100, 200]
        g = G(n , 10n, seed=1)
        s["pagerank","$g"] = @benchmarkable pagerank($g)
    end
end

for G in DGLIST
    for n in [50, 100]
        g = G(n , 5n, seed=1)
        s["betweenness_centrality","$g"] = @benchmarkable betweenness_centrality($g)
    end
end


text:
module FastGaussQuadrature

using SpecialFunctions, LinearAlgebra
cumprod(A::AbstractArray) = Base.cumprod(A, dims=1)
cumprod(A::AbstractArray, d::Int) = Base.cumprod(A, dims=d)
sum(A::AbstractArray, n::Int) = Base.sum(A, dims=n)
sum(A) = Base.sum(A)
flipdim(A, d) = reverse(A, dims=d)

export gausslegendre
export gausschebyshev
export gausslaguerre
export gausshermite
export gaussjacobi
export gausslobatto
export gaussradau
export besselroots

import SpecialFunctions: besselj, airyai, airyaiprime

include("gausslegendre.jl")
include("gausschebyshev.jl")
include("gausslaguerre.jl")
include("gausshermite.jl")
include("gaussjacobi.jl")
include("gausslobatto.jl")
include("gaussradau.jl")
include("besselroots.jl")

end


text:
using BinaryBuilder

name = "PETSc"
version = v"3.16.5"

# Collection of sources required to build PETSc. Avoid using the git repository, it will
# require building SOWING which fails in all non-linux platforms.
sources = [
    ArchiveSource("https://www.mcs.anl.gov/petsc/mirror/release-snapshots/petsc-$(version).tar.gz",
    "7de8570eeb94062752d82a83208fc2bafc77b3f515023a4c14d8ff9440e66cac"),
    DirectorySource("./bundled"),
]

# Bash recipe for building across all platforms
script = raw"""
cd $WORKSPACE/srcdir/petsc*
atomic_patch -p1 $WORKSPACE/srcdir/patches/petsc_name_mangle.patch

BLAS_LAPACK_LIB="${libdir}/libopenblas.${dlext}"

if [[ "${target}" == *-mingw* ]]; then
    #atomic_patch -p1 $WORKSPACE/srcdir/patches/fix-header-cases.patch
    MPI_LIBS="${libdir}/msmpi.${dlext}"
else
    MPI_LIBS="[${libdir}/libmpifort.${dlext},${libdir}/libmpi.${dlext}]"
fi
mkdir $libdir/petsc
build_petsc()
{

    if [[ "${3}" == "Int64" ]]; then
        USE_INT64=1
    else
        USE_INT64=0
    fi
    mkdir $libdir/petsc/${1}_${2}_${3}
    ./configure --prefix=${libdir}/petsc/${1}_${2}_${3} \
        CC=${CC} \
        FC=${FC} \
        CXX=${CXX} \
        COPTFLAGS='-O3' \
        CXXOPTFLAGS='-O3' \
        CFLAGS='-fno-stack-protector' \
        LDFLAGS="-L${libdir}" \
        FOPTFLAGS='-O3' \
        --with-64-bit-indices=${USE_INT64} \
        --with-debugging=0 \
        --with-batch \
        --with-blaslapack-lib=$BLAS_LAPACK_LIB \
        --with-blaslapack-suffix="" \
        --known-64-bit-blas-indices=0 \
        --with-mpi-lib="${MPI_LIBS}" \
        --known-mpi-int64_t=0 \
        --with-mpi-include="${includedir}" \
        --with-sowing=0 \
        --with-precision=${1} \
        --with-scalar-type=${2} \
        --PETSC_ARCH=${target}_${1}_${2}_${3}

    if [[ "${target}" == *-mingw* ]]; then
        export CPPFLAGS="-Dpetsc_EXPORTS"
    elif [[ "${target}" == powerpc64le-* ]]; then
        export CFLAGS="-fPIC"
        export FFLAGS="-fPIC"
    fi

    make -j${nproc} \
        CPPFLAGS="${CPPFLAGS}" \
        CFLAGS="${CFLAGS}" \
        FFLAGS="${FFLAGS}"
    make install

    if [[ "${target}" == *-apple* ]]; then
        mv ${libdir}/petsc/${1}_${2}_${3}/lib/libpetsc.*.*.*.${dlext} "${libdir}/petsc/${1}_${2}_${3}/lib/libpetsc_${1}_${2}_${3}.${dlext}"
        install_name_tool -id libpetsc_${1}_${2}_${3}.${dlext} ${libdir}/petsc/${1}_${2}_${3}/lib/libpetsc_${1}_${2}_${3}.${dlext}
    else # windows and linux:
        mv ${libdir}/petsc/${1}_${2}_${3}/lib/libpetsc.${dlext}.*.*.* "${libdir}/petsc/${1}_${2}_${3}/lib/libpetsc_${1}_${2}_${3}.${dlext}"
    fi
    if [[ "${target}" == *-linux* ]]; then
        patchelf --set-soname "libpetsc_${1}_${2}_${3}.${dlext}" ${libdir}/petsc/${1}_${2}_${3}/lib/libpetsc_${1}_${2}_${3}.${dlext}
    fi
    # Remove now broken links

    rm ${libdir}/petsc/${1}_${2}_${3}/lib/libpetsc.*
    # Remove PETSc.pc because petsc.pc also exists, causing conflicts on case insensitive file-systems.
    rm ${libdir}/petsc/${1}_${2}_${3}/lib/pkgconfig/PETSc.pc
    sed -i -e "s/-lpetsc/-lpetsc_${1}_${2}_${3}/g" "$libdir/petsc/${1}_${2}_${3}/lib/pkgconfig/petsc.pc"
    mv $libdir/petsc/${1}_${2}_${3}/lib/pkgconfig/petsc.pc ${prefix}/lib/pkgconfig/petsc_${1}_${2}_${3}.pc
}

build_petsc double real Int32
build_petsc single real Int32
build_petsc double complex Int32
build_petsc single complex Int32
build_petsc double real Int64
build_petsc single real Int64
build_petsc double complex Int64
build_petsc single complex Int64
"""

# We attempt to build for all defined platforms
platforms = expand_gfortran_versions(supported_platforms(exclude=[Platform("i686", "windows")]))

products = [
    # Current default build, equivalent to Float64_Real_Int32
    LibraryProduct("libpetsc_double_real_Int32", :libpetsc, "\$libdir/petsc/double_real_Int32/lib")
    LibraryProduct("libpetsc_double_real_Int32", :libpetsc_Float64_Real_Int32, "\$libdir/petsc/double_real_Int32/lib")
    LibraryProduct("libpetsc_double_real_Int64", :libpetsc_Float64_Real_Int64, "\$libdir/petsc/double_real_Int64/lib")
    LibraryProduct("libpetsc_single_real_Int64", :libpetsc_Float32_Real_Int64, "\$libdir/petsc/single_real_Int64/lib")
    LibraryProduct("libpetsc_double_complex_Int64", :libpetsc_Float64_Complex_Int64, "\$libdir/petsc/double_complex_Int64/lib")
    LibraryProduct("libpetsc_single_complex_Int64", :libpetsc_Float32_Complex_Int64, "\$libdir/petsc/single_complex_Int64/lib")
    LibraryProduct("libpetsc_single_real_Int32", :libpetsc_Float32_Real_Int32, "\$libdir/petsc/single_real_Int32/lib")
    LibraryProduct("libpetsc_double_complex_Int32", :libpetsc_Float64_Complex_Int32, "\$libdir/petsc/double_complex_Int32/lib")
    LibraryProduct("libpetsc_single_complex_Int32", :libpetsc_Float32_Complex_Int32, "\$libdir/petsc/single_complex_Int32/lib")
]

dependencies = [
    Dependency("OpenBLAS32_jll"),
    Dependency("MPICH_jll"; platforms=filter(!Sys.iswindows, platforms)),
    Dependency("MicrosoftMPI_jll"; platforms=filter(Sys.iswindows, platforms)),
    Dependency("CompilerSupportLibraries_jll"),
]

# Build the tarballs.
build_tarballs(ARGS, name, version, sources, script, platforms, products, dependencies; julia_compat="1.6", preferred_gcc_version = v"9")

max_stars_repo_name: cmhamel/Yggdrasil

max_issues_repo_name: cmhamel/Yggdrasil

max_forks_repo_name: cmhamel/Yggdrasil


text:
@reexport module Bars

using ..Foos

export Bar, greet

struct Bar
  foo::Foo
  name
end

greet(bar::Bar) = phrase(bar.foo) * bar.name

end # module


text:
function differentiate(::Type{Differentiation}, arg0::ParameterFunction, arg1::jint, arg2::jdouble)
    return jcall(Differentiation, "differentiate", ParameterFunction, (ParameterFunction, jint, jdouble), arg0, arg1, arg2)
end

function differentiate(::Type{Differentiation}, arg0::StateFunction, arg1::jint, arg2::AttitudeProvider, arg3::OrbitType, arg4::PositionAngle, arg5::jdouble, arg6::jint)
    return jcall(Differentiation, "differentiate", StateJacobian, (StateFunction, jint, AttitudeProvider, OrbitType, PositionAngle, jdouble, jint), arg0, arg1, arg2, arg3, arg4, arg5, arg6)
end


text:
@inline function ExaTronKernels.dicfs(n::Int, alpha::Float64, A,
                       L,
                       wa1,
                       wa2,
                       tx)

    nbmax = 3
    alpham = 1.0e-3
    nbfactor = 512

    zero = 0.0
    one = 1.0
    two = 2.0

    # Compute the l2 norms of the columns of A.
    nrm2!(wa1, A, n, tx)

    # Compute the scaling matrix D.
    if tx <= n
        @inbounds wa2[tx] = (wa1[tx] > zero) ? one/sqrt(wa1[tx]) : one
    end
    @synchronize

    # Determine a lower bound for the step.

    if alpha <= zero
        alphas = alpham
    else
        alphas = alpha
    end

    # Compute the initial shift.

    alpha = zero
    for i in 1:n
        if A[i,i] == zero
            alpha = alphas
        else
            alpha = max(alpha, -A[i,i]*(wa2[i]^2))
        end
    end
    if alpha > 0
        alpha = max(alpha,alphas)
    end
    # if tx <= n  # No check on ty so that each warp has alpha.
    #     @inbounds alpha = (A[tx,tx] == zero) ? alphas : max(alpha, -A[tx,tx]*(wa2[tx]^2))
    # end

    # shfl_down_sync will automatically sync threads in a warp.

    # Find the maximum alpha in a warp and put it in the first thread.
    #offset = div(blockDim().x, 2)
    # offset = 16
    # while offset > 0
    #     alpha = max(alpha, CUDA.shfl_down_sync(0xffffffff, alpha, offset))
    #     offset >>= 1
    # end
    # # Broadcast it to the entire threads in a warp.
    # alpha = CUDA.shfl_sync(0xffffffff, alpha, 1)

    # if alpha > 0
    #     alpha = max(alpha,alphas)
    # end

    # Search for an acceptable shift. During the search we decrease
    # the lower bound alphas until we determine a lower bound that
    # is not acceptaable. We then increase the shift.
    # The lower bound is decreased by nbfactor at most nbmax times.

    nb = 1
    info = 0

    while true
        if tx <= n
            @inbounds for j=1:n
                L[j,tx] = A[j,tx] * wa2[j] * wa2[tx]
            end
            if alpha != zero
                @inbounds L[tx,tx] += alpha
            end
        end
        @synchronize

        # Attempt a Cholesky factorization.
        info = dicf(n, L, tx)

        # If the factorization exists, then test for termination.
        # Otherwise increment the shift.
        if info >= 0
            # If the shift is at the lower bound, reduce the shift.
            # Otherwise undo the scaling of L and exit.
            if alpha == alphas && nb < nbmax
                alphas /= nbfactor
                alpha = alphas
                nb = nb + 1
            else
                if tx <= n
                    @inbounds for j=1:n
                        if tx >= j
                            L[tx,j] /= wa2[tx]
                            L[j,tx] = L[tx,j]
                        end
                    end
                end
                @synchronize
                return
            end
        else
            alpha = max(two*alpha,alphas)
        end
    end
    @synchronize

    return
end


text:
"""
    DiagonalStandardNormal <: AbstactMvNormal

Multivariate Normal distribution with zero mean vector and identity covariance.
"""
struct DiagonalStandardNormal <: AbstractMvNormal
    D::Int
end

const DiagStdNormal = DiagonalStandardNormal

dim(d::DiagonalStandardNormal) = d.D

rand(rng::AbstractRNG, d::DiagonalStandardNormal) = randn(rng, dim(d))
rand(rng::AbstractRNG, d::DiagonalStandardNormal, N::Int) = randn(rng, dim(d), N)

@unionise function logpdf(d::DiagonalStandardNormal, x::AbstractVector{<:Real})
    @assert dim(d) == length(x)
    return -0.5 * (dim(d) * log(2π) + sum(abs2, x))
end
@unionise function logpdf(d::DiagonalStandardNormal, X::AbstractMatrix{<:Real})
    @assert dim(d) == size(X, 1)
    return -0.5 .* ((dim(d) * log(2π)) .+ reshape(mapreducedim(abs2, +, X, 1), size(X, 2)))
end


text:
module LibHSARuntime

using hsa_rocr_jll
export hsa_rocr_jll

using CEnum

import ...AMDGPU: libhsaruntime

const HSA_EXPORT_DECORATOR = identity
const HSA_CALL = nothing

@cenum hsa_status_t::UInt32 begin
    HSA_STATUS_SUCCESS = 0
    HSA_STATUS_INFO_BREAK = 1
    HSA_STATUS_ERROR = 4096
    HSA_STATUS_ERROR_INVALID_ARGUMENT = 4097
    HSA_STATUS_ERROR_INVALID_QUEUE_CREATION = 4098
    HSA_STATUS_ERROR_INVALID_ALLOCATION = 4099
    HSA_STATUS_ERROR_INVALID_AGENT = 4100
    HSA_STATUS_ERROR_INVALID_REGION = 4101
    HSA_STATUS_ERROR_INVALID_SIGNAL = 4102
    HSA_STATUS_ERROR_INVALID_QUEUE = 4103
    HSA_STATUS_ERROR_OUT_OF_RESOURCES = 4104
    HSA_STATUS_ERROR_INVALID_PACKET_FORMAT = 4105
    HSA_STATUS_ERROR_RESOURCE_FREE = 4106
    HSA_STATUS_ERROR_NOT_INITIALIZED = 4107
    HSA_STATUS_ERROR_REFCOUNT_OVERFLOW = 4108
    HSA_STATUS_ERROR_INCOMPATIBLE_ARGUMENTS = 4109
    HSA_STATUS_ERROR_INVALID_INDEX = 4110
    HSA_STATUS_ERROR_INVALID_ISA = 4111
    HSA_STATUS_ERROR_INVALID_ISA_NAME = 4119
    HSA_STATUS_ERROR_INVALID_CODE_OBJECT = 4112
    HSA_STATUS_ERROR_INVALID_EXECUTABLE = 4113
    HSA_STATUS_ERROR_FROZEN_EXECUTABLE = 4114
    HSA_STATUS_ERROR_INVALID_SYMBOL_NAME = 4115
    HSA_STATUS_ERROR_VARIABLE_ALREADY_DEFINED = 4116
    HSA_STATUS_ERROR_VARIABLE_UNDEFINED = 4117
    HSA_STATUS_ERROR_EXCEPTION = 4118
    HSA_STATUS_ERROR_INVALID_CODE_SYMBOL = 4120
    HSA_STATUS_ERROR_INVALID_EXECUTABLE_SYMBOL = 4121
    HSA_STATUS_ERROR_INVALID_FILE = 4128
    HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER = 4129
    HSA_STATUS_ERROR_INVALID_CACHE = 4130
    HSA_STATUS_ERROR_INVALID_WAVEFRONT = 4131
    HSA_STATUS_ERROR_INVALID_SIGNAL_GROUP = 4132
    HSA_STATUS_ERROR_INVALID_RUNTIME_STATE = 4133
    HSA_STATUS_ERROR_FATAL = 4134
end

function hsa_status_string(status, status_string)
    ccall((:hsa_status_string, libhsaruntime), hsa_status_t, (hsa_status_t, Ptr{Ptr{Cchar}}), status, status_string)
end

struct hsa_dim3_s
    x::UInt32
    y::UInt32
    z::UInt32
end

const hsa_dim3_t = hsa_dim3_s

@cenum hsa_access_permission_t::UInt32 begin
    HSA_ACCESS_PERMISSION_RO = 1
    HSA_ACCESS_PERMISSION_WO = 2
    HSA_ACCESS_PERMISSION_RW = 3
end

const hsa_file_t = Cint

# no prototype is found for this function at hsa.h:359:22, please use with caution
function hsa_init()
    ccall((:hsa_init, libhsaruntime), hsa_status_t, ())
end

# no prototype is found for this function at hsa.h:381:22, please use with caution
function hsa_shut_down()
    ccall((:hsa_shut_down, libhsaruntime), hsa_status_t, ())
end

@cenum hsa_endianness_t::UInt32 begin
    HSA_ENDIANNESS_LITTLE = 0
    HSA_ENDIANNESS_BIG = 1
end

@cenum hsa_machine_model_t::UInt32 begin
    HSA_MACHINE_MODEL_SMALL = 0
    HSA_MACHINE_MODEL_LARGE = 1
end

@cenum hsa_profile_t::UInt32 begin
    HSA_PROFILE_BASE = 0
    HSA_PROFILE_FULL = 1
end

@cenum hsa_system_info_t::UInt32 begin
    HSA_SYSTEM_INFO_VERSION_MAJOR = 0
    HSA_SYSTEM_INFO_VERSION_MINOR = 1
    HSA_SYSTEM_INFO_TIMESTAMP = 2
    HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY = 3
    HSA_SYSTEM_INFO_SIGNAL_MAX_WAIT = 4
    HSA_SYSTEM_INFO_ENDIANNESS = 5
    HSA_SYSTEM_INFO_MACHINE_MODEL = 6
    HSA_SYSTEM_INFO_EXTENSIONS = 7
    HSA_AMD_SYSTEM_INFO_BUILD_VERSION = 512
end

function hsa_system_get_info(attribute, value)
    ccall((:hsa_system_get_info, libhsaruntime), hsa_status_t, (hsa_system_info_t, Ptr{Cvoid}), attribute, value)
end

@cenum hsa_extension_t::UInt32 begin
    HSA_EXTENSION_FINALIZER = 0
    HSA_EXTENSION_IMAGES = 1
    HSA_EXTENSION_PERFORMANCE_COUNTERS = 2
    HSA_EXTENSION_PROFILING_EVENTS = 3
    HSA_EXTENSION_STD_LAST = 3
    HSA_AMD_FIRST_EXTENSION = 512
    HSA_EXTENSION_AMD_PROFILER = 512
    HSA_EXTENSION_AMD_LOADER = 513
    HSA_EXTENSION_AMD_AQLPROFILE = 514
    HSA_AMD_LAST_EXTENSION = 514
end

function hsa_extension_get_name(extension, name)
    ccall((:hsa_extension_get_name, libhsaruntime), hsa_status_t, (UInt16, Ptr{Ptr{Cchar}}), extension, name)
end

function hsa_system_extension_supported(extension, version_major, version_minor, result)
    ccall((:hsa_system_extension_supported, libhsaruntime), hsa_status_t, (UInt16, UInt16, UInt16, Ptr{Bool}), extension, version_major, version_minor, result)
end

function hsa_system_major_extension_supported(extension, version_major, version_minor, result)
    ccall((:hsa_system_major_extension_supported, libhsaruntime), hsa_status_t, (UInt16, UInt16, Ptr{UInt16}, Ptr{Bool}), extension, version_major, version_minor, result)
end

function hsa_system_get_extension_table(extension, version_major, version_minor, table)
    ccall((:hsa_system_get_extension_table, libhsaruntime), hsa_status_t, (UInt16, UInt16, UInt16, Ptr{Cvoid}), extension, version_major, version_minor, table)
end

function hsa_system_get_major_extension_table(extension, version_major, table_length, table)
    ccall((:hsa_system_get_major_extension_table, libhsaruntime), hsa_status_t, (UInt16, UInt16, Csize_t, Ptr{Cvoid}), extension, version_major, table_length, table)
end

struct hsa_agent_s
    handle::UInt64
end

const hsa_agent_t = hsa_agent_s

@cenum hsa_agent_feature_t::UInt32 begin
    HSA_AGENT_FEATURE_KERNEL_DISPATCH = 1
    HSA_AGENT_FEATURE_AGENT_DISPATCH = 2
end

@cenum hsa_device_type_t::UInt32 begin
    HSA_DEVICE_TYPE_CPU = 0
    HSA_DEVICE_TYPE_GPU = 1
    HSA_DEVICE_TYPE_DSP = 2
end

@cenum hsa_default_float_rounding_mode_t::UInt32 begin
    HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT = 0
    HSA_DEFAULT_FLOAT_ROUNDING_MODE_ZERO = 1
    HSA_DEFAULT_FLOAT_ROUNDING_MODE_NEAR = 2
end

@cenum hsa_agent_info_t::UInt32 begin
    HSA_AGENT_INFO_NAME = 0
    HSA_AGENT_INFO_VENDOR_NAME = 1
    HSA_AGENT_INFO_FEATURE = 2
    HSA_AGENT_INFO_MACHINE_MODEL = 3
    HSA_AGENT_INFO_PROFILE = 4
    HSA_AGENT_INFO_DEFAULT_FLOAT_ROUNDING_MODE = 5
    HSA_AGENT_INFO_BASE_PROFILE_DEFAULT_FLOAT_ROUNDING_MODES = 23
    HSA_AGENT_INFO_FAST_F16_OPERATION = 24
    HSA_AGENT_INFO_WAVEFRONT_SIZE = 6
    HSA_AGENT_INFO_WORKGROUP_MAX_DIM = 7
    HSA_AGENT_INFO_WORKGROUP_MAX_SIZE = 8
    HSA_AGENT_INFO_GRID_MAX_DIM = 9
    HSA_AGENT_INFO_GRID_MAX_SIZE = 10
    HSA_AGENT_INFO_FBARRIER_MAX_SIZE = 11
    HSA_AGENT_INFO_QUEUES_MAX = 12
    HSA_AGENT_INFO_QUEUE_MIN_SIZE = 13
    HSA_AGENT_INFO_QUEUE_MAX_SIZE = 14
    HSA_AGENT_INFO_QUEUE_TYPE = 15
    HSA_AGENT_INFO_NODE = 16
    HSA_AGENT_INFO_DEVICE = 17
    HSA_AGENT_INFO_CACHE_SIZE = 18
    HSA_AGENT_INFO_ISA = 19
    HSA_AGENT_INFO_EXTENSIONS = 20
    HSA_AGENT_INFO_VERSION_MAJOR = 21
    HSA_AGENT_INFO_VERSION_MINOR = 22
end

function hsa_agent_get_info(agent, attribute, value)
    ccall((:hsa_agent_get_info, libhsaruntime), hsa_status_t, (hsa_agent_t, hsa_agent_info_t, Ptr{Cvoid}), agent, attribute, value)
end

function hsa_iterate_agents(callback, data)
    ccall((:hsa_iterate_agents, libhsaruntime), hsa_status_t, (Ptr{Cvoid}, Ptr{Cvoid}), callback, data)
end

@cenum hsa_exception_policy_t::UInt32 begin
    HSA_EXCEPTION_POLICY_BREAK = 1
    HSA_EXCEPTION_POLICY_DETECT = 2
end

function hsa_agent_get_exception_policies(agent, profile, mask)
    ccall((:hsa_agent_get_exception_policies, libhsaruntime), hsa_status_t, (hsa_agent_t, hsa_profile_t, Ptr{UInt16}), agent, profile, mask)
end

struct hsa_cache_s
    handle::UInt64
end

const hsa_cache_t = hsa_cache_s

@cenum hsa_cache_info_t::UInt32 begin
    HSA_CACHE_INFO_NAME_LENGTH = 0
    HSA_CACHE_INFO_NAME = 1
    HSA_CACHE_INFO_LEVEL = 2
    HSA_CACHE_INFO_SIZE = 3
end

function hsa_cache_get_info(cache, attribute, value)
    ccall((:hsa_cache_get_info, libhsaruntime), hsa_status_t, (hsa_cache_t, hsa_cache_info_t, Ptr{Cvoid}), cache, attribute, value)
end

function hsa_agent_iterate_caches(agent, callback, data)
    ccall((:hsa_agent_iterate_caches, libhsaruntime), hsa_status_t, (hsa_agent_t, Ptr{Cvoid}, Ptr{Cvoid}), agent, callback, data)
end

function hsa_agent_extension_supported(extension, agent, version_major, version_minor, result)
    ccall((:hsa_agent_extension_supported, libhsaruntime), hsa_status_t, (UInt16, hsa_agent_t, UInt16, UInt16, Ptr{Bool}), extension, agent, version_major, version_minor, result)
end

function hsa_agent_major_extension_supported(extension, agent, version_major, version_minor, result)
    ccall((:hsa_agent_major_extension_supported, libhsaruntime), hsa_status_t, (UInt16, hsa_agent_t, UInt16, Ptr{UInt16}, Ptr{Bool}), extension, agent, version_major, version_minor, result)
end

struct hsa_signal_s
    handle::UInt64
end

const hsa_signal_t = hsa_signal_s

const hsa_signal_value_t = Int64

function hsa_signal_create(initial_value, num_consumers, consumers, signal)
    ccall((:hsa_signal_create, libhsaruntime), hsa_status_t, (hsa_signal_value_t, UInt32, Ptr{hsa_agent_t}, Ptr{hsa_signal_t}), initial_value, num_consumers, consumers, signal)
end

function hsa_signal_destroy(signal)
    ccall((:hsa_signal_destroy, libhsaruntime), hsa_status_t, (hsa_signal_t,), signal)
end

function hsa_signal_load_scacquire(signal)
    ccall((:hsa_signal_load_scacquire, libhsaruntime), hsa_signal_value_t, (hsa_signal_t,), signal)
end

function hsa_signal_load_relaxed(signal)
    ccall((:hsa_signal_load_relaxed, libhsaruntime), hsa_signal_value_t, (hsa_signal_t,), signal)
end

function hsa_signal_load_acquire(signal)
    ccall((:hsa_signal_load_acquire, libhsaruntime), hsa_signal_value_t, (hsa_signal_t,), signal)
end

function hsa_signal_store_relaxed(signal, value)
    ccall((:hsa_signal_store_relaxed, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_store_screlease(signal, value)
    ccall((:hsa_signal_store_screlease, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_store_release(signal, value)
    ccall((:hsa_signal_store_release, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_silent_store_relaxed(signal, value)
    ccall((:hsa_signal_silent_store_relaxed, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_silent_store_screlease(signal, value)
    ccall((:hsa_signal_silent_store_screlease, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_exchange_scacq_screl(signal, value)
    ccall((:hsa_signal_exchange_scacq_screl, libhsaruntime), hsa_signal_value_t, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_exchange_acq_rel(signal, value)
    ccall((:hsa_signal_exchange_acq_rel, libhsaruntime), hsa_signal_value_t, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_exchange_scacquire(signal, value)
    ccall((:hsa_signal_exchange_scacquire, libhsaruntime), hsa_signal_value_t, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_exchange_acquire(signal, value)
    ccall((:hsa_signal_exchange_acquire, libhsaruntime), hsa_signal_value_t, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_exchange_relaxed(signal, value)
    ccall((:hsa_signal_exchange_relaxed, libhsaruntime), hsa_signal_value_t, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_exchange_screlease(signal, value)
    ccall((:hsa_signal_exchange_screlease, libhsaruntime), hsa_signal_value_t, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_exchange_release(signal, value)
    ccall((:hsa_signal_exchange_release, libhsaruntime), hsa_signal_value_t, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_cas_scacq_screl(signal, expected, value)
    ccall((:hsa_signal_cas_scacq_screl, libhsaruntime), hsa_signal_value_t, (hsa_signal_t, hsa_signal_value_t, hsa_signal_value_t), signal, expected, value)
end

function hsa_signal_cas_acq_rel(signal, expected, value)
    ccall((:hsa_signal_cas_acq_rel, libhsaruntime), hsa_signal_value_t, (hsa_signal_t, hsa_signal_value_t, hsa_signal_value_t), signal, expected, value)
end

function hsa_signal_cas_scacquire(signal, expected, value)
    ccall((:hsa_signal_cas_scacquire, libhsaruntime), hsa_signal_value_t, (hsa_signal_t, hsa_signal_value_t, hsa_signal_value_t), signal, expected, value)
end

function hsa_signal_cas_acquire(signal, expected, value)
    ccall((:hsa_signal_cas_acquire, libhsaruntime), hsa_signal_value_t, (hsa_signal_t, hsa_signal_value_t, hsa_signal_value_t), signal, expected, value)
end

function hsa_signal_cas_relaxed(signal, expected, value)
    ccall((:hsa_signal_cas_relaxed, libhsaruntime), hsa_signal_value_t, (hsa_signal_t, hsa_signal_value_t, hsa_signal_value_t), signal, expected, value)
end

function hsa_signal_cas_screlease(signal, expected, value)
    ccall((:hsa_signal_cas_screlease, libhsaruntime), hsa_signal_value_t, (hsa_signal_t, hsa_signal_value_t, hsa_signal_value_t), signal, expected, value)
end

function hsa_signal_cas_release(signal, expected, value)
    ccall((:hsa_signal_cas_release, libhsaruntime), hsa_signal_value_t, (hsa_signal_t, hsa_signal_value_t, hsa_signal_value_t), signal, expected, value)
end

function hsa_signal_add_scacq_screl(signal, value)
    ccall((:hsa_signal_add_scacq_screl, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_add_acq_rel(signal, value)
    ccall((:hsa_signal_add_acq_rel, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_add_scacquire(signal, value)
    ccall((:hsa_signal_add_scacquire, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_add_acquire(signal, value)
    ccall((:hsa_signal_add_acquire, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_add_relaxed(signal, value)
    ccall((:hsa_signal_add_relaxed, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_add_screlease(signal, value)
    ccall((:hsa_signal_add_screlease, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_add_release(signal, value)
    ccall((:hsa_signal_add_release, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_subtract_scacq_screl(signal, value)
    ccall((:hsa_signal_subtract_scacq_screl, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_subtract_acq_rel(signal, value)
    ccall((:hsa_signal_subtract_acq_rel, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_subtract_scacquire(signal, value)
    ccall((:hsa_signal_subtract_scacquire, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_subtract_acquire(signal, value)
    ccall((:hsa_signal_subtract_acquire, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_subtract_relaxed(signal, value)
    ccall((:hsa_signal_subtract_relaxed, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_subtract_screlease(signal, value)
    ccall((:hsa_signal_subtract_screlease, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_subtract_release(signal, value)
    ccall((:hsa_signal_subtract_release, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_and_scacq_screl(signal, value)
    ccall((:hsa_signal_and_scacq_screl, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_and_acq_rel(signal, value)
    ccall((:hsa_signal_and_acq_rel, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_and_scacquire(signal, value)
    ccall((:hsa_signal_and_scacquire, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_and_acquire(signal, value)
    ccall((:hsa_signal_and_acquire, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_and_relaxed(signal, value)
    ccall((:hsa_signal_and_relaxed, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_and_screlease(signal, value)
    ccall((:hsa_signal_and_screlease, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_and_release(signal, value)
    ccall((:hsa_signal_and_release, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_or_scacq_screl(signal, value)
    ccall((:hsa_signal_or_scacq_screl, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_or_acq_rel(signal, value)
    ccall((:hsa_signal_or_acq_rel, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_or_scacquire(signal, value)
    ccall((:hsa_signal_or_scacquire, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_or_acquire(signal, value)
    ccall((:hsa_signal_or_acquire, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_or_relaxed(signal, value)
    ccall((:hsa_signal_or_relaxed, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_or_screlease(signal, value)
    ccall((:hsa_signal_or_screlease, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_or_release(signal, value)
    ccall((:hsa_signal_or_release, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_xor_scacq_screl(signal, value)
    ccall((:hsa_signal_xor_scacq_screl, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_xor_acq_rel(signal, value)
    ccall((:hsa_signal_xor_acq_rel, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_xor_scacquire(signal, value)
    ccall((:hsa_signal_xor_scacquire, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_xor_acquire(signal, value)
    ccall((:hsa_signal_xor_acquire, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_xor_relaxed(signal, value)
    ccall((:hsa_signal_xor_relaxed, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_xor_screlease(signal, value)
    ccall((:hsa_signal_xor_screlease, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

function hsa_signal_xor_release(signal, value)
    ccall((:hsa_signal_xor_release, libhsaruntime), Cvoid, (hsa_signal_t, hsa_signal_value_t), signal, value)
end

@cenum hsa_signal_condition_t::UInt32 begin
    HSA_SIGNAL_CONDITION_EQ = 0
    HSA_SIGNAL_CONDITION_NE = 1
    HSA_SIGNAL_CONDITION_LT = 2
    HSA_SIGNAL_CONDITION_GTE = 3
end

@cenum hsa_wait_state_t::UInt32 begin
    HSA_WAIT_STATE_BLOCKED = 0
    HSA_WAIT_STATE_ACTIVE = 1
end

function hsa_signal_wait_scacquire(signal, condition, compare_value, timeout_hint, wait_state_hint)
    ccall((:hsa_signal_wait_scacquire, libhsaruntime), hsa_signal_value_t, (hsa_signal_t, hsa_signal_condition_t, hsa_signal_value_t, UInt64, hsa_wait_state_t), signal, condition, compare_value, timeout_hint, wait_state_hint)
end

function hsa_signal_wait_relaxed(signal, condition, compare_value, timeout_hint, wait_state_hint)
    ccall((:hsa_signal_wait_relaxed, libhsaruntime), hsa_signal_value_t, (hsa_signal_t, hsa_signal_condition_t, hsa_signal_value_t, UInt64, hsa_wait_state_t), signal, condition, compare_value, timeout_hint, wait_state_hint)
end

function hsa_signal_wait_acquire(signal, condition, compare_value, timeout_hint, wait_state_hint)
    ccall((:hsa_signal_wait_acquire, libhsaruntime), hsa_signal_value_t, (hsa_signal_t, hsa_signal_condition_t, hsa_signal_value_t, UInt64, hsa_wait_state_t), signal, condition, compare_value, timeout_hint, wait_state_hint)
end

struct hsa_signal_group_s
    handle::UInt64
end

const hsa_signal_group_t = hsa_signal_group_s

function hsa_signal_group_create(num_signals, signals, num_consumers, consumers, signal_group)
    ccall((:hsa_signal_group_create, libhsaruntime), hsa_status_t, (UInt32, Ptr{hsa_signal_t}, UInt32, Ptr{hsa_agent_t}, Ptr{hsa_signal_group_t}), num_signals, signals, num_consumers, consumers, signal_group)
end

function hsa_signal_group_destroy(signal_group)
    ccall((:hsa_signal_group_destroy, libhsaruntime), hsa_status_t, (hsa_signal_group_t,), signal_group)
end

function hsa_signal_group_wait_any_scacquire(signal_group, conditions, compare_values, wait_state_hint, signal, value)
    ccall((:hsa_signal_group_wait_any_scacquire, libhsaruntime), hsa_status_t, (hsa_signal_group_t, Ptr{hsa_signal_condition_t}, Ptr{hsa_signal_value_t}, hsa_wait_state_t, Ptr{hsa_signal_t}, Ptr{hsa_signal_value_t}), signal_group, conditions, compare_values, wait_state_hint, signal, value)
end

function hsa_signal_group_wait_any_relaxed(signal_group, conditions, compare_values, wait_state_hint, signal, value)
    ccall((:hsa_signal_group_wait_any_relaxed, libhsaruntime), hsa_status_t, (hsa_signal_group_t, Ptr{hsa_signal_condition_t}, Ptr{hsa_signal_value_t}, hsa_wait_state_t, Ptr{hsa_signal_t}, Ptr{hsa_signal_value_t}), signal_group, conditions, compare_values, wait_state_hint, signal, value)
end

struct hsa_region_s
    handle::UInt64
end

const hsa_region_t = hsa_region_s

@cenum hsa_queue_type_t::UInt32 begin
    HSA_QUEUE_TYPE_MULTI = 0
    HSA_QUEUE_TYPE_SINGLE = 1
    HSA_QUEUE_TYPE_COOPERATIVE = 2
end

const hsa_queue_type32_t = UInt32

@cenum hsa_queue_feature_t::UInt32 begin
    HSA_QUEUE_FEATURE_KERNEL_DISPATCH = 1
    HSA_QUEUE_FEATURE_AGENT_DISPATCH = 2
end

struct hsa_queue_s
    type::hsa_queue_type32_t
    features::UInt32
    base_address::Ptr{Cvoid}
    doorbell_signal::hsa_signal_t
    size::UInt32
    reserved1::UInt32
    id::UInt64
end

const hsa_queue_t = hsa_queue_s

function hsa_queue_create(agent, size, type, callback, data, private_segment_size, group_segment_size, queue)
    ccall((:hsa_queue_create, libhsaruntime), hsa_status_t, (hsa_agent_t, UInt32, hsa_queue_type32_t, Ptr{Cvoid}, Ptr{Cvoid}, UInt32, UInt32, Ptr{Ptr{hsa_queue_t}}), agent, size, type, callback, data, private_segment_size, group_segment_size, queue)
end

function hsa_soft_queue_create(region, size, type, features, doorbell_signal, queue)
    ccall((:hsa_soft_queue_create, libhsaruntime), hsa_status_t, (hsa_region_t, UInt32, hsa_queue_type32_t, UInt32, hsa_signal_t, Ptr{Ptr{hsa_queue_t}}), region, size, type, features, doorbell_signal, queue)
end

function hsa_queue_destroy(queue)
    ccall((:hsa_queue_destroy, libhsaruntime), hsa_status_t, (Ptr{hsa_queue_t},), queue)
end

function hsa_queue_inactivate(queue)
    ccall((:hsa_queue_inactivate, libhsaruntime), hsa_status_t, (Ptr{hsa_queue_t},), queue)
end

function hsa_queue_load_read_index_acquire(queue)
    ccall((:hsa_queue_load_read_index_acquire, libhsaruntime), UInt64, (Ptr{hsa_queue_t},), queue)
end

function hsa_queue_load_read_index_scacquire(queue)
    ccall((:hsa_queue_load_read_index_scacquire, libhsaruntime), UInt64, (Ptr{hsa_queue_t},), queue)
end

function hsa_queue_load_read_index_relaxed(queue)
    ccall((:hsa_queue_load_read_index_relaxed, libhsaruntime), UInt64, (Ptr{hsa_queue_t},), queue)
end

function hsa_queue_load_write_index_acquire(queue)
    ccall((:hsa_queue_load_write_index_acquire, libhsaruntime), UInt64, (Ptr{hsa_queue_t},), queue)
end

function hsa_queue_load_write_index_scacquire(queue)
    ccall((:hsa_queue_load_write_index_scacquire, libhsaruntime), UInt64, (Ptr{hsa_queue_t},), queue)
end

function hsa_queue_load_write_index_relaxed(queue)
    ccall((:hsa_queue_load_write_index_relaxed, libhsaruntime), UInt64, (Ptr{hsa_queue_t},), queue)
end

function hsa_queue_store_write_index_relaxed(queue, value)
    ccall((:hsa_queue_store_write_index_relaxed, libhsaruntime), Cvoid, (Ptr{hsa_queue_t}, UInt64), queue, value)
end

function hsa_queue_store_write_index_release(queue, value)
    ccall((:hsa_queue_store_write_index_release, libhsaruntime), Cvoid, (Ptr{hsa_queue_t}, UInt64), queue, value)
end

function hsa_queue_store_write_index_screlease(queue, value)
    ccall((:hsa_queue_store_write_index_screlease, libhsaruntime), Cvoid, (Ptr{hsa_queue_t}, UInt64), queue, value)
end

function hsa_queue_cas_write_index_acq_rel(queue, expected, value)
    ccall((:hsa_queue_cas_write_index_acq_rel, libhsaruntime), UInt64, (Ptr{hsa_queue_t}, UInt64, UInt64), queue, expected, value)
end

function hsa_queue_cas_write_index_scacq_screl(queue, expected, value)
    ccall((:hsa_queue_cas_write_index_scacq_screl, libhsaruntime), UInt64, (Ptr{hsa_queue_t}, UInt64, UInt64), queue, expected, value)
end

function hsa_queue_cas_write_index_acquire(queue, expected, value)
    ccall((:hsa_queue_cas_write_index_acquire, libhsaruntime), UInt64, (Ptr{hsa_queue_t}, UInt64, UInt64), queue, expected, value)
end

function hsa_queue_cas_write_index_scacquire(queue, expected, value)
    ccall((:hsa_queue_cas_write_index_scacquire, libhsaruntime), UInt64, (Ptr{hsa_queue_t}, UInt64, UInt64), queue, expected, value)
end

function hsa_queue_cas_write_index_relaxed(queue, expected, value)
    ccall((:hsa_queue_cas_write_index_relaxed, libhsaruntime), UInt64, (Ptr{hsa_queue_t}, UInt64, UInt64), queue, expected, value)
end

function hsa_queue_cas_write_index_release(queue, expected, value)
    ccall((:hsa_queue_cas_write_index_release, libhsaruntime), UInt64, (Ptr{hsa_queue_t}, UInt64, UInt64), queue, expected, value)
end

function hsa_queue_cas_write_index_screlease(queue, expected, value)
    ccall((:hsa_queue_cas_write_index_screlease, libhsaruntime), UInt64, (Ptr{hsa_queue_t}, UInt64, UInt64), queue, expected, value)
end

function hsa_queue_add_write_index_acq_rel(queue, value)
    ccall((:hsa_queue_add_write_index_acq_rel, libhsaruntime), UInt64, (Ptr{hsa_queue_t}, UInt64), queue, value)
end

function hsa_queue_add_write_index_scacq_screl(queue, value)
    ccall((:hsa_queue_add_write_index_scacq_screl, libhsaruntime), UInt64, (Ptr{hsa_queue_t}, UInt64), queue, value)
end

function hsa_queue_add_write_index_acquire(queue, value)
    ccall((:hsa_queue_add_write_index_acquire, libhsaruntime), UInt64, (Ptr{hsa_queue_t}, UInt64), queue, value)
end

function hsa_queue_add_write_index_scacquire(queue, value)
    ccall((:hsa_queue_add_write_index_scacquire, libhsaruntime), UInt64, (Ptr{hsa_queue_t}, UInt64), queue, value)
end

function hsa_queue_add_write_index_relaxed(queue, value)
    ccall((:hsa_queue_add_write_index_relaxed, libhsaruntime), UInt64, (Ptr{hsa_queue_t}, UInt64), queue, value)
end

function hsa_queue_add_write_index_release(queue, value)
    ccall((:hsa_queue_add_write_index_release, libhsaruntime), UInt64, (Ptr{hsa_queue_t}, UInt64), queue, value)
end

function hsa_queue_add_write_index_screlease(queue, value)
    ccall((:hsa_queue_add_write_index_screlease, libhsaruntime), UInt64, (Ptr{hsa_queue_t}, UInt64), queue, value)
end

function hsa_queue_store_read_index_relaxed(queue, value)
    ccall((:hsa_queue_store_read_index_relaxed, libhsaruntime), Cvoid, (Ptr{hsa_queue_t}, UInt64), queue, value)
end

function hsa_queue_store_read_index_release(queue, value)
    ccall((:hsa_queue_store_read_index_release, libhsaruntime), Cvoid, (Ptr{hsa_queue_t}, UInt64), queue, value)
end

function hsa_queue_store_read_index_screlease(queue, value)
    ccall((:hsa_queue_store_read_index_screlease, libhsaruntime), Cvoid, (Ptr{hsa_queue_t}, UInt64), queue, value)
end

@cenum hsa_packet_type_t::UInt32 begin
    HSA_PACKET_TYPE_VENDOR_SPECIFIC = 0
    HSA_PACKET_TYPE_INVALID = 1
    HSA_PACKET_TYPE_KERNEL_DISPATCH = 2
    HSA_PACKET_TYPE_BARRIER_AND = 3
    HSA_PACKET_TYPE_AGENT_DISPATCH = 4
    HSA_PACKET_TYPE_BARRIER_OR = 5
end

@cenum hsa_fence_scope_t::UInt32 begin
    HSA_FENCE_SCOPE_NONE = 0
    HSA_FENCE_SCOPE_AGENT = 1
    HSA_FENCE_SCOPE_SYSTEM = 2
end

@cenum hsa_packet_header_t::UInt32 begin
    HSA_PACKET_HEADER_TYPE = 0
    HSA_PACKET_HEADER_BARRIER = 8
    HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE = 9
    HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE = 9
    HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE = 11
    HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE = 11
end

@cenum hsa_packet_header_width_t::UInt32 begin
    HSA_PACKET_HEADER_WIDTH_TYPE = 8
    HSA_PACKET_HEADER_WIDTH_BARRIER = 1
    HSA_PACKET_HEADER_WIDTH_SCACQUIRE_FENCE_SCOPE = 2
    HSA_PACKET_HEADER_WIDTH_ACQUIRE_FENCE_SCOPE = 2
    HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE = 2
    HSA_PACKET_HEADER_WIDTH_RELEASE_FENCE_SCOPE = 2
end

@cenum hsa_kernel_dispatch_packet_setup_t::UInt32 begin
    HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS = 0
end

@cenum hsa_kernel_dispatch_packet_setup_width_t::UInt32 begin
    HSA_KERNEL_DISPATCH_PACKET_SETUP_WIDTH_DIMENSIONS = 2
end

struct hsa_kernel_dispatch_packet_s
    header::UInt16
    setup::UInt16
    workgroup_size_x::UInt16
    workgroup_size_y::UInt16
    workgroup_size_z::UInt16
    reserved0::UInt16
    grid_size_x::UInt32
    grid_size_y::UInt32
    grid_size_z::UInt32
    private_segment_size::UInt32
    group_segment_size::UInt32
    kernel_object::UInt64
    kernarg_address::Ptr{Cvoid}
    reserved2::UInt64
    completion_signal::hsa_signal_t
end

const hsa_kernel_dispatch_packet_t = hsa_kernel_dispatch_packet_s

struct hsa_agent_dispatch_packet_s
    header::UInt16
    type::UInt16
    reserved0::UInt32
    return_address::Ptr{Cvoid}
    arg::NTuple{4, UInt64}
    reserved2::UInt64
    completion_signal::hsa_signal_t
end

const hsa_agent_dispatch_packet_t = hsa_agent_dispatch_packet_s

struct hsa_barrier_and_packet_s
    header::UInt16
    reserved0::UInt16
    reserved1::UInt32
    dep_signal::NTuple{5, hsa_signal_t}
    reserved2::UInt64
    completion_signal::hsa_signal_t
end

const hsa_barrier_and_packet_t = hsa_barrier_and_packet_s

struct hsa_barrier_or_packet_s
    header::UInt16
    reserved0::UInt16
    reserved1::UInt32
    dep_signal::NTuple{5, hsa_signal_t}
    reserved2::UInt64
    completion_signal::hsa_signal_t
end

const hsa_barrier_or_packet_t = hsa_barrier_or_packet_s

@cenum hsa_region_segment_t::UInt32 begin
    HSA_REGION_SEGMENT_GLOBAL = 0
    HSA_REGION_SEGMENT_READONLY = 1
    HSA_REGION_SEGMENT_PRIVATE = 2
    HSA_REGION_SEGMENT_GROUP = 3
    HSA_REGION_SEGMENT_KERNARG = 4
end

@cenum hsa_region_global_flag_t::UInt32 begin
    HSA_REGION_GLOBAL_FLAG_KERNARG = 1
    HSA_REGION_GLOBAL_FLAG_FINE_GRAINED = 2
    HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED = 4
end

@cenum hsa_region_info_t::UInt32 begin
    HSA_REGION_INFO_SEGMENT = 0
    HSA_REGION_INFO_GLOBAL_FLAGS = 1
    HSA_REGION_INFO_SIZE = 2
    HSA_REGION_INFO_ALLOC_MAX_SIZE = 4
    HSA_REGION_INFO_ALLOC_MAX_PRIVATE_WORKGROUP_SIZE = 8
    HSA_REGION_INFO_RUNTIME_ALLOC_ALLOWED = 5
    HSA_REGION_INFO_RUNTIME_ALLOC_GRANULE = 6
    HSA_REGION_INFO_RUNTIME_ALLOC_ALIGNMENT = 7
end

function hsa_region_get_info(region, attribute, value)
    ccall((:hsa_region_get_info, libhsaruntime), hsa_status_t, (hsa_region_t, hsa_region_info_t, Ptr{Cvoid}), region, attribute, value)
end

function hsa_agent_iterate_regions(agent, callback, data)
    ccall((:hsa_agent_iterate_regions, libhsaruntime), hsa_status_t, (hsa_agent_t, Ptr{Cvoid}, Ptr{Cvoid}), agent, callback, data)
end

function hsa_memory_allocate(region, size, ptr)
    ccall((:hsa_memory_allocate, libhsaruntime), hsa_status_t, (hsa_region_t, Csize_t, Ptr{Ptr{Cvoid}}), region, size, ptr)
end

function hsa_memory_free(ptr)
    ccall((:hsa_memory_free, libhsaruntime), hsa_status_t, (Ptr{Cvoid},), ptr)
end

function hsa_memory_copy(dst, src, size)
    ccall((:hsa_memory_copy, libhsaruntime), hsa_status_t, (Ptr{Cvoid}, Ptr{Cvoid}, Csize_t), dst, src, size)
end

function hsa_memory_assign_agent(ptr, agent, access)
    ccall((:hsa_memory_assign_agent, libhsaruntime), hsa_status_t, (Ptr{Cvoid}, hsa_agent_t, hsa_access_permission_t), ptr, agent, access)
end

function hsa_memory_register(ptr, size)
    ccall((:hsa_memory_register, libhsaruntime), hsa_status_t, (Ptr{Cvoid}, Csize_t), ptr, size)
end

function hsa_memory_deregister(ptr, size)
    ccall((:hsa_memory_deregister, libhsaruntime), hsa_status_t, (Ptr{Cvoid}, Csize_t), ptr, size)
end

struct hsa_isa_s
    handle::UInt64
end

const hsa_isa_t = hsa_isa_s

function hsa_isa_from_name(name, isa)
    ccall((:hsa_isa_from_name, libhsaruntime), hsa_status_t, (Ptr{Cchar}, Ptr{hsa_isa_t}), name, isa)
end

function hsa_agent_iterate_isas(agent, callback, data)
    ccall((:hsa_agent_iterate_isas, libhsaruntime), hsa_status_t, (hsa_agent_t, Ptr{Cvoid}, Ptr{Cvoid}), agent, callback, data)
end

@cenum hsa_isa_info_t::UInt32 begin
    HSA_ISA_INFO_NAME_LENGTH = 0
    HSA_ISA_INFO_NAME = 1
    HSA_ISA_INFO_CALL_CONVENTION_COUNT = 2
    HSA_ISA_INFO_CALL_CONVENTION_INFO_WAVEFRONT_SIZE = 3
    HSA_ISA_INFO_CALL_CONVENTION_INFO_WAVEFRONTS_PER_COMPUTE_UNIT = 4
    HSA_ISA_INFO_MACHINE_MODELS = 5
    HSA_ISA_INFO_PROFILES = 6
    HSA_ISA_INFO_DEFAULT_FLOAT_ROUNDING_MODES = 7
    HSA_ISA_INFO_BASE_PROFILE_DEFAULT_FLOAT_ROUNDING_MODES = 8
    HSA_ISA_INFO_FAST_F16_OPERATION = 9
    HSA_ISA_INFO_WORKGROUP_MAX_DIM = 12
    HSA_ISA_INFO_WORKGROUP_MAX_SIZE = 13
    HSA_ISA_INFO_GRID_MAX_DIM = 14
    HSA_ISA_INFO_GRID_MAX_SIZE = 16
    HSA_ISA_INFO_FBARRIER_MAX_SIZE = 17
end

function hsa_isa_get_info(isa, attribute, index, value)
    ccall((:hsa_isa_get_info, libhsaruntime), hsa_status_t, (hsa_isa_t, hsa_isa_info_t, UInt32, Ptr{Cvoid}), isa, attribute, index, value)
end

function hsa_isa_get_info_alt(isa, attribute, value)
    ccall((:hsa_isa_get_info_alt, libhsaruntime), hsa_status_t, (hsa_isa_t, hsa_isa_info_t, Ptr{Cvoid}), isa, attribute, value)
end

function hsa_isa_get_exception_policies(isa, profile, mask)
    ccall((:hsa_isa_get_exception_policies, libhsaruntime), hsa_status_t, (hsa_isa_t, hsa_profile_t, Ptr{UInt16}), isa, profile, mask)
end

@cenum hsa_fp_type_t::UInt32 begin
    HSA_FP_TYPE_16 = 1
    HSA_FP_TYPE_32 = 2
    HSA_FP_TYPE_64 = 4
end

@cenum hsa_flush_mode_t::UInt32 begin
    HSA_FLUSH_MODE_FTZ = 1
    HSA_FLUSH_MODE_NON_FTZ = 2
end

@cenum hsa_round_method_t::UInt32 begin
    HSA_ROUND_METHOD_SINGLE = 1
    HSA_ROUND_METHOD_DOUBLE = 2
end

function hsa_isa_get_round_method(isa, fp_type, flush_mode, round_method)
    ccall((:hsa_isa_get_round_method, libhsaruntime), hsa_status_t, (hsa_isa_t, hsa_fp_type_t, hsa_flush_mode_t, Ptr{hsa_round_method_t}), isa, fp_type, flush_mode, round_method)
end

struct hsa_wavefront_s
    handle::UInt64
end

const hsa_wavefront_t = hsa_wavefront_s

@cenum hsa_wavefront_info_t::UInt32 begin
    HSA_WAVEFRONT_INFO_SIZE = 0
end

function hsa_wavefront_get_info(wavefront, attribute, value)
    ccall((:hsa_wavefront_get_info, libhsaruntime), hsa_status_t, (hsa_wavefront_t, hsa_wavefront_info_t, Ptr{Cvoid}), wavefront, attribute, value)
end

function hsa_isa_iterate_wavefronts(isa, callback, data)
    ccall((:hsa_isa_iterate_wavefronts, libhsaruntime), hsa_status_t, (hsa_isa_t, Ptr{Cvoid}, Ptr{Cvoid}), isa, callback, data)
end

function hsa_isa_compatible(code_object_isa, agent_isa, result)
    ccall((:hsa_isa_compatible, libhsaruntime), hsa_status_t, (hsa_isa_t, hsa_isa_t, Ptr{Bool}), code_object_isa, agent_isa, result)
end

struct hsa_code_object_reader_s
    handle::UInt64
end

const hsa_code_object_reader_t = hsa_code_object_reader_s

function hsa_code_object_reader_create_from_file(file, code_object_reader)
    ccall((:hsa_code_object_reader_create_from_file, libhsaruntime), hsa_status_t, (hsa_file_t, Ptr{hsa_code_object_reader_t}), file, code_object_reader)
end

function hsa_code_object_reader_create_from_memory(code_object, size, code_object_reader)
    ccall((:hsa_code_object_reader_create_from_memory, libhsaruntime), hsa_status_t, (Ptr{Cvoid}, Csize_t, Ptr{hsa_code_object_reader_t}), code_object, size, code_object_reader)
end

function hsa_code_object_reader_destroy(code_object_reader)
    ccall((:hsa_code_object_reader_destroy, libhsaruntime), hsa_status_t, (hsa_code_object_reader_t,), code_object_reader)
end

struct hsa_executable_s
    handle::UInt64
end

const hsa_executable_t = hsa_executable_s

@cenum hsa_executable_state_t::UInt32 begin
    HSA_EXECUTABLE_STATE_UNFROZEN = 0
    HSA_EXECUTABLE_STATE_FROZEN = 1
end

function hsa_executable_create(profile, executable_state, options, executable)
    ccall((:hsa_executable_create, libhsaruntime), hsa_status_t, (hsa_profile_t, hsa_executable_state_t, Ptr{Cchar}, Ptr{hsa_executable_t}), profile, executable_state, options, executable)
end

function hsa_executable_create_alt(profile, default_float_rounding_mode, options, executable)
    ccall((:hsa_executable_create_alt, libhsaruntime), hsa_status_t, (hsa_profile_t, hsa_default_float_rounding_mode_t, Ptr{Cchar}, Ptr{hsa_executable_t}), profile, default_float_rounding_mode, options, executable)
end

function hsa_executable_destroy(executable)
    ccall((:hsa_executable_destroy, libhsaruntime), hsa_status_t, (hsa_executable_t,), executable)
end

struct hsa_loaded_code_object_s
    handle::UInt64
end

const hsa_loaded_code_object_t = hsa_loaded_code_object_s

function hsa_executable_load_program_code_object(executable, code_object_reader, options, loaded_code_object)
    ccall((:hsa_executable_load_program_code_object, libhsaruntime), hsa_status_t, (hsa_executable_t, hsa_code_object_reader_t, Ptr{Cchar}, Ptr{hsa_loaded_code_object_t}), executable, code_object_reader, options, loaded_code_object)
end

function hsa_executable_load_agent_code_object(executable, agent, code_object_reader, options, loaded_code_object)
    ccall((:hsa_executable_load_agent_code_object, libhsaruntime), hsa_status_t, (hsa_executable_t, hsa_agent_t, hsa_code_object_reader_t, Ptr{Cchar}, Ptr{hsa_loaded_code_object_t}), executable, agent, code_object_reader, options, loaded_code_object)
end

function hsa_executable_freeze(executable, options)
    ccall((:hsa_executable_freeze, libhsaruntime), hsa_status_t, (hsa_executable_t, Ptr{Cchar}), executable, options)
end

@cenum hsa_executable_info_t::UInt32 begin
    HSA_EXECUTABLE_INFO_PROFILE = 1
    HSA_EXECUTABLE_INFO_STATE = 2
    HSA_EXECUTABLE_INFO_DEFAULT_FLOAT_ROUNDING_MODE = 3
end

function hsa_executable_get_info(executable, attribute, value)
    ccall((:hsa_executable_get_info, libhsaruntime), hsa_status_t, (hsa_executable_t, hsa_executable_info_t, Ptr{Cvoid}), executable, attribute, value)
end

function hsa_executable_global_variable_define(executable, variable_name, address)
    ccall((:hsa_executable_global_variable_define, libhsaruntime), hsa_status_t, (hsa_executable_t, Ptr{Cchar}, Ptr{Cvoid}), executable, variable_name, address)
end

function hsa_executable_agent_global_variable_define(executable, agent, variable_name, address)
    ccall((:hsa_executable_agent_global_variable_define, libhsaruntime), hsa_status_t, (hsa_executable_t, hsa_agent_t, Ptr{Cchar}, Ptr{Cvoid}), executable, agent, variable_name, address)
end

function hsa_executable_readonly_variable_define(executable, agent, variable_name, address)
    ccall((:hsa_executable_readonly_variable_define, libhsaruntime), hsa_status_t, (hsa_executable_t, hsa_agent_t, Ptr{Cchar}, Ptr{Cvoid}), executable, agent, variable_name, address)
end

function hsa_executable_validate(executable, result)
    ccall((:hsa_executable_validate, libhsaruntime), hsa_status_t, (hsa_executable_t, Ptr{UInt32}), executable, result)
end

function hsa_executable_validate_alt(executable, options, result)
    ccall((:hsa_executable_validate_alt, libhsaruntime), hsa_status_t, (hsa_executable_t, Ptr{Cchar}, Ptr{UInt32}), executable, options, result)
end

struct hsa_executable_symbol_s
    handle::UInt64
end

const hsa_executable_symbol_t = hsa_executable_symbol_s

function hsa_executable_get_symbol(executable, module_name, symbol_name, agent, call_convention, symbol)
    ccall((:hsa_executable_get_symbol, libhsaruntime), hsa_status_t, (hsa_executable_t, Ptr{Cchar}, Ptr{Cchar}, hsa_agent_t, Int32, Ptr{hsa_executable_symbol_t}), executable, module_name, symbol_name, agent, call_convention, symbol)
end

function hsa_executable_get_symbol_by_name(executable, symbol_name, agent, symbol)
    ccall((:hsa_executable_get_symbol_by_name, libhsaruntime), hsa_status_t, (hsa_executable_t, Ptr{Cchar}, Ptr{hsa_agent_t}, Ptr{hsa_executable_symbol_t}), executable, symbol_name, agent, symbol)
end

@cenum hsa_symbol_kind_t::UInt32 begin
    HSA_SYMBOL_KIND_VARIABLE = 0
    HSA_SYMBOL_KIND_KERNEL = 1
    HSA_SYMBOL_KIND_INDIRECT_FUNCTION = 2
end

@cenum hsa_symbol_linkage_t::UInt32 begin
    HSA_SYMBOL_LINKAGE_MODULE = 0
    HSA_SYMBOL_LINKAGE_PROGRAM = 1
end

@cenum hsa_variable_allocation_t::UInt32 begin
    HSA_VARIABLE_ALLOCATION_AGENT = 0
    HSA_VARIABLE_ALLOCATION_PROGRAM = 1
end

@cenum hsa_variable_segment_t::UInt32 begin
    HSA_VARIABLE_SEGMENT_GLOBAL = 0
    HSA_VARIABLE_SEGMENT_READONLY = 1
end

@cenum hsa_executable_symbol_info_t::UInt32 begin
    HSA_EXECUTABLE_SYMBOL_INFO_TYPE = 0
    HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH = 1
    HSA_EXECUTABLE_SYMBOL_INFO_NAME = 2
    HSA_EXECUTABLE_SYMBOL_INFO_MODULE_NAME_LENGTH = 3
    HSA_EXECUTABLE_SYMBOL_INFO_MODULE_NAME = 4
    HSA_EXECUTABLE_SYMBOL_INFO_AGENT = 20
    HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS = 21
    HSA_EXECUTABLE_SYMBOL_INFO_LINKAGE = 5
    HSA_EXECUTABLE_SYMBOL_INFO_IS_DEFINITION = 17
    HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ALLOCATION = 6
    HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SEGMENT = 7
    HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ALIGNMENT = 8
    HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE = 9
    HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_IS_CONST = 10
    HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT = 22
    HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE = 11
    HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_ALIGNMENT = 12
    HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE = 13
    HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE = 14
    HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_DYNAMIC_CALLSTACK = 15
    HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_CALL_CONVENTION = 18
    HSA_EXECUTABLE_SYMBOL_INFO_INDIRECT_FUNCTION_OBJECT = 23
    HSA_EXECUTABLE_SYMBOL_INFO_INDIRECT_FUNCTION_CALL_CONVENTION = 16
end

function hsa_executable_symbol_get_info(executable_symbol, attribute, value)
    ccall((:hsa_executable_symbol_get_info, libhsaruntime), hsa_status_t, (hsa_executable_symbol_t, hsa_executable_symbol_info_t, Ptr{Cvoid}), executable_symbol, attribute, value)
end

function hsa_executable_iterate_symbols(executable, callback, data)
    ccall((:hsa_executable_iterate_symbols, libhsaruntime), hsa_status_t, (hsa_executable_t, Ptr{Cvoid}, Ptr{Cvoid}), executable, callback, data)
end

function hsa_executable_iterate_agent_symbols(executable, agent, callback, data)
    ccall((:hsa_executable_iterate_agent_symbols, libhsaruntime), hsa_status_t, (hsa_executable_t, hsa_agent_t, Ptr{Cvoid}, Ptr{Cvoid}), executable, agent, callback, data)
end

function hsa_executable_iterate_program_symbols(executable, callback, data)
    ccall((:hsa_executable_iterate_program_symbols, libhsaruntime), hsa_status_t, (hsa_executable_t, Ptr{Cvoid}, Ptr{Cvoid}), executable, callback, data)
end

struct hsa_code_object_s
    handle::UInt64
end

const hsa_code_object_t = hsa_code_object_s

struct hsa_callback_data_s
    handle::UInt64
end

const hsa_callback_data_t = hsa_callback_data_s

function hsa_code_object_serialize(code_object, alloc_callback, callback_data, options, serialized_code_object, serialized_code_object_size)
    ccall((:hsa_code_object_serialize, libhsaruntime), hsa_status_t, (hsa_code_object_t, Ptr{Cvoid}, hsa_callback_data_t, Ptr{Cchar}, Ptr{Ptr{Cvoid}}, Ptr{Csize_t}), code_object, alloc_callback, callback_data, options, serialized_code_object, serialized_code_object_size)
end

function hsa_code_object_deserialize(serialized_code_object, serialized_code_object_size, options, code_object)
    ccall((:hsa_code_object_deserialize, libhsaruntime), hsa_status_t, (Ptr{Cvoid}, Csize_t, Ptr{Cchar}, Ptr{hsa_code_object_t}), serialized_code_object, serialized_code_object_size, options, code_object)
end

function hsa_code_object_destroy(code_object)
    ccall((:hsa_code_object_destroy, libhsaruntime), hsa_status_t, (hsa_code_object_t,), code_object)
end

@cenum hsa_code_object_type_t::UInt32 begin
    HSA_CODE_OBJECT_TYPE_PROGRAM = 0
end

@cenum hsa_code_object_info_t::UInt32 begin
    HSA_CODE_OBJECT_INFO_VERSION = 0
    HSA_CODE_OBJECT_INFO_TYPE = 1
    HSA_CODE_OBJECT_INFO_ISA = 2
    HSA_CODE_OBJECT_INFO_MACHINE_MODEL = 3
    HSA_CODE_OBJECT_INFO_PROFILE = 4
    HSA_CODE_OBJECT_INFO_DEFAULT_FLOAT_ROUNDING_MODE = 5
end

function hsa_code_object_get_info(code_object, attribute, value)
    ccall((:hsa_code_object_get_info, libhsaruntime), hsa_status_t, (hsa_code_object_t, hsa_code_object_info_t, Ptr{Cvoid}), code_object, attribute, value)
end

function hsa_executable_load_code_object(executable, agent, code_object, options)
    ccall((:hsa_executable_load_code_object, libhsaruntime), hsa_status_t, (hsa_executable_t, hsa_agent_t, hsa_code_object_t, Ptr{Cchar}), executable, agent, code_object, options)
end

struct hsa_code_symbol_s
    handle::UInt64
end

const hsa_code_symbol_t = hsa_code_symbol_s

function hsa_code_object_get_symbol(code_object, symbol_name, symbol)
    ccall((:hsa_code_object_get_symbol, libhsaruntime), hsa_status_t, (hsa_code_object_t, Ptr{Cchar}, Ptr{hsa_code_symbol_t}), code_object, symbol_name, symbol)
end

function hsa_code_object_get_symbol_from_name(code_object, module_name, symbol_name, symbol)
    ccall((:hsa_code_object_get_symbol_from_name, libhsaruntime), hsa_status_t, (hsa_code_object_t, Ptr{Cchar}, Ptr{Cchar}, Ptr{hsa_code_symbol_t}), code_object, module_name, symbol_name, symbol)
end

@cenum hsa_code_symbol_info_t::UInt32 begin
    HSA_CODE_SYMBOL_INFO_TYPE = 0
    HSA_CODE_SYMBOL_INFO_NAME_LENGTH = 1
    HSA_CODE_SYMBOL_INFO_NAME = 2
    HSA_CODE_SYMBOL_INFO_MODULE_NAME_LENGTH = 3
    HSA_CODE_SYMBOL_INFO_MODULE_NAME = 4
    HSA_CODE_SYMBOL_INFO_LINKAGE = 5
    HSA_CODE_SYMBOL_INFO_IS_DEFINITION = 17
    HSA_CODE_SYMBOL_INFO_VARIABLE_ALLOCATION = 6
    HSA_CODE_SYMBOL_INFO_VARIABLE_SEGMENT = 7
    HSA_CODE_SYMBOL_INFO_VARIABLE_ALIGNMENT = 8
    HSA_CODE_SYMBOL_INFO_VARIABLE_SIZE = 9
    HSA_CODE_SYMBOL_INFO_VARIABLE_IS_CONST = 10
    HSA_CODE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE = 11
    HSA_CODE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_ALIGNMENT = 12
    HSA_CODE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE = 13
    HSA_CODE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE = 14
    HSA_CODE_SYMBOL_INFO_KERNEL_DYNAMIC_CALLSTACK = 15
    HSA_CODE_SYMBOL_INFO_KERNEL_CALL_CONVENTION = 18
    HSA_CODE_SYMBOL_INFO_INDIRECT_FUNCTION_CALL_CONVENTION = 16
end

function hsa_code_symbol_get_info(code_symbol, attribute, value)
    ccall((:hsa_code_symbol_get_info, libhsaruntime), hsa_status_t, (hsa_code_symbol_t, hsa_code_symbol_info_t, Ptr{Cvoid}), code_symbol, attribute, value)
end

function hsa_code_object_iterate_symbols(code_object, callback, data)
    ccall((:hsa_code_object_iterate_symbols, libhsaruntime), hsa_status_t, (hsa_code_object_t, Ptr{Cvoid}, Ptr{Cvoid}), code_object, callback, data)
end

@cenum var"##Ctag#317"::UInt32 begin
    HSA_EXT_STATUS_ERROR_IMAGE_FORMAT_UNSUPPORTED = 12288
    HSA_EXT_STATUS_ERROR_IMAGE_SIZE_UNSUPPORTED = 12289
    HSA_EXT_STATUS_ERROR_IMAGE_PITCH_UNSUPPORTED = 12290
    HSA_EXT_STATUS_ERROR_SAMPLER_DESCRIPTOR_UNSUPPORTED = 12291
end

@cenum var"##Ctag#318"::UInt32 begin
    HSA_EXT_AGENT_INFO_IMAGE_1D_MAX_ELEMENTS = 12288
    HSA_EXT_AGENT_INFO_IMAGE_1DA_MAX_ELEMENTS = 12289
    HSA_EXT_AGENT_INFO_IMAGE_1DB_MAX_ELEMENTS = 12290
    HSA_EXT_AGENT_INFO_IMAGE_2D_MAX_ELEMENTS = 12291
    HSA_EXT_AGENT_INFO_IMAGE_2DA_MAX_ELEMENTS = 12292
    HSA_EXT_AGENT_INFO_IMAGE_2DDEPTH_MAX_ELEMENTS = 12293
    HSA_EXT_AGENT_INFO_IMAGE_2DADEPTH_MAX_ELEMENTS = 12294
    HSA_EXT_AGENT_INFO_IMAGE_3D_MAX_ELEMENTS = 12295
    HSA_EXT_AGENT_INFO_IMAGE_ARRAY_MAX_LAYERS = 12296
    HSA_EXT_AGENT_INFO_MAX_IMAGE_RD_HANDLES = 12297
    HSA_EXT_AGENT_INFO_MAX_IMAGE_RORW_HANDLES = 12298
    HSA_EXT_AGENT_INFO_MAX_SAMPLER_HANDLERS = 12299
    HSA_EXT_AGENT_INFO_IMAGE_LINEAR_ROW_PITCH_ALIGNMENT = 12300
end

struct hsa_ext_image_s
    handle::UInt64
end

const hsa_ext_image_t = hsa_ext_image_s

@cenum hsa_ext_image_geometry_t::UInt32 begin
    HSA_EXT_IMAGE_GEOMETRY_1D = 0
    HSA_EXT_IMAGE_GEOMETRY_2D = 1
    HSA_EXT_IMAGE_GEOMETRY_3D = 2
    HSA_EXT_IMAGE_GEOMETRY_1DA = 3
    HSA_EXT_IMAGE_GEOMETRY_2DA = 4
    HSA_EXT_IMAGE_GEOMETRY_1DB = 5
    HSA_EXT_IMAGE_GEOMETRY_2DDEPTH = 6
    HSA_EXT_IMAGE_GEOMETRY_2DADEPTH = 7
end

@cenum hsa_ext_image_channel_type_t::UInt32 begin
    HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT8 = 0
    HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT16 = 1
    HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT8 = 2
    HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT16 = 3
    HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT24 = 4
    HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_SHORT_555 = 5
    HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_SHORT_565 = 6
    HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_SHORT_101010 = 7
    HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8 = 8
    HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16 = 9
    HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32 = 10
    HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8 = 11
    HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16 = 12
    HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32 = 13
    HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT = 14
    HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT = 15
end

const hsa_ext_image_channel_type32_t = UInt32

@cenum hsa_ext_image_channel_order_t::UInt32 begin
    HSA_EXT_IMAGE_CHANNEL_ORDER_A = 0
    HSA_EXT_IMAGE_CHANNEL_ORDER_R = 1
    HSA_EXT_IMAGE_CHANNEL_ORDER_RX = 2
    HSA_EXT_IMAGE_CHANNEL_ORDER_RG = 3
    HSA_EXT_IMAGE_CHANNEL_ORDER_RGX = 4
    HSA_EXT_IMAGE_CHANNEL_ORDER_RA = 5
    HSA_EXT_IMAGE_CHANNEL_ORDER_RGB = 6
    HSA_EXT_IMAGE_CHANNEL_ORDER_RGBX = 7
    HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA = 8
    HSA_EXT_IMAGE_CHANNEL_ORDER_BGRA = 9
    HSA_EXT_IMAGE_CHANNEL_ORDER_ARGB = 10
    HSA_EXT_IMAGE_CHANNEL_ORDER_ABGR = 11
    HSA_EXT_IMAGE_CHANNEL_ORDER_SRGB = 12
    HSA_EXT_IMAGE_CHANNEL_ORDER_SRGBX = 13
    HSA_EXT_IMAGE_CHANNEL_ORDER_SRGBA = 14
    HSA_EXT_IMAGE_CHANNEL_ORDER_SBGRA = 15
    HSA_EXT_IMAGE_CHANNEL_ORDER_INTENSITY = 16
    HSA_EXT_IMAGE_CHANNEL_ORDER_LUMINANCE = 17
    HSA_EXT_IMAGE_CHANNEL_ORDER_DEPTH = 18
    HSA_EXT_IMAGE_CHANNEL_ORDER_DEPTH_STENCIL = 19
end

const hsa_ext_image_channel_order32_t = UInt32

struct hsa_ext_image_format_s
    channel_type::hsa_ext_image_channel_type32_t
    channel_order::hsa_ext_image_channel_order32_t
end

const hsa_ext_image_format_t = hsa_ext_image_format_s

struct hsa_ext_image_descriptor_s
    geometry::hsa_ext_image_geometry_t
    width::Csize_t
    height::Csize_t
    depth::Csize_t
    array_size::Csize_t
    format::hsa_ext_image_format_t
end

const hsa_ext_image_descriptor_t = hsa_ext_image_descriptor_s

@cenum hsa_ext_image_capability_t::UInt32 begin
    HSA_EXT_IMAGE_CAPABILITY_NOT_SUPPORTED = 0
    HSA_EXT_IMAGE_CAPABILITY_READ_ONLY = 1
    HSA_EXT_IMAGE_CAPABILITY_WRITE_ONLY = 2
    HSA_EXT_IMAGE_CAPABILITY_READ_WRITE = 4
    HSA_EXT_IMAGE_CAPABILITY_READ_MODIFY_WRITE = 8
    HSA_EXT_IMAGE_CAPABILITY_ACCESS_INVARIANT_DATA_LAYOUT = 16
end

@cenum hsa_ext_image_data_layout_t::UInt32 begin
    HSA_EXT_IMAGE_DATA_LAYOUT_OPAQUE = 0
    HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR = 1
end

function hsa_ext_image_get_capability(agent, geometry, image_format, capability_mask)
    ccall((:hsa_ext_image_get_capability, libhsaruntime), hsa_status_t, (hsa_agent_t, hsa_ext_image_geometry_t, Ptr{hsa_ext_image_format_t}, Ptr{UInt32}), agent, geometry, image_format, capability_mask)
end

function hsa_ext_image_get_capability_with_layout(agent, geometry, image_format, image_data_layout, capability_mask)
    ccall((:hsa_ext_image_get_capability_with_layout, libhsaruntime), hsa_status_t, (hsa_agent_t, hsa_ext_image_geometry_t, Ptr{hsa_ext_image_format_t}, hsa_ext_image_data_layout_t, Ptr{UInt32}), agent, geometry, image_format, image_data_layout, capability_mask)
end

struct hsa_ext_image_data_info_s
    size::Csize_t
    alignment::Csize_t
end

const hsa_ext_image_data_info_t = hsa_ext_image_data_info_s

function hsa_ext_image_data_get_info(agent, image_descriptor, access_permission, image_data_info)
    ccall((:hsa_ext_image_data_get_info, libhsaruntime), hsa_status_t, (hsa_agent_t, Ptr{hsa_ext_image_descriptor_t}, hsa_access_permission_t, Ptr{hsa_ext_image_data_info_t}), agent, image_descriptor, access_permission, image_data_info)
end

function hsa_ext_image_data_get_info_with_layout(agent, image_descriptor, access_permission, image_data_layout, image_data_row_pitch, image_data_slice_pitch, image_data_info)
    ccall((:hsa_ext_image_data_get_info_with_layout, libhsaruntime), hsa_status_t, (hsa_agent_t, Ptr{hsa_ext_image_descriptor_t}, hsa_access_permission_t, hsa_ext_image_data_layout_t, Csize_t, Csize_t, Ptr{hsa_ext_image_data_info_t}), agent, image_descriptor, access_permission, image_data_layout, image_data_row_pitch, image_data_slice_pitch, image_data_info)
end

function hsa_ext_image_create(agent, image_descriptor, image_data, access_permission, image)
    ccall((:hsa_ext_image_create, libhsaruntime), hsa_status_t, (hsa_agent_t, Ptr{hsa_ext_image_descriptor_t}, Ptr{Cvoid}, hsa_access_permission_t, Ptr{hsa_ext_image_t}), agent, image_descriptor, image_data, access_permission, image)
end

function hsa_ext_image_create_with_layout(agent, image_descriptor, image_data, access_permission, image_data_layout, image_data_row_pitch, image_data_slice_pitch, image)
    ccall((:hsa_ext_image_create_with_layout, libhsaruntime), hsa_status_t, (hsa_agent_t, Ptr{hsa_ext_image_descriptor_t}, Ptr{Cvoid}, hsa_access_permission_t, hsa_ext_image_data_layout_t, Csize_t, Csize_t, Ptr{hsa_ext_image_t}), agent, image_descriptor, image_data, access_permission, image_data_layout, image_data_row_pitch, image_data_slice_pitch, image)
end

function hsa_ext_image_destroy(agent, image)
    ccall((:hsa_ext_image_destroy, libhsaruntime), hsa_status_t, (hsa_agent_t, hsa_ext_image_t), agent, image)
end

function hsa_ext_image_copy(agent, src_image, src_offset, dst_image, dst_offset, range)
    ccall((:hsa_ext_image_copy, libhsaruntime), hsa_status_t, (hsa_agent_t, hsa_ext_image_t, Ptr{hsa_dim3_t}, hsa_ext_image_t, Ptr{hsa_dim3_t}, Ptr{hsa_dim3_t}), agent, src_image, src_offset, dst_image, dst_offset, range)
end

struct hsa_ext_image_region_s
    offset::hsa_dim3_t
    range::hsa_dim3_t
end

const hsa_ext_image_region_t = hsa_ext_image_region_s

function hsa_ext_image_import(agent, src_memory, src_row_pitch, src_slice_pitch, dst_image, image_region)
    ccall((:hsa_ext_image_import, libhsaruntime), hsa_status_t, (hsa_agent_t, Ptr{Cvoid}, Csize_t, Csize_t, hsa_ext_image_t, Ptr{hsa_ext_image_region_t}), agent, src_memory, src_row_pitch, src_slice_pitch, dst_image, image_region)
end

function hsa_ext_image_export(agent, src_image, dst_memory, dst_row_pitch, dst_slice_pitch, image_region)
    ccall((:hsa_ext_image_export, libhsaruntime), hsa_status_t, (hsa_agent_t, hsa_ext_image_t, Ptr{Cvoid}, Csize_t, Csize_t, Ptr{hsa_ext_image_region_t}), agent, src_image, dst_memory, dst_row_pitch, dst_slice_pitch, image_region)
end

function hsa_ext_image_clear(agent, image, data, image_region)
    ccall((:hsa_ext_image_clear, libhsaruntime), hsa_status_t, (hsa_agent_t, hsa_ext_image_t, Ptr{Cvoid}, Ptr{hsa_ext_image_region_t}), agent, image, data, image_region)
end

struct hsa_ext_sampler_s
    handle::UInt64
end

const hsa_ext_sampler_t = hsa_ext_sampler_s

@cenum hsa_ext_sampler_addressing_mode_t::UInt32 begin
    HSA_EXT_SAMPLER_ADDRESSING_MODE_UNDEFINED = 0
    HSA_EXT_SAMPLER_ADDRESSING_MODE_CLAMP_TO_EDGE = 1
    HSA_EXT_SAMPLER_ADDRESSING_MODE_CLAMP_TO_BORDER = 2
    HSA_EXT_SAMPLER_ADDRESSING_MODE_REPEAT = 3
    HSA_EXT_SAMPLER_ADDRESSING_MODE_MIRRORED_REPEAT = 4
end

const hsa_ext_sampler_addressing_mode32_t = UInt32

@cenum hsa_ext_sampler_coordinate_mode_t::UInt32 begin
    HSA_EXT_SAMPLER_COORDINATE_MODE_UNNORMALIZED = 0
    HSA_EXT_SAMPLER_COORDINATE_MODE_NORMALIZED = 1
end

const hsa_ext_sampler_coordinate_mode32_t = UInt32

@cenum hsa_ext_sampler_filter_mode_t::UInt32 begin
    HSA_EXT_SAMPLER_FILTER_MODE_NEAREST = 0
    HSA_EXT_SAMPLER_FILTER_MODE_LINEAR = 1
end

const hsa_ext_sampler_filter_mode32_t = UInt32

struct hsa_ext_sampler_descriptor_s
    coordinate_mode::hsa_ext_sampler_coordinate_mode32_t
    filter_mode::hsa_ext_sampler_filter_mode32_t
    address_mode::hsa_ext_sampler_addressing_mode32_t
end

const hsa_ext_sampler_descriptor_t = hsa_ext_sampler_descriptor_s

function hsa_ext_sampler_create(agent, sampler_descriptor, sampler)
    ccall((:hsa_ext_sampler_create, libhsaruntime), hsa_status_t, (hsa_agent_t, Ptr{hsa_ext_sampler_descriptor_t}, Ptr{hsa_ext_sampler_t}), agent, sampler_descriptor, sampler)
end

function hsa_ext_sampler_destroy(agent, sampler)
    ccall((:hsa_ext_sampler_destroy, libhsaruntime), hsa_status_t, (hsa_agent_t, hsa_ext_sampler_t), agent, sampler)
end

struct hsa_ext_images_1_00_pfn_s
    hsa_ext_image_get_capability::Ptr{Cvoid}
    hsa_ext_image_data_get_info::Ptr{Cvoid}
    hsa_ext_image_create::Ptr{Cvoid}
    hsa_ext_image_destroy::Ptr{Cvoid}
    hsa_ext_image_copy::Ptr{Cvoid}
    hsa_ext_image_import::Ptr{Cvoid}
    hsa_ext_image_export::Ptr{Cvoid}
    hsa_ext_image_clear::Ptr{Cvoid}
    hsa_ext_sampler_create::Ptr{Cvoid}
    hsa_ext_sampler_destroy::Ptr{Cvoid}
end

const hsa_ext_images_1_00_pfn_t = hsa_ext_images_1_00_pfn_s

struct hsa_ext_images_1_pfn_s
    hsa_ext_image_get_capability::Ptr{Cvoid}
    hsa_ext_image_data_get_info::Ptr{Cvoid}
    hsa_ext_image_create::Ptr{Cvoid}
    hsa_ext_image_destroy::Ptr{Cvoid}
    hsa_ext_image_copy::Ptr{Cvoid}
    hsa_ext_image_import::Ptr{Cvoid}
    hsa_ext_image_export::Ptr{Cvoid}
    hsa_ext_image_clear::Ptr{Cvoid}
    hsa_ext_sampler_create::Ptr{Cvoid}
    hsa_ext_sampler_destroy::Ptr{Cvoid}
    hsa_ext_image_get_capability_with_layout::Ptr{Cvoid}
    hsa_ext_image_data_get_info_with_layout::Ptr{Cvoid}
    hsa_ext_image_create_with_layout::Ptr{Cvoid}
end

const hsa_ext_images_1_pfn_t = hsa_ext_images_1_pfn_s

const hsa_signal_condition32_t = UInt32

@cenum hsa_amd_packet_type_t::UInt32 begin
    HSA_AMD_PACKET_TYPE_BARRIER_VALUE = 2
end

const hsa_amd_packet_type8_t = UInt8

struct hsa_amd_packet_header_s
    header::UInt16
    AmdFormat::hsa_amd_packet_type8_t
    reserved::UInt8
end

const hsa_amd_vendor_packet_header_t = hsa_amd_packet_header_s

struct hsa_amd_barrier_value_packet_s
    header::hsa_amd_vendor_packet_header_t
    reserved0::UInt32
    signal::hsa_signal_t
    value::hsa_signal_value_t
    mask::hsa_signal_value_t
    cond::hsa_signal_condition32_t
    reserved1::UInt32
    reserved2::UInt64
    reserved3::UInt64
    completion_signal::hsa_signal_t
end

const hsa_amd_barrier_value_packet_t = hsa_amd_barrier_value_packet_s

@cenum var"##Ctag#328"::UInt32 begin
    HSA_STATUS_ERROR_INVALID_MEMORY_POOL = 40
    HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION = 41
    HSA_STATUS_ERROR_ILLEGAL_INSTRUCTION = 42
end

@cenum hsa_amd_agent_info_s::UInt32 begin
    HSA_AMD_AGENT_INFO_CHIP_ID = 40960
    HSA_AMD_AGENT_INFO_CACHELINE_SIZE = 40961
    HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT = 40962
    HSA_AMD_AGENT_INFO_MAX_CLOCK_FREQUENCY = 40963
    HSA_AMD_AGENT_INFO_DRIVER_NODE_ID = 40964
    HSA_AMD_AGENT_INFO_MAX_ADDRESS_WATCH_POINTS = 40965
    HSA_AMD_AGENT_INFO_BDFID = 40966
    HSA_AMD_AGENT_INFO_MEMORY_WIDTH = 40967
    HSA_AMD_AGENT_INFO_MEMORY_MAX_FREQUENCY = 40968
    HSA_AMD_AGENT_INFO_PRODUCT_NAME = 40969
    HSA_AMD_AGENT_INFO_MAX_WAVES_PER_CU = 40970
    HSA_AMD_AGENT_INFO_NUM_SIMDS_PER_CU = 40971
    HSA_AMD_AGENT_INFO_NUM_SHADER_ENGINES = 40972
    HSA_AMD_AGENT_INFO_NUM_SHADER_ARRAYS_PER_SE = 40973
    HSA_AMD_AGENT_INFO_HDP_FLUSH = 40974
    HSA_AMD_AGENT_INFO_DOMAIN = 40975
    HSA_AMD_AGENT_INFO_COOPERATIVE_QUEUES = 40976
    HSA_AMD_AGENT_INFO_UUID = 40977
    HSA_AMD_AGENT_INFO_ASIC_REVISION = 40978
end

const hsa_amd_agent_info_t = hsa_amd_agent_info_s

struct hsa_amd_hdp_flush_s
    HDP_MEM_FLUSH_CNTL::Ptr{UInt32}
    HDP_REG_FLUSH_CNTL::Ptr{UInt32}
end

const hsa_amd_hdp_flush_t = hsa_amd_hdp_flush_s

@cenum hsa_amd_region_info_s::UInt32 begin
    HSA_AMD_REGION_INFO_HOST_ACCESSIBLE = 40960
    HSA_AMD_REGION_INFO_BASE = 40961
    HSA_AMD_REGION_INFO_BUS_WIDTH = 40962
    HSA_AMD_REGION_INFO_MAX_CLOCK_FREQUENCY = 40963
end

const hsa_amd_region_info_t = hsa_amd_region_info_s

@cenum hsa_amd_coherency_type_s::UInt32 begin
    HSA_AMD_COHERENCY_TYPE_COHERENT = 0
    HSA_AMD_COHERENCY_TYPE_NONCOHERENT = 1
end

const hsa_amd_coherency_type_t = hsa_amd_coherency_type_s

function hsa_amd_coherency_get_type(agent, type)
    ccall((:hsa_amd_coherency_get_type, libhsaruntime), hsa_status_t, (hsa_agent_t, Ptr{hsa_amd_coherency_type_t}), agent, type)
end

function hsa_amd_coherency_set_type(agent, type)
    ccall((:hsa_amd_coherency_set_type, libhsaruntime), hsa_status_t, (hsa_agent_t, hsa_amd_coherency_type_t), agent, type)
end

struct hsa_amd_profiling_dispatch_time_s
    start::UInt64
    _end::UInt64
end

const hsa_amd_profiling_dispatch_time_t = hsa_amd_profiling_dispatch_time_s

struct hsa_amd_profiling_async_copy_time_s
    start::UInt64
    _end::UInt64
end

const hsa_amd_profiling_async_copy_time_t = hsa_amd_profiling_async_copy_time_s

function hsa_amd_profiling_set_profiler_enabled(queue, enable)
    ccall((:hsa_amd_profiling_set_profiler_enabled, libhsaruntime), hsa_status_t, (Ptr{hsa_queue_t}, Cint), queue, enable)
end

function hsa_amd_profiling_async_copy_enable(enable)
    ccall((:hsa_amd_profiling_async_copy_enable, libhsaruntime), hsa_status_t, (Bool,), enable)
end

function hsa_amd_profiling_get_dispatch_time(agent, signal, time)
    ccall((:hsa_amd_profiling_get_dispatch_time, libhsaruntime), hsa_status_t, (hsa_agent_t, hsa_signal_t, Ptr{hsa_amd_profiling_dispatch_time_t}), agent, signal, time)
end

function hsa_amd_profiling_get_async_copy_time(signal, time)
    ccall((:hsa_amd_profiling_get_async_copy_time, libhsaruntime), hsa_status_t, (hsa_signal_t, Ptr{hsa_amd_profiling_async_copy_time_t}), signal, time)
end

function hsa_amd_profiling_convert_tick_to_system_domain(agent, agent_tick, system_tick)
    ccall((:hsa_amd_profiling_convert_tick_to_system_domain, libhsaruntime), hsa_status_t, (hsa_agent_t, UInt64, Ptr{UInt64}), agent, agent_tick, system_tick)
end

@cenum hsa_amd_signal_attribute_t::UInt32 begin
    HSA_AMD_SIGNAL_AMD_GPU_ONLY = 1
    HSA_AMD_SIGNAL_IPC = 2
end

function hsa_amd_signal_create(initial_value, num_consumers, consumers, attributes, signal)
    ccall((:hsa_amd_signal_create, libhsaruntime), hsa_status_t, (hsa_signal_value_t, UInt32, Ptr{hsa_agent_t}, UInt64, Ptr{hsa_signal_t}), initial_value, num_consumers, consumers, attributes, signal)
end

function hsa_amd_signal_value_pointer(signal, value_ptr)
    ccall((:hsa_amd_signal_value_pointer, libhsaruntime), hsa_status_t, (hsa_signal_t, Ptr{Ptr{hsa_signal_value_t}}), signal, value_ptr)
end

# typedef bool ( * hsa_amd_signal_handler ) ( hsa_signal_value_t value , void * arg )
const hsa_amd_signal_handler = Ptr{Cvoid}

function hsa_amd_signal_async_handler(signal, cond, value, handler, arg)
    ccall((:hsa_amd_signal_async_handler, libhsaruntime), hsa_status_t, (hsa_signal_t, hsa_signal_condition_t, hsa_signal_value_t, hsa_amd_signal_handler, Ptr{Cvoid}), signal, cond, value, handler, arg)
end

function hsa_amd_async_function(callback, arg)
    ccall((:hsa_amd_async_function, libhsaruntime), hsa_status_t, (Ptr{Cvoid}, Ptr{Cvoid}), callback, arg)
end

function hsa_amd_signal_wait_any(signal_count, signals, conds, values, timeout_hint, wait_hint, satisfying_value)
    ccall((:hsa_amd_signal_wait_any, libhsaruntime), UInt32, (UInt32, Ptr{hsa_signal_t}, Ptr{hsa_signal_condition_t}, Ptr{hsa_signal_value_t}, UInt64, hsa_wait_state_t, Ptr{hsa_signal_value_t}), signal_count, signals, conds, values, timeout_hint, wait_hint, satisfying_value)
end

function hsa_amd_image_get_info_max_dim(agent, attribute, value)
    ccall((:hsa_amd_image_get_info_max_dim, libhsaruntime), hsa_status_t, (hsa_agent_t, hsa_agent_info_t, Ptr{Cvoid}), agent, attribute, value)
end

function hsa_amd_queue_cu_set_mask(queue, num_cu_mask_count, cu_mask)
    ccall((:hsa_amd_queue_cu_set_mask, libhsaruntime), hsa_status_t, (Ptr{hsa_queue_t}, UInt32, Ptr{UInt32}), queue, num_cu_mask_count, cu_mask)
end

@cenum hsa_amd_segment_t::UInt32 begin
    HSA_AMD_SEGMENT_GLOBAL = 0
    HSA_AMD_SEGMENT_READONLY = 1
    HSA_AMD_SEGMENT_PRIVATE = 2
    HSA_AMD_SEGMENT_GROUP = 3
end

struct hsa_amd_memory_pool_s
    handle::UInt64
end

const hsa_amd_memory_pool_t = hsa_amd_memory_pool_s

@cenum hsa_amd_memory_pool_global_flag_s::UInt32 begin
    HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT = 1
    HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_FINE_GRAINED = 2
    HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED = 4
end

const hsa_amd_memory_pool_global_flag_t = hsa_amd_memory_pool_global_flag_s

@cenum hsa_amd_memory_pool_info_t::UInt32 begin
    HSA_AMD_MEMORY_POOL_INFO_SEGMENT = 0
    HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS = 1
    HSA_AMD_MEMORY_POOL_INFO_SIZE = 2
    HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALLOWED = 5
    HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE = 6
    HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALIGNMENT = 7
    HSA_AMD_MEMORY_POOL_INFO_ACCESSIBLE_BY_ALL = 15
    HSA_AMD_MEMORY_POOL_INFO_ALLOC_MAX_SIZE = 16
end

function hsa_amd_memory_pool_get_info(memory_pool, attribute, value)
    ccall((:hsa_amd_memory_pool_get_info, libhsaruntime), hsa_status_t, (hsa_amd_memory_pool_t, hsa_amd_memory_pool_info_t, Ptr{Cvoid}), memory_pool, attribute, value)
end

function hsa_amd_agent_iterate_memory_pools(agent, callback, data)
    ccall((:hsa_amd_agent_iterate_memory_pools, libhsaruntime), hsa_status_t, (hsa_agent_t, Ptr{Cvoid}, Ptr{Cvoid}), agent, callback, data)
end

function hsa_amd_memory_pool_allocate(memory_pool, size, flags, ptr)
    ccall((:hsa_amd_memory_pool_allocate, libhsaruntime), hsa_status_t, (hsa_amd_memory_pool_t, Csize_t, UInt32, Ptr{Ptr{Cvoid}}), memory_pool, size, flags, ptr)
end

function hsa_amd_memory_pool_free(ptr)
    ccall((:hsa_amd_memory_pool_free, libhsaruntime), hsa_status_t, (Ptr{Cvoid},), ptr)
end

function hsa_amd_memory_async_copy(dst, dst_agent, src, src_agent, size, num_dep_signals, dep_signals, completion_signal)
    ccall((:hsa_amd_memory_async_copy, libhsaruntime), hsa_status_t, (Ptr{Cvoid}, hsa_agent_t, Ptr{Cvoid}, hsa_agent_t, Csize_t, UInt32, Ptr{hsa_signal_t}, hsa_signal_t), dst, dst_agent, src, src_agent, size, num_dep_signals, dep_signals, completion_signal)
end

struct hsa_pitched_ptr_s
    base::Ptr{Cvoid}
    pitch::Csize_t
    slice::Csize_t
end

const hsa_pitched_ptr_t = hsa_pitched_ptr_s

@cenum hsa_amd_copy_direction_t::UInt32 begin
    hsaHostToHost = 0
    hsaHostToDevice = 1
    hsaDeviceToHost = 2
    hsaDeviceToDevice = 3
end

function hsa_amd_memory_async_copy_rect(dst, dst_offset, src, src_offset, range, copy_agent, dir, num_dep_signals, dep_signals, completion_signal)
    ccall((:hsa_amd_memory_async_copy_rect, libhsaruntime), hsa_status_t, (Ptr{hsa_pitched_ptr_t}, Ptr{hsa_dim3_t}, Ptr{hsa_pitched_ptr_t}, Ptr{hsa_dim3_t}, Ptr{hsa_dim3_t}, hsa_agent_t, hsa_amd_copy_direction_t, UInt32, Ptr{hsa_signal_t}, hsa_signal_t), dst, dst_offset, src, src_offset, range, copy_agent, dir, num_dep_signals, dep_signals, completion_signal)
end

@cenum hsa_amd_memory_pool_access_t::UInt32 begin
    HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED = 0
    HSA_AMD_MEMORY_POOL_ACCESS_ALLOWED_BY_DEFAULT = 1
    HSA_AMD_MEMORY_POOL_ACCESS_DISALLOWED_BY_DEFAULT = 2
end

@cenum hsa_amd_link_info_type_t::UInt32 begin
    HSA_AMD_LINK_INFO_TYPE_HYPERTRANSPORT = 0
    HSA_AMD_LINK_INFO_TYPE_QPI = 1
    HSA_AMD_LINK_INFO_TYPE_PCIE = 2
    HSA_AMD_LINK_INFO_TYPE_INFINBAND = 3
    HSA_AMD_LINK_INFO_TYPE_XGMI = 4
end

struct hsa_amd_memory_pool_link_info_s
    min_latency::UInt32
    max_latency::UInt32
    min_bandwidth::UInt32
    max_bandwidth::UInt32
    atomic_support_32bit::Bool
    atomic_support_64bit::Bool
    coherent_support::Bool
    link_type::hsa_amd_link_info_type_t
    numa_distance::UInt32
end

const hsa_amd_memory_pool_link_info_t = hsa_amd_memory_pool_link_info_s

@cenum hsa_amd_agent_memory_pool_info_t::UInt32 begin
    HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS = 0
    HSA_AMD_AGENT_MEMORY_POOL_INFO_NUM_LINK_HOPS = 1
    HSA_AMD_AGENT_MEMORY_POOL_INFO_LINK_INFO = 2
end

function hsa_amd_agent_memory_pool_get_info(agent, memory_pool, attribute, value)
    ccall((:hsa_amd_agent_memory_pool_get_info, libhsaruntime), hsa_status_t, (hsa_agent_t, hsa_amd_memory_pool_t, hsa_amd_agent_memory_pool_info_t, Ptr{Cvoid}), agent, memory_pool, attribute, value)
end

function hsa_amd_agents_allow_access(num_agents, agents, flags, ptr)
    ccall((:hsa_amd_agents_allow_access, libhsaruntime), hsa_status_t, (UInt32, Ptr{hsa_agent_t}, Ptr{UInt32}, Ptr{Cvoid}), num_agents, agents, flags, ptr)
end

function hsa_amd_memory_pool_can_migrate(src_memory_pool, dst_memory_pool, result)
    ccall((:hsa_amd_memory_pool_can_migrate, libhsaruntime), hsa_status_t, (hsa_amd_memory_pool_t, hsa_amd_memory_pool_t, Ptr{Bool}), src_memory_pool, dst_memory_pool, result)
end

function hsa_amd_memory_migrate(ptr, memory_pool, flags)
    ccall((:hsa_amd_memory_migrate, libhsaruntime), hsa_status_t, (Ptr{Cvoid}, hsa_amd_memory_pool_t, UInt32), ptr, memory_pool, flags)
end

function hsa_amd_memory_lock(host_ptr, size, agents, num_agent, agent_ptr)
    ccall((:hsa_amd_memory_lock, libhsaruntime), hsa_status_t, (Ptr{Cvoid}, Csize_t, Ptr{hsa_agent_t}, Cint, Ptr{Ptr{Cvoid}}), host_ptr, size, agents, num_agent, agent_ptr)
end

function hsa_amd_memory_lock_to_pool(host_ptr, size, agents, num_agent, pool, flags, agent_ptr)
    ccall((:hsa_amd_memory_lock_to_pool, libhsaruntime), hsa_status_t, (Ptr{Cvoid}, Csize_t, Ptr{hsa_agent_t}, Cint, hsa_amd_memory_pool_t, UInt32, Ptr{Ptr{Cvoid}}), host_ptr, size, agents, num_agent, pool, flags, agent_ptr)
end

function hsa_amd_memory_unlock(host_ptr)
    ccall((:hsa_amd_memory_unlock, libhsaruntime), hsa_status_t, (Ptr{Cvoid},), host_ptr)
end

function hsa_amd_memory_fill(ptr, value, count)
    ccall((:hsa_amd_memory_fill, libhsaruntime), hsa_status_t, (Ptr{Cvoid}, UInt32, Csize_t), ptr, value, count)
end

function hsa_amd_interop_map_buffer(num_agents, agents, interop_handle, flags, size, ptr, metadata_size, metadata)
    ccall((:hsa_amd_interop_map_buffer, libhsaruntime), hsa_status_t, (UInt32, Ptr{hsa_agent_t}, Cint, UInt32, Ptr{Csize_t}, Ptr{Ptr{Cvoid}}, Ptr{Csize_t}, Ptr{Ptr{Cvoid}}), num_agents, agents, interop_handle, flags, size, ptr, metadata_size, metadata)
end

function hsa_amd_interop_unmap_buffer(ptr)
    ccall((:hsa_amd_interop_unmap_buffer, libhsaruntime), hsa_status_t, (Ptr{Cvoid},), ptr)
end

struct hsa_amd_image_descriptor_s
    version::UInt32
    deviceID::UInt32
    data::NTuple{1, UInt32}
end

const hsa_amd_image_descriptor_t = hsa_amd_image_descriptor_s

function hsa_amd_image_create(agent, image_descriptor, image_layout, image_data, access_permission, image)
    ccall((:hsa_amd_image_create, libhsaruntime), hsa_status_t, (hsa_agent_t, Ptr{hsa_ext_image_descriptor_t}, Ptr{hsa_amd_image_descriptor_t}, Ptr{Cvoid}, hsa_access_permission_t, Ptr{hsa_ext_image_t}), agent, image_descriptor, image_layout, image_data, access_permission, image)
end

@cenum hsa_amd_pointer_type_t::UInt32 begin
    HSA_EXT_POINTER_TYPE_UNKNOWN = 0
    HSA_EXT_POINTER_TYPE_HSA = 1
    HSA_EXT_POINTER_TYPE_LOCKED = 2
    HSA_EXT_POINTER_TYPE_GRAPHICS = 3
    HSA_EXT_POINTER_TYPE_IPC = 4
end

struct hsa_amd_pointer_info_s
    size::UInt32
    type::hsa_amd_pointer_type_t
    agentBaseAddress::Ptr{Cvoid}
    hostBaseAddress::Ptr{Cvoid}
    sizeInBytes::Csize_t
    userData::Ptr{Cvoid}
    agentOwner::hsa_agent_t
end

const hsa_amd_pointer_info_t = hsa_amd_pointer_info_s

function hsa_amd_pointer_info(ptr, info, alloc, num_agents_accessible, accessible)
    ccall((:hsa_amd_pointer_info, libhsaruntime), hsa_status_t, (Ptr{Cvoid}, Ptr{hsa_amd_pointer_info_t}, Ptr{Cvoid}, Ptr{UInt32}, Ptr{Ptr{hsa_agent_t}}), ptr, info, alloc, num_agents_accessible, accessible)
end

function hsa_amd_pointer_info_set_userdata(ptr, userdata)
    ccall((:hsa_amd_pointer_info_set_userdata, libhsaruntime), hsa_status_t, (Ptr{Cvoid}, Ptr{Cvoid}), ptr, userdata)
end

struct hsa_amd_ipc_memory_s
    handle::NTuple{8, UInt32}
end

const hsa_amd_ipc_memory_t = hsa_amd_ipc_memory_s

function hsa_amd_ipc_memory_create(ptr, len, handle)
    ccall((:hsa_amd_ipc_memory_create, libhsaruntime), hsa_status_t, (Ptr{Cvoid}, Csize_t, Ptr{hsa_amd_ipc_memory_t}), ptr, len, handle)
end

function hsa_amd_ipc_memory_attach(handle, len, num_agents, mapping_agents, mapped_ptr)
    ccall((:hsa_amd_ipc_memory_attach, libhsaruntime), hsa_status_t, (Ptr{hsa_amd_ipc_memory_t}, Csize_t, UInt32, Ptr{hsa_agent_t}, Ptr{Ptr{Cvoid}}), handle, len, num_agents, mapping_agents, mapped_ptr)
end

function hsa_amd_ipc_memory_detach(mapped_ptr)
    ccall((:hsa_amd_ipc_memory_detach, libhsaruntime), hsa_status_t, (Ptr{Cvoid},), mapped_ptr)
end

const hsa_amd_ipc_signal_t = hsa_amd_ipc_memory_t

function hsa_amd_ipc_signal_create(signal, handle)
    ccall((:hsa_amd_ipc_signal_create, libhsaruntime), hsa_status_t, (hsa_signal_t, Ptr{hsa_amd_ipc_signal_t}), signal, handle)
end

function hsa_amd_ipc_signal_attach(handle, signal)
    ccall((:hsa_amd_ipc_signal_attach, libhsaruntime), hsa_status_t, (Ptr{hsa_amd_ipc_signal_t}, Ptr{hsa_signal_t}), handle, signal)
end

@cenum hsa_amd_event_type_s::UInt32 begin
    HSA_AMD_GPU_MEMORY_FAULT_EVENT = 0
end

const hsa_amd_event_type_t = hsa_amd_event_type_s

@cenum hsa_amd_memory_fault_reason_t::Int32 begin
    HSA_AMD_MEMORY_FAULT_PAGE_NOT_PRESENT = 1
    HSA_AMD_MEMORY_FAULT_READ_ONLY = 2
    HSA_AMD_MEMORY_FAULT_NX = 4
    HSA_AMD_MEMORY_FAULT_HOST_ONLY = 8
    HSA_AMD_MEMORY_FAULT_DRAMECC = 16
    HSA_AMD_MEMORY_FAULT_IMPRECISE = 32
    HSA_AMD_MEMORY_FAULT_SRAMECC = 64
    HSA_AMD_MEMORY_FAULT_HANG = -2147483648
end

struct hsa_amd_gpu_memory_fault_info_s
    agent::hsa_agent_t
    virtual_address::UInt64
    fault_reason_mask::UInt32
end

const hsa_amd_gpu_memory_fault_info_t = hsa_amd_gpu_memory_fault_info_s

struct hsa_amd_event_s
    data::NTuple{32, UInt8}
end

function Base.getproperty(x::Ptr{hsa_amd_event_s}, f::Symbol)
    f === :event_type && return Ptr{hsa_amd_event_type_t}(x + 0)
    f === :memory_fault && return Ptr{hsa_amd_gpu_memory_fault_info_t}(x + 8)
    return getfield(x, f)
end

function Base.getproperty(x::hsa_amd_event_s, f::Symbol)
    r = Ref{hsa_amd_event_s}(x)
    ptr = Base.unsafe_convert(Ptr{hsa_amd_event_s}, r)
    fptr = getproperty(ptr, f)
    GC.@preserve r unsafe_load(fptr)
end

function Base.setproperty!(x::Ptr{hsa_amd_event_s}, f::Symbol, v)
    unsafe_store!(getproperty(x, f), v)
end

const hsa_amd_event_t = hsa_amd_event_s

# typedef hsa_status_t ( * hsa_amd_system_event_callback_t ) ( const hsa_amd_event_t * event , void * data )
const hsa_amd_system_event_callback_t = Ptr{Cvoid}

function hsa_amd_register_system_event_handler(callback, data)
    ccall((:hsa_amd_register_system_event_handler, libhsaruntime), hsa_status_t, (hsa_amd_system_event_callback_t, Ptr{Cvoid}), callback, data)
end

@cenum hsa_amd_queue_priority_s::UInt32 begin
    HSA_AMD_QUEUE_PRIORITY_LOW = 0
    HSA_AMD_QUEUE_PRIORITY_NORMAL = 1
    HSA_AMD_QUEUE_PRIORITY_HIGH = 2
end

const hsa_amd_queue_priority_t = hsa_amd_queue_priority_s

function hsa_amd_queue_set_priority(queue, priority)
    ccall((:hsa_amd_queue_set_priority, libhsaruntime), hsa_status_t, (Ptr{hsa_queue_t}, hsa_amd_queue_priority_t), queue, priority)
end

# typedef void ( * hsa_amd_deallocation_callback_t ) ( void * ptr , void * user_data )
const hsa_amd_deallocation_callback_t = Ptr{Cvoid}

function hsa_amd_register_deallocation_callback(ptr, callback, user_data)
    ccall((:hsa_amd_register_deallocation_callback, libhsaruntime), hsa_status_t, (Ptr{Cvoid}, hsa_amd_deallocation_callback_t, Ptr{Cvoid}), ptr, callback, user_data)
end

function hsa_amd_deregister_deallocation_callback(ptr, callback)
    ccall((:hsa_amd_deregister_deallocation_callback, libhsaruntime), hsa_status_t, (Ptr{Cvoid}, hsa_amd_deallocation_callback_t), ptr, callback)
end

# Skipping MacroDefinition: HSA_EXPORT_DECORATOR __attribute__ ( ( visibility ( "default" ) ) )

const HSA_API_EXPORT = HSA_EXPORT_DECORATOR(HSA_CALL)

const HSA_API_IMPORT = HSA_CALL

const HSA_API = HSA_API_IMPORT

const HSA_VERSION_1_0 = 1

const HSA_AMD_INTERFACE_VERSION_MAJOR = 1

const HSA_AMD_INTERFACE_VERSION_MINOR = 0

# exports
const PREFIXES = ["HSA", "hsa_"]
for name in names(@__MODULE__; all=true), prefix in PREFIXES
    if startswith(string(name), prefix)
        @eval export $name
    end
end

end # module