Skip to content

Instantly share code, notes, and snippets.

@Artem-B
Created September 11, 2023 23:42
Show Gist options
  • Star 0 You must be signed in to star a gist
  • Fork 0 You must be signed in to fork a gist
  • Save Artem-B/fe6092aa228d146ba4001bb95041b5f2 to your computer and use it in GitHub Desktop.
Save Artem-B/fe6092aa228d146ba4001bb95041b5f2 to your computer and use it in GitHub Desktop.
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