diff options
| -rwxr-xr-x | configure.py | 3 | ||||
| -rw-r--r-- | tgsi/lib/OVERRIDES | 0 | ||||
| -rw-r--r-- | tgsi/lib/SOURCES | 1 | ||||
| -rw-r--r-- | tgsi/lib/workitem/get_local_id.ll | 18 |
4 files changed, 21 insertions, 1 deletions
diff --git a/configure.py b/configure.py index b96faee..59cd0a8 100755 --- a/configure.py +++ b/configure.py @@ -103,9 +103,10 @@ available_targets = { 'nvptx64--' : { 'devices' : [{'gpu' : '', 'aliases' : []} ]}, 'nvptx--nvidiacl' : { 'devices' : [{'gpu' : '', 'aliases' : []} ]}, 'nvptx64--nvidiacl' : { 'devices' : [{'gpu' : '', 'aliases' : []} ]}, + 'tgsi--' : { 'devices' : [{'gpu' : '', 'aliases' : []} ]}, } -default_targets = ['nvptx--nvidiacl', 'nvptx64--nvidiacl', 'r600--', 'amdgcn--', 'amdgcn--amdhsa'] +default_targets = ['nvptx--nvidiacl', 'nvptx64--nvidiacl', 'r600--', 'amdgcn--', 'amdgcn--amdhsa', 'tgsi--'] targets = args if not targets: diff --git a/tgsi/lib/OVERRIDES b/tgsi/lib/OVERRIDES new file mode 100644 index 0000000..e69de29 --- /dev/null +++ b/tgsi/lib/OVERRIDES diff --git a/tgsi/lib/SOURCES b/tgsi/lib/SOURCES new file mode 100644 index 0000000..c116b86 --- /dev/null +++ b/tgsi/lib/SOURCES @@ -0,0 +1 @@ +workitem/get_local_id.ll diff --git a/tgsi/lib/workitem/get_local_id.ll b/tgsi/lib/workitem/get_local_id.ll new file mode 100644 index 0000000..4f807b3 --- /dev/null +++ b/tgsi/lib/workitem/get_local_id.ll @@ -0,0 +1,18 @@ +declare i32 @llvm.tgsi.read.threadid.x() nounwind readnone +declare i32 @llvm.tgsi.read.threadid.y() nounwind readnone +declare i32 @llvm.tgsi.read.threadid.z() nounwind readnone + +define i32 @get_local_id(i32 %dim) nounwind readnone alwaysinline { + switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim] +x_dim: + %x = call i32 @llvm.tgsi.read.threadid.x() nounwind readnone + ret i32 %x +y_dim: + %y = call i32 @llvm.tgsi.read.threadid.y() nounwind readnone + ret i32 %y +z_dim: + %z = call i32 @llvm.tgsi.read.threadid.z() nounwind readnone + ret i32 %z +default: + ret i32 0 +} |