Created
September 11, 2023 23:42
-
-
Save Artem-B/fe6092aa228d146ba4001bb95041b5f2 to your computer and use it in GitHub Desktop.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| from itertools import product | |
| from string import Template | |
| from itertools import product, count | |
| from string import Template | |
| from absl import app | |
| from typing import Sequence | |
| types = [ | |
| "char", "signed char", "char1", "char2", "char4", "unsigned char", "uchar1", | |
| "uchar2", "uchar4", "short", "short1", "short2", "short4", "ushort", | |
| "ushort1", "ushort2", "ushort4", "int", "int1", "int2", "int4", "uint", | |
| "uint1", "uint2", "uint4", "float", "float1", "float2", "float4" | |
| ] | |
| # pyformat: disable | |
| ops = [ | |
| ["tex1D", "T *", "cudaTextureObject_t", "float"], | |
| ["tex1D", "cudaTextureObject_t", "float"], | |
| ["tex1D", "texture<T, cudaTextureType1D, cudaReadModeElementType>", "float"], | |
| ["tex1D", "texture<T, cudaTextureType1D, cudaReadModeNormalizedFloat>", "float"], | |
| ["tex1DGrad", "T *", "cudaTextureObject_t", "float", "float", "float"], | |
| ["tex1DGrad", "cudaTextureObject_t", "float", "float", "float"], | |
| ["tex1DGrad", "texture<T, cudaTextureType1D, cudaReadModeElementType>", "float", "float", "float"], | |
| ["tex1DGrad", "texture<T, cudaTextureType1D, cudaReadModeNormalizedFloat>", "float", "float", "float"], | |
| ["tex1DLayered", "T *", "cudaTextureObject_t", "float", "int"], | |
| ["tex1DLayered", "cudaTextureObject_t", "float", "int"], | |
| ["tex1DLayered", "texture<T, cudaTextureType1DLayered, cudaReadModeElementType>", "float", "int"], | |
| ["tex1DLayered", "texture<T, cudaTextureType1DLayered, cudaReadModeNormalizedFloat>", "float", "int"], | |
| ["tex1DLayeredGrad", "T *", "cudaTextureObject_t", "float", "int", "float", "float"], | |
| ["tex1DLayeredGrad", "cudaTextureObject_t", "float", "int", "float", "float"], | |
| ["tex1DLayeredGrad", "texture<T, cudaTextureType1DLayered, cudaReadModeElementType>", "float", "int", "float", "float"], | |
| ["tex1DLayeredGrad", "texture<T, cudaTextureType1DLayered, cudaReadModeNormalizedFloat>", "float", "int", "float", "float"], | |
| ["tex1DLayeredLod", "T *", "cudaTextureObject_t", "float", "int", "float"], | |
| ["tex1DLayeredLod", "cudaTextureObject_t", "float", "int", "float"], | |
| ["tex1DLayeredLod", "texture<T, cudaTextureType1DLayered, cudaReadModeElementType>", "float", "int", "float"], | |
| ["tex1DLayeredLod", "texture<T, cudaTextureType1DLayered, cudaReadModeNormalizedFloat>", "float", "int", "float"], | |
| ["tex1DLod", "T *", "cudaTextureObject_t", "float", "float"], | |
| ["tex1DLod", "cudaTextureObject_t", "float", "float"], | |
| ["tex1DLod", "texture<T, cudaTextureType1D, cudaReadModeElementType>", "float", "float"], | |
| ["tex1DLod", "texture<T, cudaTextureType1D, cudaReadModeNormalizedFloat>", "float", "float"], | |
| ["tex1Dfetch", "T *", "cudaTextureObject_t", "int"], | |
| ["tex1Dfetch", "cudaTextureObject_t", "int"], | |
| ["tex1Dfetch", "texture<T, cudaTextureType1D, cudaReadModeElementType>", "int"], | |
| ["tex1Dfetch", "texture<T, cudaTextureType1D, cudaReadModeNormalizedFloat>", "int"], | |
| ["tex2D", "T *", "cudaTextureObject_t", "float", "float"], | |
| ["tex2D", "T *", "cudaTextureObject_t", "float", "float"], | |
| ["tex2D", "cudaTextureObject_t", "float", "float"], | |
| ["tex2D", "cudaTextureObject_t", "float", "float", "bool*"], | |
| ["tex2D", "texture<T, cudaTextureType2D, cudaReadModeElementType>", "float", "float"], | |
| ["tex2D", "texture<T, cudaTextureType2D, cudaReadModeNormalizedFloat>", "float", "float"], | |
| ["tex2DGrad", "T *", "cudaTextureObject_t", "float", "float", "float2", "float2"], | |
| ["tex2DGrad", "T *", "cudaTextureObject_t", "float", "float", "float2", "float2", "bool*"], | |
| ["tex2DGrad", "cudaTextureObject_t", "float", "float", "float2", "float2"], | |
| ["tex2DGrad", "cudaTextureObject_t", "float", "float", "float2", "float2", "bool*"], | |
| ["tex2DGrad", "texture<T, cudaTextureType2D, cudaReadModeElementType>", "float", "float", "float2", "float2"], | |
| ["tex2DGrad", "texture<T, cudaTextureType2D, cudaReadModeNormalizedFloat>", "float", "float", "float2", "float2"], | |
| ["tex2DLayered", "T *", "cudaTextureObject_t", "float", "float", "int"], | |
| ["tex2DLayered", "T *", "cudaTextureObject_t", "float", "float", "int", "bool*"], | |
| ["tex2DLayered", "cudaTextureObject_t", "float", "float", "int"], | |
| ["tex2DLayered", "cudaTextureObject_t", "float", "float", "int", "bool*"], | |
| ["tex2DLayered", "texture<T, cudaTextureType2DLayered, cudaReadModeElementType>", "float", "float", "int"], | |
| ["tex2DLayered", "texture<T, cudaTextureType2DLayered, cudaReadModeNormalizedFloat>", "float", "float", "int"], | |
| ["tex2DLayeredGrad", "T * ", "cudaTextureObject_t", "float", "float", "int", "float2", "float2"], | |
| ["tex2DLayeredGrad", "T * ", "cudaTextureObject_t", "float", "float", "int", "float2", "float2", "bool*"], | |
| ["tex2DLayeredGrad", "cudaTextureObject_t", "float", "float", "int", "float2", "float2"], | |
| ["tex2DLayeredGrad", "cudaTextureObject_t", "float", "float", "int", "float2", "float2", "bool*"], | |
| ["tex2DLayeredGrad", "texture<T, cudaTextureType2DLayered, cudaReadModeElementType>", "float", "float", "int", "float2", "float2"], | |
| ["tex2DLayeredGrad", "texture<T, cudaTextureType2DLayered, cudaReadModeNormalizedFloat>", "float", "float", "int", "float2", "float2"], | |
| ["tex2DLayeredLod", "T *", "cudaTextureObject_t", "float", "float", "int", "float"], | |
| ["tex2DLayeredLod", "T *", "cudaTextureObject_t", "float", "float", "int", "float", "bool*"], | |
| ["tex2DLayeredLod", "cudaTextureObject_t", "float", "float", "int", "float"], | |
| ["tex2DLayeredLod", "cudaTextureObject_t", "float", "float", "int", "float", "bool*"], | |
| ["tex2DLayeredLod", "texture<T, cudaTextureType2DLayered, cudaReadModeElementType>", "float", "float", "int", "float"], | |
| ["tex2DLayeredLod", "texture<T, cudaTextureType2DLayered, cudaReadModeNormalizedFloat>", "float", "float", "int", "float"], | |
| ["tex2DLod", "T *", "cudaTextureObject_t", "float", "float", "float"], | |
| ["tex2DLod", "T *", "cudaTextureObject_t", "float", "float", "float", "bool*"], | |
| ["tex2DLod", "cudaTextureObject_t", "float", "float", "float"], | |
| ["tex2DLod", "cudaTextureObject_t", "float", "float", "float", "bool*"], | |
| ["tex2DLod", "texture<T, cudaTextureType2D, cudaReadModeElementType>", "float", "float", "float"], | |
| ["tex2DLod", "texture<T, cudaTextureType2D, cudaReadModeNormalizedFloat>", "float", "float", "float"], | |
| ["tex2Dgather", "T *", "cudaTextureObject_t", "float", "float", "bool*, int"], | |
| ["tex2Dgather", "T *", "cudaTextureObject_t", "float", "float", "int"], | |
| ["tex2Dgather", "cudaTextureObject_t to, float", "float", "bool*, int"], | |
| ["tex2Dgather", "cudaTextureObject_t to, float", "float", "int"], | |
| ["tex2Dgather", "texture<T, cudaTextureType2D, cudaReadModeElementType>", "float", "float", "int"], | |
| ["tex2Dgather", "texture<T, cudaTextureType2D, cudaReadModeNormalizedFloat>", "float", "float", "int"], | |
| ["tex3D", "T *", "cudaTextureObject_t", "float", "float", "float"], | |
| ["tex3D", "T *", "cudaTextureObject_t", "float", "float", "float"], | |
| ["tex3D", "cudaTextureObject_t", "float", "float", "float"], | |
| ["tex3D", "cudaTextureObject_t", "float", "float", "float", "bool*"], | |
| ["tex3D", "texture<T, cudaTextureType3D, cudaReadModeElementType>", "float", "float", "float"], | |
| ["tex3D", "texture<T, cudaTextureType3D, cudaReadModeNormalizedFloat>", "float", "float", "float"], | |
| ["tex3DGrad", "T *", "cudaTextureObject_t", "float", "float", "float", "float4", "float4"], | |
| ["tex3DGrad", "T *", "cudaTextureObject_t", "float", "float", "float", "float4", "float4", "bool*"], | |
| ["tex3DGrad", "cudaTextureObject_t", "float", "float", "float", "float4", "float4"], | |
| ["tex3DGrad", "cudaTextureObject_t", "float", "float", "float", "float4", "float4", "bool*"], | |
| ["tex3DGrad", "texture<T, cudaTextureType3D, cudaReadModeElementType>", "float", "float", "float", "float4", "float4"], | |
| ["tex3DGrad", "texture<T, cudaTextureType3D, cudaReadModeNormalizedFloat>", "float", "float", "float", "float4", "float4"], | |
| ["tex3DLod", "T *", "cudaTextureObject_t", "float", "float", "float", "float"], | |
| ["tex3DLod", "T *", "cudaTextureObject_t", "float", "float", "float", "float", "bool*"], | |
| ["tex3DLod", "cudaTextureObject_t", "float", "float", "float", "float"], | |
| ["tex3DLod", "cudaTextureObject_t", "float", "float", "float", "float", "bool*"], | |
| ["tex3DLod", "texture<T, cudaTextureType3D, cudaReadModeElementType>", "float", "float", "float", "float"], | |
| ["tex3DLod", "texture<T, cudaTextureType3D, cudaReadModeNormalizedFloat>", "float", "float", "float", "float"], | |
| ["texCubemap", "T *", "cudaTextureObject_t", "float", "float", "float"], | |
| ["texCubemap", "cudaTextureObject_t", "float", "float", "float"], | |
| ["texCubemap", "texture<T, cudaTextureTypeCubemap, cudaReadModeElementType>", "float", "float", "float"], | |
| ["texCubemap", "texture<T, cudaTextureTypeCubemap, cudaReadModeNormalizedFloat>", "float", "float", "float"], | |
| ["texCubemapGrad", "T *", "cudaTextureObject_t", "float", "float", "float", "float4", "float4"], | |
| ["texCubemapGrad", "cudaTextureObject_t", "float", "float", "float", "float4", "float4"], | |
| ["texCubemapGrad", "texture<T, cudaTextureTypeCubemap, cudaReadModeElementType>", "float", "float", "float", "float4", "float4"], | |
| ["texCubemapGrad", "texture<T, cudaTextureTypeCubemap, cudaReadModeNormalizedFloat>", "float", "float", "float", "float4", "float4"], | |
| ["texCubemapLayered", "T *", "cudaTextureObject_t", "float", "float", "float", "int"], | |
| ["texCubemapLayered", "cudaTextureObject_t", "float", "float", "float", "int"], | |
| ["texCubemapLayered", "texture<T, cudaTextureTypeCubemapLayered, cudaReadModeElementType>", "float", "float", "float", "int"], | |
| ["texCubemapLayered", "texture<T, cudaTextureTypeCubemapLayered, cudaReadModeNormalizedFloat>", "float", "float", "float", "int"], | |
| ["texCubemapLayeredGrad", "T *", "cudaTextureObject_t", "float", "float", "float", "int", "float4", "float4"], | |
| ["texCubemapLayeredGrad", "cudaTextureObject_t", "float", "float", "float", "int", "float4", "float4"], | |
| ["texCubemapLayeredGrad", "texture<T, cudaTextureTypeCubemapLayered, cudaReadModeElementType>", "float", "float", "float", "int", "float4", "float4"], | |
| ["texCubemapLayeredGrad", "texture<T, cudaTextureTypeCubemapLayered, cudaReadModeNormalizedFloat>", "float", "float", "float", "int", "float4", "float4"], | |
| ["texCubemapLayeredLod", "T *", "cudaTextureObject_t", "float", "float", "float", "int", "float"], | |
| ["texCubemapLayeredLod", "cudaTextureObject_t", "float", "float", "float", "int", "float"], | |
| ["texCubemapLayeredLod", "texture<T, cudaTextureTypeCubemapLayered, cudaReadModeElementType>", "float", "float", "float", "int", "float"], | |
| ["texCubemapLayeredLod", "texture<T, cudaTextureTypeCubemapLayered, cudaReadModeNormalizedFloat>", "float", "float", "float", "int", "float"], | |
| ["texCubemapLod", "T *", "cudaTextureObject_t", "float", "float", "float", "float"], | |
| ["texCubemapLod", "cudaTextureObject_t", "float", "float", "float", "float"], | |
| ["texCubemapLod", "texture<T, cudaTextureTypeCubemap, cudaReadModeElementType>", "float", "float", "float", "float"], | |
| ["texCubemapLod", "texture<T, cudaTextureTypeCubemap, cudaReadModeNormalizedFloat>", "float", "float", "float", "float"], | |
| ] | |
| # pyformat: enable | |
| global_counter = count() | |
| def gen_id(prefix): | |
| return "%s_%d" % (prefix, next(global_counter)) | |
| def gen_test(op, t): | |
| #print(op,t) | |
| types = op[1:] | |
| if types[0].startswith("T") or types[0].startswith("texture<T"): | |
| types[0] = types[0].replace("T", t, 1) | |
| op_name = op[0] | |
| callargs = ["v%d" % v for v in range(len(types))] | |
| funcargs = ["%s %s" % (t, v) for t, v in zip(types, callargs)] | |
| result = """ | |
| // {params} | |
| __device__ auto {name}({signature}){{ | |
| return {call}; | |
| }} | |
| """.format( | |
| name=gen_id("test_%s_%s" % (op_name, t.replace(" ", "_"))), | |
| signature=", ".join(funcargs), | |
| call="%s<%s>(%s)" % (op_name, t, ", ".join(callargs)), | |
| params=[t, op]) | |
| return result | |
| def gen_op_direct(op, t): | |
| return gen_test(op, t) | |
| def gen_op(op, t): | |
| return gen_op_direct(op, t) | |
| def gen_all(): | |
| print(""" | |
| #include <texture_fetch_functions.h> | |
| #include <texture_indirect_functions.h> | |
| """) | |
| for op, t in product(ops, types): | |
| # Normalized float is available for char/short only. | |
| if op[1].endswith("cudaReadModeNormalizedFloat>") and t in [ | |
| "int", "int1", "int2", "int4", "uint", "uint1", "uint2", "uint4", | |
| "float", "float1", "float2", "float4" | |
| ]: | |
| continue | |
| print(gen_op(op, t)) | |
| def main(argv: Sequence[str]) -> None: | |
| if len(argv) > 1: | |
| raise app.UsageError("Too many command-line arguments.") | |
| gen_all() | |
| if __name__ == "__main__": | |
| app.run(main) |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment