added possibility to create registers
Some checks are pending
CI / Julia ${{ matrix.version }} - ${{ matrix.os }} - ${{ matrix.arch }} - ${{ github.event_name }} (x64, ubuntu-latest, 1.10) (push) Waiting to run
CI / Julia ${{ matrix.version }} - ${{ matrix.os }} - ${{ matrix.arch }} - ${{ github.event_name }} (x64, ubuntu-latest, 1.6) (push) Waiting to run
CI / Julia ${{ matrix.version }} - ${{ matrix.os }} - ${{ matrix.arch }} - ${{ github.event_name }} (x64, ubuntu-latest, pre) (push) Waiting to run

This commit is contained in:
Daniel 2024-10-19 13:44:38 +02:00
parent de73d83d9e
commit 1f6b40b750

View File

@ -87,14 +87,14 @@ function transpile(expression::ExpressionProcessing.PostfixType)
println(ptxBuffer, "{") println(ptxBuffer, "{")
# Register definition # Register definition
println(ptxBuffer, get_register_definitions(1, 5)) # TODO: Actually calculate the number of needed registers and extend to more register kinds
# Parameter loading # Parameter loading
println(ptxBuffer, get_guard_clause()) println(ptxBuffer, get_guard_clause())
# Code goes here # Code goes here
# exit jump location # exit jump location
print(ptxBuffer, exitJumpLocationMarker) print(ptxBuffer, exitJumpLocationMarker); println(ptxBuffer, ": ret;")
println(ptxBuffer, ": ret;")
println(ptxBuffer, "}") println(ptxBuffer, "}")
println(String(take!(ptxBuffer))) println(String(take!(ptxBuffer)))
end end
@ -132,7 +132,7 @@ end
Constructs the PTX code used for handling the case where too many threads are started. Constructs the PTX code used for handling the case where too many threads are started.
Assumes the following: Assumes the following:
- There are the unused ```32 bit``` registers ```r1, r2, r3, r4``` - There are the unused ```32 bit``` registers ```r1, r2, r3, r4 (index of the variable set)```
- There is an unused ```predicate``` register ```p1``` - There is an unused ```predicate``` register ```p1```
- The ```32 bit``` register ```r5``` contains the number of variable sets - The ```32 bit``` register ```r5``` contains the number of variable sets
" "
@ -154,6 +154,25 @@ function get_guard_clause()::String
return String(take!(guardBuffer)) return String(take!(guardBuffer))
end end
function get_register_definitions(nrPred::Int, nr32Bit::Int):::String
registersBuffer = IOBuffer()
if nrPred > 0
print(registersBuffer, ".reg .pred")
print(registersBuffer, " %p<")
print(registersBuffer, nrPred + 1)
println(registersBuffer, ">;")
end
if nr32Bit > 0
print(registersBuffer, ".reg .b32")
print(registersBuffer, " %r<")
print(registersBuffer, nr32Bit + 1)
println(registersBuffer, ">;")
end
return String(take!(registersBuffer))
end
function type_to_cuda_type(type::DataType)::String function type_to_cuda_type(type::DataType)::String
if type == Int64 if type == Int64
return ".s64" return ".s64"