diff options
| author | Hans de Goede <hdegoede@redhat.com> | 2016-01-26 12:22:29 +0100 |
|---|---|---|
| committer | Hans de Goede <hdegoede@redhat.com> | 2016-06-28 11:53:40 +0200 |
| commit | 7cc2bc083fdd5d1268b2c0b17ece3b31acfac1d1 (patch) | |
| tree | e33adb822334172c6a757f7a2ab7839cf76a3d14 /include | |
| parent | 0f498eab65a9e3cfb4c0f26bd3b5f20bb243eacf (diff) | |
TGSI: initial Intrinsic support
Add Intrinsic support for get_local_id() and friends.
Signed-off-by: Hans de Goede <hdegoede@redhat.com>
Diffstat (limited to 'include')
| -rw-r--r-- | include/llvm/IR/Intrinsics.td | 1 | ||||
| -rw-r--r-- | include/llvm/IR/IntrinsicsTGSI.td | 35 |
2 files changed, 36 insertions, 0 deletions
diff --git a/include/llvm/IR/Intrinsics.td b/include/llvm/IR/Intrinsics.td index b50f07f3263..40686728f57 100644 --- a/include/llvm/IR/Intrinsics.td +++ b/include/llvm/IR/Intrinsics.td @@ -691,3 +691,4 @@ include "llvm/IR/IntrinsicsAMDGPU.td" include "llvm/IR/IntrinsicsBPF.td" include "llvm/IR/IntrinsicsSystemZ.td" include "llvm/IR/IntrinsicsWebAssembly.td" +include "llvm/IR/IntrinsicsTGSI.td" diff --git a/include/llvm/IR/IntrinsicsTGSI.td b/include/llvm/IR/IntrinsicsTGSI.td new file mode 100644 index 00000000000..1bb6c9b6309 --- /dev/null +++ b/include/llvm/IR/IntrinsicsTGSI.td @@ -0,0 +1,35 @@ +//===--- IntrinsicsTGSI.td - Defines TGSI intrinsics -----*- tablegen -*---===// +// +// The LLVM Compiler Infrastructure +// +// This file is distributed under the University of Illinois Open Source +// License. See LICENSE.TXT for details. +// +//===----------------------------------------------------------------------===// +// +// This file defines all of the TGSI-specific intrinsics. +// +//===----------------------------------------------------------------------===// + +let TargetPrefix = "tgsi" in { + +class TGSIReadPreloadRegisterIntrinsic<string name> + : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>, + GCCBuiltin<name>; + +multiclass TGSIReadPreloadRegisterIntrinsic_xyz<string prefix> { + def _x : TGSIReadPreloadRegisterIntrinsic<!strconcat(prefix, "_x")>; + def _y : TGSIReadPreloadRegisterIntrinsic<!strconcat(prefix, "_y")>; + def _z : TGSIReadPreloadRegisterIntrinsic<!strconcat(prefix, "_z")>; +} + +defm int_tgsi_read_blockid : TGSIReadPreloadRegisterIntrinsic_xyz < + "__builtin_tgsi_read_blockid">; +defm int_tgsi_read_blocksize : TGSIReadPreloadRegisterIntrinsic_xyz < + "__builtin_tgsi_read_blocksize">; +defm int_tgsi_read_gridsize : TGSIReadPreloadRegisterIntrinsic_xyz < + "__builtin_tgsi_read_gridsize">; +defm int_tgsi_read_threadid : TGSIReadPreloadRegisterIntrinsic_xyz < + "__builtin_tgsi_read_threadid">; + +} // End TargetPrefix = "tgsi" |