Created
September 11, 2023 23:42
-
-
Save Artem-B/fe6092aa228d146ba4001bb95041b5f2 to your computer and use it in GitHub Desktop.
This file contains 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