Skip to content

Commit

Permalink
Update the wrapper scripts.
Browse files Browse the repository at this point in the history
  • Loading branch information
@maleadt
maleadt committed Feb 6, 2020
1 parent 5cb8107 commit 9ccfc4b
Show file tree
Hide file tree
Showing 2 changed files with 77 additions and 47 deletions.
28 changes: 23 additions & 5 deletions res/wrap.jl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -103,21 +103,36 @@ mutable struct State
edits::Vector{Edit}
end

# insert `@check` before each `ccall` when it returns a checked type
# insert `@checked` before each function with a `ccall` returning a checked type`
const checked_types = [
"CUresult",
]
function insert_check(x, state)
if x isa CSTParser.EXPR && x.typ == CSTParser.Call && x.args[1].val == "ccall"
if x isa CSTParser.EXPR && x.typ == CSTParser.FunctionDef
_, def, body, _ = x.args
@assert body isa CSTParser.EXPR && body.typ == CSTParser.Block
@assert length(body.args) == 1

# Clang.jl-generated ccalls should be directly part of a function definition
call = body.args[1]
@assert call isa CSTParser.EXPR && call.typ == CSTParser.Call && call.args[1].val == "ccall"

# get the ccall return type
rv = x.args[5]
rv = call.args[5]

if rv.val in checked_types
push!(state.edits, Edit(state.offset, "@check "))
push!(state.edits, Edit(state.offset, "@checked "))
end
end
end

# rewrite ordinary `ccall`s to `@runtime_ccall`
function rewrite_ccall(x, state)
if x isa CSTParser.EXPR && x.typ == CSTParser.Call && x.args[1].val == "ccall"
push!(state.edits, Edit(state.offset, "@runtime_"))
end
end


## indenting passes

Expand Down Expand Up @@ -171,7 +186,7 @@ function wrap_at_comma(x, state, indent, offset, column)
end

function indent_ccall(x, state)
if x isa CSTParser.EXPR && x.typ == CSTParser.Call && x.args[1].val == "ccall"
if x isa CSTParser.EXPR && x.typ == CSTParser.MacroCall && x.args[1].args[2].val == "runtime_ccall"
# figure out how much to indent by looking at where the expr starts
line = findlast(y -> state.offset >= y[2], state.lines) # index, not the actual number
line_indent, line_offset = state.lines[line]
Expand Down Expand Up @@ -235,6 +250,9 @@ function process(name, headers...; kwargs...)
state.offset = 0
pass(ast, state, insert_check)

state.offset = 0
pass(ast, state, rewrite_ccall)

# apply
state.offset = 0
sort!(state.edits, lt = (a,b) -> first(a.loc) < first(b.loc), rev = true)
Expand Down
96 changes: 54 additions & 42 deletions src/libcuda.jl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -154,7 +154,7 @@ end

@checked function cuCtxSynchronize()
@runtime_ccall((:cuCtxSynchronize, libcuda), CUresult, ())
end
end

@checked function cuCtxSetLimit(limit, value)
@runtime_ccall((:cuCtxSetLimit, libcuda), CUresult,
Expand Down Expand Up @@ -276,18 +276,19 @@ end
numOptions, options, optionValues, stateOut)
end

@checked function cuLinkAddData_v2(state, type, data, size, name, numOptions, options, optionValues)
@checked function cuLinkAddData_v2(state, type, data, size, name, numOptions, options,
optionValues)
@runtime_ccall((:cuLinkAddData_v2, libcuda), CUresult,
(CUlinkState, CUjitInputType, Ptr{Cvoid}, Csize_t, Cstring, UInt32,
Ptr{CUjit_option}, Ptr{Ptr{Cvoid}}),
state, type, data, size, name, numOptions, options, optionValues)
state, type, data, size, name, numOptions, options, optionValues)
end

@checked function cuLinkAddFile_v2(state, type, path, numOptions, options, optionValues)
@runtime_ccall((:cuLinkAddFile_v2, libcuda), CUresult,
(CUlinkState, CUjitInputType, Cstring, UInt32, Ptr{CUjit_option},
Ptr{Ptr{Cvoid}}),
state, type, path, numOptions, options, optionValues)
state, type, path, numOptions, options, optionValues)
end

@checked function cuLinkComplete(state, cubinOut, sizeOut)
Expand Down Expand Up @@ -512,7 +513,8 @@ end
dst, src, ByteCount, hStream)
end

@checked function cuMemcpyPeerAsync(dstDevice, dstContext, srcDevice, srcContext, ByteCount, hStream)
@checked function cuMemcpyPeerAsync(dstDevice, dstContext, srcDevice, srcContext,
ByteCount, hStream)
@runtime_ccall((:cuMemcpyPeerAsync, libcuda), CUresult,
(CUdeviceptr, CUcontext, CUdeviceptr, CUcontext, Csize_t, CUstream),
dstDevice, dstContext, srcDevice, srcContext, ByteCount, hStream)
Expand Down Expand Up @@ -710,11 +712,12 @@ end
data, dataSize, attribute, devPtr, count)
end

@checked function cuMemRangeGetAttributes(data, dataSizes, attributes, numAttributes, devPtr, count)
@checked function cuMemRangeGetAttributes(data, dataSizes, attributes, numAttributes,
devPtr, count)
@runtime_ccall((:cuMemRangeGetAttributes, libcuda), CUresult,
(Ptr{Ptr{Cvoid}}, Ptr{Csize_t}, Ptr{CUmem_range_attribute}, Csize_t,
CUdeviceptr, Csize_t),
data, dataSizes, attributes, numAttributes, devPtr, count)
data, dataSizes, attributes, numAttributes, devPtr, count)
end

@checked function cuPointerSetAttribute(value, attribute, ptr)
Expand Down Expand Up @@ -871,14 +874,14 @@ end
@runtime_ccall((:cuExternalMemoryGetMappedBuffer, libcuda), CUresult,
(Ptr{CUdeviceptr}, CUexternalMemory,
Ptr{CUDA_EXTERNAL_MEMORY_BUFFER_DESC}),
devPtr, extMem, bufferDesc)
devPtr, extMem, bufferDesc)
end

@checked function cuExternalMemoryGetMappedMipmappedArray(mipmap, extMem, mipmapDesc)
@runtime_ccall((:cuExternalMemoryGetMappedMipmappedArray, libcuda), CUresult,
(Ptr{CUmipmappedArray}, CUexternalMemory,
Ptr{CUDA_EXTERNAL_MEMORY_MIPMAPPED_ARRAY_DESC}),
mipmap, extMem, mipmapDesc)
mipmap, extMem, mipmapDesc)
end

@checked function cuDestroyExternalMemory(extMem)
Expand All @@ -893,18 +896,19 @@ end
extSem_out, semHandleDesc)
end

@checked function cuSignalExternalSemaphoresAsync(extSemArray, paramsArray, numExtSems, stream)
@checked function cuSignalExternalSemaphoresAsync(extSemArray, paramsArray, numExtSems,
stream)
@runtime_ccall((:cuSignalExternalSemaphoresAsync, libcuda), CUresult,
(Ptr{CUexternalSemaphore}, Ptr{CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS},
UInt32, CUstream),
extSemArray, paramsArray, numExtSems, stream)
extSemArray, paramsArray, numExtSems, stream)
end

@checked function cuWaitExternalSemaphoresAsync(extSemArray, paramsArray, numExtSems, stream)
@runtime_ccall((:cuWaitExternalSemaphoresAsync, libcuda), CUresult,
(Ptr{CUexternalSemaphore}, Ptr{CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS},
UInt32, CUstream),
extSemArray, paramsArray, numExtSems, stream)
extSemArray, paramsArray, numExtSems, stream)
end

@checked function cuDestroyExternalSemaphore(extSem)
Expand Down Expand Up @@ -961,22 +965,23 @@ end
hfunc, config)
end

@checked function cuLaunchKernel(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ,
sharedMemBytes, hStream, kernelParams, extra)
@checked function cuLaunchKernel(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY,
blockDimZ, sharedMemBytes, hStream, kernelParams, extra)
@runtime_ccall((:cuLaunchKernel, libcuda), CUresult,
(CUfunction, UInt32, UInt32, UInt32, UInt32, UInt32, UInt32, UInt32,
CUstream, Ptr{Ptr{Cvoid}}, Ptr{Ptr{Cvoid}}),
f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ,
sharedMemBytes, hStream, kernelParams, extra)
f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ,
sharedMemBytes, hStream, kernelParams, extra)
end

@checked function cuLaunchCooperativeKernel(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY,
blockDimZ, sharedMemBytes, hStream, kernelParams)
@checked function cuLaunchCooperativeKernel(f, gridDimX, gridDimY, gridDimZ, blockDimX,
blockDimY, blockDimZ, sharedMemBytes, hStream,
kernelParams)
@runtime_ccall((:cuLaunchCooperativeKernel, libcuda), CUresult,
(CUfunction, UInt32, UInt32, UInt32, UInt32, UInt32, UInt32, UInt32,
CUstream, Ptr{Ptr{Cvoid}}),
f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ,
sharedMemBytes, hStream, kernelParams)
f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ,
sharedMemBytes, hStream, kernelParams)
end

@checked function cuLaunchCooperativeKernelMultiDevice(launchParamsList, numDevices, flags)
Expand Down Expand Up @@ -1057,7 +1062,8 @@ end
phGraph, flags)
end

@checked function cuGraphAddKernelNode(phGraphNode, hGraph, dependencies, numDependencies, nodeParams)
@checked function cuGraphAddKernelNode(phGraphNode, hGraph, dependencies, numDependencies,
nodeParams)
@runtime_ccall((:cuGraphAddKernelNode, libcuda), CUresult,
(Ptr{CUgraphNode}, CUgraph, Ptr{CUgraphNode}, Csize_t,
Ptr{CUDA_KERNEL_NODE_PARAMS}),
Expand Down Expand Up @@ -1097,11 +1103,11 @@ end
end

@checked function cuGraphAddMemsetNode(phGraphNode, hGraph, dependencies, numDependencies,
memsetParams, ctx)
memsetParams, ctx)
@runtime_ccall((:cuGraphAddMemsetNode, libcuda), CUresult,
(Ptr{CUgraphNode}, CUgraph, Ptr{CUgraphNode}, Csize_t,
Ptr{CUDA_MEMSET_NODE_PARAMS}, CUcontext),
phGraphNode, hGraph, dependencies, numDependencies, memsetParams, ctx)
phGraphNode, hGraph, dependencies, numDependencies, memsetParams, ctx)
end

@checked function cuGraphMemsetNodeGetParams(hNode, nodeParams)
Expand All @@ -1116,11 +1122,12 @@ end
hNode, nodeParams)
end

@checked function cuGraphAddHostNode(phGraphNode, hGraph, dependencies, numDependencies, nodeParams)
@checked function cuGraphAddHostNode(phGraphNode, hGraph, dependencies, numDependencies,
nodeParams)
@runtime_ccall((:cuGraphAddHostNode, libcuda), CUresult,
(Ptr{CUgraphNode}, CUgraph, Ptr{CUgraphNode}, Csize_t,
Ptr{CUDA_HOST_NODE_PARAMS}),
phGraphNode, hGraph, dependencies, numDependencies, nodeParams)
phGraphNode, hGraph, dependencies, numDependencies, nodeParams)
end

@checked function cuGraphHostNodeGetParams(hNode, nodeParams)
Expand All @@ -1135,8 +1142,8 @@ end
hNode, nodeParams)
end

@checked function cuGraphAddChildGraphNode(phGraphNode, hGraph, dependencies, numDependencies,
childGraph)
@checked function cuGraphAddChildGraphNode(phGraphNode, hGraph, dependencies,
numDependencies, childGraph)
@runtime_ccall((:cuGraphAddChildGraphNode, libcuda), CUresult,
(Ptr{CUgraphNode}, CUgraph, Ptr{CUgraphNode}, Csize_t, CUgraph),
phGraphNode, hGraph, dependencies, numDependencies, childGraph)
Expand Down Expand Up @@ -1257,30 +1264,33 @@ end
numBlocks, func, blockSize, dynamicSMemSize)
end

@checked function cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(numBlocks, func, blockSize,
dynamicSMemSize, flags)
@checked function cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(numBlocks, func,
blockSize,
dynamicSMemSize,
flags)
@runtime_ccall((:cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags, libcuda), CUresult,
(Ptr{Cint}, CUfunction, Cint, Csize_t, UInt32),
numBlocks, func, blockSize, dynamicSMemSize, flags)
end

@checked function cuOccupancyMaxPotentialBlockSize(minGridSize, blockSize, func,
blockSizeToDynamicSMemSize, dynamicSMemSize,
blockSizeLimit)
blockSizeToDynamicSMemSize,
dynamicSMemSize, blockSizeLimit)
@runtime_ccall((:cuOccupancyMaxPotentialBlockSize, libcuda), CUresult,
(Ptr{Cint}, Ptr{Cint}, CUfunction, CUoccupancyB2DSize, Csize_t, Cint),
minGridSize, blockSize, func, blockSizeToDynamicSMemSize, dynamicSMemSize,
blockSizeLimit)
minGridSize, blockSize, func, blockSizeToDynamicSMemSize,
dynamicSMemSize, blockSizeLimit)
end

@checked function cuOccupancyMaxPotentialBlockSizeWithFlags(minGridSize, blockSize, func,
blockSizeToDynamicSMemSize,
dynamicSMemSize, blockSizeLimit, flags)
blockSizeToDynamicSMemSize,
dynamicSMemSize,
blockSizeLimit, flags)
@runtime_ccall((:cuOccupancyMaxPotentialBlockSizeWithFlags, libcuda), CUresult,
(Ptr{Cint}, Ptr{Cint}, CUfunction, CUoccupancyB2DSize, Csize_t, Cint,
UInt32),
minGridSize, blockSize, func, blockSizeToDynamicSMemSize, dynamicSMemSize,
blockSizeLimit, flags)
minGridSize, blockSize, func, blockSizeToDynamicSMemSize,
dynamicSMemSize, blockSizeLimit, flags)
end

@checked function cuTexRefSetArray(hTexRef, hArray, Flags)
Expand Down Expand Up @@ -1337,7 +1347,8 @@ end
hTexRef, bias)
end

@checked function cuTexRefSetMipmapLevelClamp(hTexRef, minMipmapLevelClamp, maxMipmapLevelClamp)
@checked function cuTexRefSetMipmapLevelClamp(hTexRef, minMipmapLevelClamp,
maxMipmapLevelClamp)
@runtime_ccall((:cuTexRefSetMipmapLevelClamp, libcuda), CUresult,
(CUtexref, Cfloat, Cfloat),
hTexRef, minMipmapLevelClamp, maxMipmapLevelClamp)
Expand Down Expand Up @@ -1409,7 +1420,8 @@ end
pbias, hTexRef)
end

@checked function cuTexRefGetMipmapLevelClamp(pminMipmapLevelClamp, pmaxMipmapLevelClamp, hTexRef)
@checked function cuTexRefGetMipmapLevelClamp(pminMipmapLevelClamp, pmaxMipmapLevelClamp,
hTexRef)
@runtime_ccall((:cuTexRefGetMipmapLevelClamp, libcuda), CUresult,
(Ptr{Cfloat}, Ptr{Cfloat}, CUtexref),
pminMipmapLevelClamp, pmaxMipmapLevelClamp, hTexRef)
Expand Down Expand Up @@ -1461,7 +1473,7 @@ end
@runtime_ccall((:cuTexObjectCreate, libcuda), CUresult,
(Ptr{CUtexObject}, Ptr{CUDA_RESOURCE_DESC}, Ptr{CUDA_TEXTURE_DESC},
Ptr{CUDA_RESOURCE_VIEW_DESC}),
pTexObject, pResDesc, pTexDesc, pResViewDesc)
pTexObject, pResDesc, pTexDesc, pResViewDesc)
end

@checked function cuTexObjectDestroy(texObject)
Expand Down Expand Up @@ -1589,8 +1601,8 @@ end

@checked function cuProfilerStart()
@runtime_ccall((:cuProfilerStart, libcuda), CUresult, ())
end
end

@checked function cuProfilerStop()
@runtime_ccall((:cuProfilerStop, libcuda), CUresult, ())
end
end

0 comments on commit 9ccfc4b

Please sign in to comment.