Vague CompilationError

Working on a kernel I got a CompilationError when trying to launch it, and while the CompilationError message points to the "problem" it doesn't specify what the problem is. I eventually figured out that it doesn't like when you call load with other set but mask set to None. I don't mind that restriction, but it would be nice if the CompilationError message said something to that effect. I didn't see any specifics in the error message. Perhaps it's buried in the traceback, but the trackback gets truncated in Jupyter notebooks because it's so exceedingly long.

Triton version 2.0.0.dev20221202

The kernel:

@triton.jit
def _clip_embed_fused(
	A,         # *Pointer* to first input vector (257xN).
	Out,       # *Pointer* to output vector (257xN).
	Position,  # *Pointer* position vector (257xN).
	Weight,    # *Pointer* weight vector (N).
	Bias,      # *Pointer* bias vector (N).
	eps: tl.constexpr,
):
	# position of elements processed by this program
	row = tl.program_id(0)
	Out += row * 1024
	A += row * 1024
	Position += row * 1024

	# loads
	cols = tl.arange(0, 1024)
	mask = cols < 1024
	a = tl.load(A + cols, other=0.).to(tl.float32)
	weight = tl.load(Weight + cols)
	bias = tl.load(Bias + cols)
	position = tl.load(Position + cols, other=0.).to(tl.float32)

	# position embedding
	a += position

	# compute mean
	mean = tl.sum(a, axis=0) / 1024

	# compute variance
	var = tl.sum((a - mean) * (a - mean), axis=0) / 1024
	rstd = 1 / tl.sqrt(var + eps)

	# multiply by weight and add bias
	a_hat = (a - mean) * rstd
	out = a_hat * weight + bias

	# write-back
	tl.store(Out + cols, out)

The error:

{
	"name": "CompilationError",
	"message": "at 18:17:\ndef _clip_embed_fused(\n\tA,         # *Pointer* to first input vector (257xN).\n\tOut,       # *Pointer* to output vector (257xN).\n\tPosition,  # *Pointer* position vector (257xN).\n\tWeight,    # *Pointer* weight vector (N).\n\tBias,      # *Pointer* bias vector (N).\n\teps: tl.constexpr,\n):\n\t# position of elements processed by this program\n\trow = tl.program_id(0)\n\tOut += row * 1024\n\tA += row * 1024\n\tPosition += row * 1024\n\n\t# loads\n\tcols = tl.arange(0, 1024)\n\tmask = cols < 1024\n\ta = tl.load(A + cols, other=0.).to(tl.float32)\n                 ^",
	"stack": "\u001b[0;31m---------------------------------------------------------------------------\u001b[0m\n\u001b[0;31mKeyError\u001b[0m                                  Traceback (most recent call last)\nFile \u001b[0;32m<string>:21\u001b[0m, in \u001b[0;36m_clip_embed_fused\u001b[0;34m(A, Out, Position, Weight, Bias, eps, grid, num_warps, num_stages, extern_libs, stream, warmup)\u001b[0m\n\n\u001b[0;31mKeyError\u001b[0m: ('2-.-0-.-0-d82511111ad128294e9d31a6ac684238-2b0c5161c53c71b37ae20a9996ee4bb8-c1f92808b4e4644c1732e8338187ac87-d962222789c30252d492a16cca3bf467-12f7ac1ca211e037f62a7c0c323d9990-5c5e32ff210f3b7f56c98ca29917c25e-06f0df2d61979d629033f4a22eff5198-0dd03b0bd512a184b3512b278d9dfa59-d35ab04ae841e2714a253c523530b071', (torch.float32, torch.float32, torch.float32, torch.float32, torch.float32), (1e-05,), (True, True, True, True, True))\n\nDuring handling of the above exception, another exception occurred:\n\n\u001b[0;31mValueError\u001b[0m                                Traceback (most recent call last)\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/site-packages/triton/compiler.py:838\u001b[0m, in \u001b[0;36mmake_triton_ir\u001b[0;34m(fn, signature, specialization, constants)\u001b[0m\n\u001b[1;32m    837\u001b[0m \u001b[39mtry\u001b[39;00m:\n\u001b[0;32m--> 838\u001b[0m     generator\u001b[39m.\u001b[39;49mvisit(fn\u001b[39m.\u001b[39;49mparse())\n\u001b[1;32m    839\u001b[0m \u001b[39mexcept\u001b[39;00m \u001b[39mException\u001b[39;00m \u001b[39mas\u001b[39;00m e:\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/site-packages/triton/compiler.py:771\u001b[0m, in \u001b[0;36mCodeGenerator.visit\u001b[0;34m(self, node)\u001b[0m\n\u001b[1;32m    770\u001b[0m warnings\u001b[39m.\u001b[39msimplefilter(\u001b[39m\"\u001b[39m\u001b[39mignore\u001b[39m\u001b[39m\"\u001b[39m, \u001b[39mPendingDeprecationWarning\u001b[39;00m)  \u001b[39m# python 3.8\u001b[39;00m\n\u001b[0;32m--> 771\u001b[0m \u001b[39mreturn\u001b[39;00m \u001b[39msuper\u001b[39;49m()\u001b[39m.\u001b[39;49mvisit(node)\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/ast.py:418\u001b[0m, in \u001b[0;36mNodeVisitor.visit\u001b[0;34m(self, node)\u001b[0m\n\u001b[1;32m    417\u001b[0m visitor \u001b[39m=\u001b[39m \u001b[39mgetattr\u001b[39m(\u001b[39mself\u001b[39m, method, \u001b[39mself\u001b[39m\u001b[39m.\u001b[39mgeneric_visit)\n\u001b[0;32m--> 418\u001b[0m \u001b[39mreturn\u001b[39;00m visitor(node)\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/site-packages/triton/compiler.py:260\u001b[0m, in \u001b[0;36mCodeGenerator.visit_Module\u001b[0;34m(self, node)\u001b[0m\n\u001b[1;32m    259\u001b[0m \u001b[39mdef\u001b[39;00m \u001b[39mvisit_Module\u001b[39m(\u001b[39mself\u001b[39m, node):\n\u001b[0;32m--> 260\u001b[0m     ast\u001b[39m.\u001b[39;49mNodeVisitor\u001b[39m.\u001b[39;49mgeneric_visit(\u001b[39mself\u001b[39;49m, node)\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/ast.py:426\u001b[0m, in \u001b[0;36mNodeVisitor.generic_visit\u001b[0;34m(self, node)\u001b[0m\n\u001b[1;32m    425\u001b[0m         \u001b[39mif\u001b[39;00m \u001b[39misinstance\u001b[39m(item, AST):\n\u001b[0;32m--> 426\u001b[0m             \u001b[39mself\u001b[39;49m\u001b[39m.\u001b[39;49mvisit(item)\n\u001b[1;32m    427\u001b[0m \u001b[39melif\u001b[39;00m \u001b[39misinstance\u001b[39m(value, AST):\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/site-packages/triton/compiler.py:771\u001b[0m, in \u001b[0;36mCodeGenerator.visit\u001b[0;34m(self, node)\u001b[0m\n\u001b[1;32m    770\u001b[0m warnings\u001b[39m.\u001b[39msimplefilter(\u001b[39m\"\u001b[39m\u001b[39mignore\u001b[39m\u001b[39m\"\u001b[39m, \u001b[39mPendingDeprecationWarning\u001b[39;00m)  \u001b[39m# python 3.8\u001b[39;00m\n\u001b[0;32m--> 771\u001b[0m \u001b[39mreturn\u001b[39;00m \u001b[39msuper\u001b[39;49m()\u001b[39m.\u001b[39;49mvisit(node)\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/ast.py:418\u001b[0m, in \u001b[0;36mNodeVisitor.visit\u001b[0;34m(self, node)\u001b[0m\n\u001b[1;32m    417\u001b[0m visitor \u001b[39m=\u001b[39m \u001b[39mgetattr\u001b[39m(\u001b[39mself\u001b[39m, method, \u001b[39mself\u001b[39m\u001b[39m.\u001b[39mgeneric_visit)\n\u001b[0;32m--> 418\u001b[0m \u001b[39mreturn\u001b[39;00m visitor(node)\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/site-packages/triton/compiler.py:320\u001b[0m, in \u001b[0;36mCodeGenerator.visit_FunctionDef\u001b[0;34m(self, node)\u001b[0m\n\u001b[1;32m    319\u001b[0m \u001b[39m# visit function body\u001b[39;00m\n\u001b[0;32m--> 320\u001b[0m has_ret \u001b[39m=\u001b[39m \u001b[39mself\u001b[39;49m\u001b[39m.\u001b[39;49mvisit_compound_statement(node\u001b[39m.\u001b[39;49mbody)\n\u001b[1;32m    321\u001b[0m \u001b[39m# finalize\u001b[39;00m\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/site-packages/triton/compiler.py:254\u001b[0m, in \u001b[0;36mCodeGenerator.visit_compound_statement\u001b[0;34m(self, stmts)\u001b[0m\n\u001b[1;32m    253\u001b[0m \u001b[39mfor\u001b[39;00m stmt \u001b[39min\u001b[39;00m stmts:\n\u001b[0;32m--> 254\u001b[0m     \u001b[39mself\u001b[39m\u001b[39m.\u001b[39mlast_ret \u001b[39m=\u001b[39m \u001b[39mself\u001b[39;49m\u001b[39m.\u001b[39;49mvisit(stmt)\n\u001b[1;32m    255\u001b[0m     \u001b[39mif\u001b[39;00m \u001b[39misinstance\u001b[39m(stmt, ast\u001b[39m.\u001b[39mReturn):\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/site-packages/triton/compiler.py:771\u001b[0m, in \u001b[0;36mCodeGenerator.visit\u001b[0;34m(self, node)\u001b[0m\n\u001b[1;32m    770\u001b[0m warnings\u001b[39m.\u001b[39msimplefilter(\u001b[39m\"\u001b[39m\u001b[39mignore\u001b[39m\u001b[39m\"\u001b[39m, \u001b[39mPendingDeprecationWarning\u001b[39;00m)  \u001b[39m# python 3.8\u001b[39;00m\n\u001b[0;32m--> 771\u001b[0m \u001b[39mreturn\u001b[39;00m \u001b[39msuper\u001b[39;49m()\u001b[39m.\u001b[39;49mvisit(node)\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/ast.py:418\u001b[0m, in \u001b[0;36mNodeVisitor.visit\u001b[0;34m(self, node)\u001b[0m\n\u001b[1;32m    417\u001b[0m visitor \u001b[39m=\u001b[39m \u001b[39mgetattr\u001b[39m(\u001b[39mself\u001b[39m, method, \u001b[39mself\u001b[39m\u001b[39m.\u001b[39mgeneric_visit)\n\u001b[0;32m--> 418\u001b[0m \u001b[39mreturn\u001b[39;00m visitor(node)\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/site-packages/triton/compiler.py:367\u001b[0m, in \u001b[0;36mCodeGenerator.visit_Assign\u001b[0;34m(self, node)\u001b[0m\n\u001b[1;32m    366\u001b[0m names \u001b[39m=\u001b[39m _names[\u001b[39m0\u001b[39m]\n\u001b[0;32m--> 367\u001b[0m values \u001b[39m=\u001b[39m \u001b[39mself\u001b[39;49m\u001b[39m.\u001b[39;49mvisit(node\u001b[39m.\u001b[39;49mvalue)\n\u001b[1;32m    368\u001b[0m \u001b[39mif\u001b[39;00m \u001b[39mnot\u001b[39;00m \u001b[39misinstance\u001b[39m(names, \u001b[39mtuple\u001b[39m):\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/site-packages/triton/compiler.py:771\u001b[0m, in \u001b[0;36mCodeGenerator.visit\u001b[0;34m(self, node)\u001b[0m\n\u001b[1;32m    770\u001b[0m warnings\u001b[39m.\u001b[39msimplefilter(\u001b[39m\"\u001b[39m\u001b[39mignore\u001b[39m\u001b[39m\"\u001b[39m, \u001b[39mPendingDeprecationWarning\u001b[39;00m)  \u001b[39m# python 3.8\u001b[39;00m\n\u001b[0;32m--> 771\u001b[0m \u001b[39mreturn\u001b[39;00m \u001b[39msuper\u001b[39;49m()\u001b[39m.\u001b[39;49mvisit(node)\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/ast.py:418\u001b[0m, in \u001b[0;36mNodeVisitor.visit\u001b[0;34m(self, node)\u001b[0m\n\u001b[1;32m    417\u001b[0m visitor \u001b[39m=\u001b[39m \u001b[39mgetattr\u001b[39m(\u001b[39mself\u001b[39m, method, \u001b[39mself\u001b[39m\u001b[39m.\u001b[39mgeneric_visit)\n\u001b[0;32m--> 418\u001b[0m \u001b[39mreturn\u001b[39;00m visitor(node)\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/site-packages/triton/compiler.py:673\u001b[0m, in \u001b[0;36mCodeGenerator.visit_Call\u001b[0;34m(self, node)\u001b[0m\n\u001b[1;32m    672\u001b[0m \u001b[39mdef\u001b[39;00m \u001b[39mvisit_Call\u001b[39m(\u001b[39mself\u001b[39m, node):\n\u001b[0;32m--> 673\u001b[0m     fn \u001b[39m=\u001b[39m \u001b[39mself\u001b[39;49m\u001b[39m.\u001b[39;49mvisit(node\u001b[39m.\u001b[39;49mfunc)\n\u001b[1;32m    674\u001b[0m     \u001b[39mif\u001b[39;00m \u001b[39misinstance\u001b[39m(fn, triton\u001b[39m.\u001b[39mlanguage\u001b[39m.\u001b[39mconstexpr):\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/site-packages/triton/compiler.py:771\u001b[0m, in \u001b[0;36mCodeGenerator.visit\u001b[0;34m(self, node)\u001b[0m\n\u001b[1;32m    770\u001b[0m warnings\u001b[39m.\u001b[39msimplefilter(\u001b[39m\"\u001b[39m\u001b[39mignore\u001b[39m\u001b[39m\"\u001b[39m, \u001b[39mPendingDeprecationWarning\u001b[39;00m)  \u001b[39m# python 3.8\u001b[39;00m\n\u001b[0;32m--> 771\u001b[0m \u001b[39mreturn\u001b[39;00m \u001b[39msuper\u001b[39;49m()\u001b[39m.\u001b[39;49mvisit(node)\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/ast.py:418\u001b[0m, in \u001b[0;36mNodeVisitor.visit\u001b[0;34m(self, node)\u001b[0m\n\u001b[1;32m    417\u001b[0m visitor \u001b[39m=\u001b[39m \u001b[39mgetattr\u001b[39m(\u001b[39mself\u001b[39m, method, \u001b[39mself\u001b[39m\u001b[39m.\u001b[39mgeneric_visit)\n\u001b[0;32m--> 418\u001b[0m \u001b[39mreturn\u001b[39;00m visitor(node)\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/site-packages/triton/compiler.py:754\u001b[0m, in \u001b[0;36mCodeGenerator.visit_Attribute\u001b[0;34m(self, node)\u001b[0m\n\u001b[1;32m    753\u001b[0m \u001b[39mdef\u001b[39;00m \u001b[39mvisit_Attribute\u001b[39m(\u001b[39mself\u001b[39m, node):\n\u001b[0;32m--> 754\u001b[0m     lhs \u001b[39m=\u001b[39m \u001b[39mself\u001b[39;49m\u001b[39m.\u001b[39;49mvisit(node\u001b[39m.\u001b[39;49mvalue)\n\u001b[1;32m    755\u001b[0m     \u001b[39mreturn\u001b[39;00m \u001b[39mgetattr\u001b[39m(lhs, node\u001b[39m.\u001b[39mattr)\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/site-packages/triton/compiler.py:771\u001b[0m, in \u001b[0;36mCodeGenerator.visit\u001b[0;34m(self, node)\u001b[0m\n\u001b[1;32m    770\u001b[0m warnings\u001b[39m.\u001b[39msimplefilter(\u001b[39m\"\u001b[39m\u001b[39mignore\u001b[39m\u001b[39m\"\u001b[39m, \u001b[39mPendingDeprecationWarning\u001b[39;00m)  \u001b[39m# python 3.8\u001b[39;00m\n\u001b[0;32m--> 771\u001b[0m \u001b[39mreturn\u001b[39;00m \u001b[39msuper\u001b[39;49m()\u001b[39m.\u001b[39;49mvisit(node)\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/ast.py:418\u001b[0m, in \u001b[0;36mNodeVisitor.visit\u001b[0;34m(self, node)\u001b[0m\n\u001b[1;32m    417\u001b[0m visitor \u001b[39m=\u001b[39m \u001b[39mgetattr\u001b[39m(\u001b[39mself\u001b[39m, method, \u001b[39mself\u001b[39m\u001b[39m.\u001b[39mgeneric_visit)\n\u001b[0;32m--> 418\u001b[0m \u001b[39mreturn\u001b[39;00m visitor(node)\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/site-packages/triton/compiler.py:710\u001b[0m, in \u001b[0;36mCodeGenerator.visit_Call\u001b[0;34m(self, node)\u001b[0m\n\u001b[1;32m    709\u001b[0m \u001b[39mif\u001b[39;00m sys\u001b[39m.\u001b[39mmodules[fn\u001b[39m.\u001b[39m\u001b[39m__module__\u001b[39m] \u001b[39mis\u001b[39;00m triton\u001b[39m.\u001b[39mlanguage\u001b[39m.\u001b[39mcore \u001b[39mor\u001b[39;00m \u001b[39misinstance\u001b[39m(fn, triton\u001b[39m.\u001b[39mlanguage\u001b[39m.\u001b[39mextern\u001b[39m.\u001b[39mExternalFunction):\n\u001b[0;32m--> 710\u001b[0m     ret \u001b[39m=\u001b[39m fn(\u001b[39m*\u001b[39;49margs, _builder\u001b[39m=\u001b[39;49m\u001b[39mself\u001b[39;49m\u001b[39m.\u001b[39;49mbuilder, \u001b[39m*\u001b[39;49m\u001b[39m*\u001b[39;49mkws)\n\u001b[1;32m    711\u001b[0m \u001b[39mif\u001b[39;00m fn \u001b[39min\u001b[39;00m \u001b[39mself\u001b[39m\u001b[39m.\u001b[39mvalue_constructor\u001b[39m.\u001b[39mbuiltins\u001b[39m.\u001b[39mvalues():\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/site-packages/triton/language/core.py:46\u001b[0m, in \u001b[0;36mbuiltin.<locals>.wrapper\u001b[0;34m(*args, **kwargs)\u001b[0m\n\u001b[1;32m     45\u001b[0m     \u001b[39mraise\u001b[39;00m \u001b[39mValueError\u001b[39;00m(\u001b[39m\"\u001b[39m\u001b[39mDid you forget to add @triton.jit ? (`_builder` argument must be provided outside of JIT functions.)\u001b[39m\u001b[39m\"\u001b[39m)\n\u001b[0;32m---> 46\u001b[0m \u001b[39mreturn\u001b[39;00m fn(\u001b[39m*\u001b[39;49margs, \u001b[39m*\u001b[39;49m\u001b[39m*\u001b[39;49mkwargs)\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/site-packages/triton/language/core.py:813\u001b[0m, in \u001b[0;36mload\u001b[0;34m(pointer, mask, other, cache_modifier, eviction_policy, volatile, _builder)\u001b[0m\n\u001b[1;32m    812\u001b[0m volatile \u001b[39m=\u001b[39m _constexpr_to_value(volatile)\n\u001b[0;32m--> 813\u001b[0m \u001b[39mreturn\u001b[39;00m semantic\u001b[39m.\u001b[39;49mload(pointer, mask, other, cache_modifier, eviction_policy, volatile, _builder)\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/site-packages/triton/language/semantic.py:759\u001b[0m, in \u001b[0;36mload\u001b[0;34m(ptr, mask, other, cache_modifier, eviction_policy, is_volatile, builder)\u001b[0m\n\u001b[1;32m    758\u001b[0m \u001b[39mif\u001b[39;00m \u001b[39mnot\u001b[39;00m mask:\n\u001b[0;32m--> 759\u001b[0m     \u001b[39mraise\u001b[39;00m \u001b[39mValueError\u001b[39;00m(\u001b[39m\"\u001b[39m\u001b[39m`other` cannot be provided without `mask`\u001b[39m\u001b[39m\"\u001b[39m)\n\u001b[1;32m    761\u001b[0m \u001b[39mif\u001b[39;00m \u001b[39mnot\u001b[39;00m other:\n\n\u001b[0;31mValueError\u001b[0m: `other` cannot be provided without `mask`\n\nThe above exception was the direct cause of the following exception:\n\n\u001b[0;31mCompilationError\u001b[0m                          Traceback (most recent call last)\nCell \u001b[0;32mIn[78], line 11\u001b[0m\n\u001b[1;32m      8\u001b[0m clip \u001b[39m=\u001b[39m CLIPVisionEmbeddings()\u001b[39m.\u001b[39mto(\u001b[39m'\u001b[39m\u001b[39mcuda\u001b[39m\u001b[39m'\u001b[39m)\n\u001b[1;32m     10\u001b[0m \u001b[39m# forward pass\u001b[39;00m\n\u001b[0;32m---> 11\u001b[0m y_fus1 \u001b[39m=\u001b[39m fast_clip_vision_embeddings(x, clip\u001b[39m.\u001b[39;49mclass_embedding, clip\u001b[39m.\u001b[39;49mpatch_embedding\u001b[39m.\u001b[39;49mweight, clip\u001b[39m.\u001b[39;49mposition_embedding\u001b[39m.\u001b[39;49mweight, weight, bias, eps\u001b[39m=\u001b[39;49m\u001b[39m1e-5\u001b[39;49m)\n\u001b[1;32m     12\u001b[0m y_ref2 \u001b[39m=\u001b[39m clip_vision_embeddings2(x, clip\u001b[39m.\u001b[39mclass_embedding, clip\u001b[39m.\u001b[39mpatch_embedding\u001b[39m.\u001b[39mweight, clip\u001b[39m.\u001b[39mposition_embedding\u001b[39m.\u001b[39mweight, weight, bias, eps\u001b[39m=\u001b[39m\u001b[39m1e-5\u001b[39m)\n\u001b[1;32m     13\u001b[0m y_ref \u001b[39m=\u001b[39m clip_vision_embeddings(x, clip\u001b[39m.\u001b[39mclass_embedding, clip\u001b[39m.\u001b[39mpatch_embedding\u001b[39m.\u001b[39mweight, clip\u001b[39m.\u001b[39mposition_embedding\u001b[39m.\u001b[39mweight, weight, bias, eps\u001b[39m=\u001b[39m\u001b[39m1e-5\u001b[39m)\n\nCell \u001b[0;32mIn[77], line 144\u001b[0m, in \u001b[0;36mfast_clip_vision_embeddings\u001b[0;34m(pixel_values, class_embedding, patch_embedding, position_embedding, ln_weight, ln_bias, eps)\u001b[0m\n\u001b[1;32m    142\u001b[0m embeddings \u001b[39m=\u001b[39m torch\u001b[39m.\u001b[39mcat([class_embeds, patch_embeds], dim\u001b[39m=\u001b[39m\u001b[39m1\u001b[39m)  \u001b[39m# shape = [*, 257, 1024]\u001b[39;00m\n\u001b[1;32m    143\u001b[0m \u001b[39massert\u001b[39;00m embeddings\u001b[39m.\u001b[39mis_contiguous()\n\u001b[0;32m--> 144\u001b[0m embeddings \u001b[39m=\u001b[39m clip_positions_and_layernorm(embeddings, position_embedding, ln_weight, ln_bias, eps)\n\u001b[1;32m    146\u001b[0m \u001b[39mreturn\u001b[39;00m embeddings\n\nCell \u001b[0;32mIn[77], line 125\u001b[0m, in \u001b[0;36mclip_positions_and_layernorm\u001b[0;34m(x, positions, weight, bias, eps)\u001b[0m\n\u001b[1;32m    123\u001b[0m num_warps \u001b[39m=\u001b[39m \u001b[39mmin\u001b[39m(\u001b[39mmax\u001b[39m(\u001b[39m1024\u001b[39m \u001b[39m/\u001b[39m\u001b[39m/\u001b[39m \u001b[39m256\u001b[39m, \u001b[39m1\u001b[39m), \u001b[39m8\u001b[39m)\n\u001b[1;32m    124\u001b[0m \u001b[39m# enqueue kernel\u001b[39;00m\n\u001b[0;32m--> 125\u001b[0m _clip_embed_fused[(M,)](x_arg, y, positions, weight, bias,\n\u001b[1;32m    126\u001b[0m \t\t\t\t\t\t\teps, num_warps\u001b[39m=\u001b[39;49mnum_warps)\n\u001b[1;32m    128\u001b[0m \u001b[39mreturn\u001b[39;00m y\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/site-packages/triton/runtime/jit.py:106\u001b[0m, in \u001b[0;36mKernelInterface.__getitem__.<locals>.launcher\u001b[0;34m(*args, **kwargs)\u001b[0m\n\u001b[1;32m    105\u001b[0m \u001b[39mdef\u001b[39;00m \u001b[39mlauncher\u001b[39m(\u001b[39m*\u001b[39margs, \u001b[39m*\u001b[39m\u001b[39m*\u001b[39mkwargs):\n\u001b[0;32m--> 106\u001b[0m     \u001b[39mreturn\u001b[39;00m \u001b[39mself\u001b[39;49m\u001b[39m.\u001b[39;49mrun(\u001b[39m*\u001b[39;49margs, grid\u001b[39m=\u001b[39;49mgrid, \u001b[39m*\u001b[39;49m\u001b[39m*\u001b[39;49mkwargs)\n\nFile \u001b[0;32m<string>:41\u001b[0m, in \u001b[0;36m_clip_embed_fused\u001b[0;34m(A, Out, Position, Weight, Bias, eps, grid, num_warps, num_stages, extern_libs, stream, warmup)\u001b[0m\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/site-packages/triton/compiler.py:1256\u001b[0m, in \u001b[0;36mcompile\u001b[0;34m(fn, signature, device, constants, num_warps, num_stages, extern_libs, configs, cc, warm_cache_only)\u001b[0m\n\u001b[1;32m   1250\u001b[0m llir_name \u001b[39m=\u001b[39m \u001b[39mf\u001b[39m\u001b[39m\"\u001b[39m\u001b[39m{\u001b[39;00mname\u001b[39m}\u001b[39;00m\u001b[39m.llir\u001b[39m\u001b[39m\"\u001b[39m\n\u001b[1;32m   1251\u001b[0m \u001b[39mif\u001b[39;00m \u001b[39mnot\u001b[39;00m fn_cache_manager\u001b[39m.\u001b[39mhas_file(cubin_name) \u001b[39mor\u001b[39;00m \\\n\u001b[1;32m   1252\u001b[0m    \u001b[39mnot\u001b[39;00m fn_cache_manager\u001b[39m.\u001b[39mhas_file(data_name) \u001b[39mor\u001b[39;00m \\\n\u001b[1;32m   1253\u001b[0m    \u001b[39mnot\u001b[39;00m fn_cache_manager\u001b[39m.\u001b[39mhas_file(ptx_name) \u001b[39mor\u001b[39;00m \\\n\u001b[1;32m   1254\u001b[0m    \u001b[39mnot\u001b[39;00m fn_cache_manager\u001b[39m.\u001b[39mhas_file(ttir_name) \u001b[39mor\u001b[39;00m \\\n\u001b[1;32m   1255\u001b[0m    \u001b[39mnot\u001b[39;00m fn_cache_manager\u001b[39m.\u001b[39mhas_file(llir_name):\n\u001b[0;32m-> 1256\u001b[0m     asm, shared, kernel_name \u001b[39m=\u001b[39m _compile(fn, signature, device, constants, configs[\u001b[39m0\u001b[39;49m], num_warps, num_stages,\n\u001b[1;32m   1257\u001b[0m                                         extern_libs, \u001b[39m\"\u001b[39;49m\u001b[39mcubin\u001b[39;49m\u001b[39m\"\u001b[39;49m, cc)\n\u001b[1;32m   1258\u001b[0m     metadata \u001b[39m=\u001b[39m {\u001b[39m\"\u001b[39m\u001b[39mname\u001b[39m\u001b[39m\"\u001b[39m: kernel_name, \u001b[39m\"\u001b[39m\u001b[39mshared\u001b[39m\u001b[39m\"\u001b[39m: shared, \u001b[39m\"\u001b[39m\u001b[39mnum_warps\u001b[39m\u001b[39m\"\u001b[39m: num_warps, \u001b[39m\"\u001b[39m\u001b[39mnum_stages\u001b[39m\u001b[39m\"\u001b[39m: num_stages}\n\u001b[1;32m   1259\u001b[0m     fn_cache_manager\u001b[39m.\u001b[39mput(asm[\u001b[39m\"\u001b[39m\u001b[39mcubin\u001b[39m\u001b[39m\"\u001b[39m], cubin_name)\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/site-packages/triton/compiler.py:892\u001b[0m, in \u001b[0;36m_compile\u001b[0;34m(fn, signature, device, constants, specialization, num_warps, num_stages, extern_libs, output, cc)\u001b[0m\n\u001b[1;32m    889\u001b[0m \u001b[39massert\u001b[39;00m output \u001b[39min\u001b[39;00m valid_outputs, \u001b[39m\"\u001b[39m\u001b[39moutput should be one of [\u001b[39m\u001b[39m%s\u001b[39;00m\u001b[39m], but get \u001b[39m\u001b[39m\\\"\u001b[39;00m\u001b[39m%s\u001b[39;00m\u001b[39m\\\"\u001b[39;00m\u001b[39m\"\u001b[39m \u001b[39m%\u001b[39m (\u001b[39m'\u001b[39m\u001b[39m,\u001b[39m\u001b[39m'\u001b[39m\u001b[39m.\u001b[39mjoin(valid_outputs), output)\n\u001b[1;32m    891\u001b[0m \u001b[39m# triton-ir\u001b[39;00m\n\u001b[0;32m--> 892\u001b[0m module, _ \u001b[39m=\u001b[39m make_triton_ir(fn, signature, specialization, constants)\n\u001b[1;32m    893\u001b[0m \u001b[39mif\u001b[39;00m output \u001b[39m==\u001b[39m \u001b[39m\"\u001b[39m\u001b[39mttir\u001b[39m\u001b[39m\"\u001b[39m:\n\u001b[1;32m    894\u001b[0m     \u001b[39mreturn\u001b[39;00m module\n\nFile \u001b[0;32m~/anaconda3/envs/tmpenv2/lib/python3.10/site-packages/triton/compiler.py:843\u001b[0m, in \u001b[0;36mmake_triton_ir\u001b[0;34m(fn, signature, specialization, constants)\u001b[0m\n\u001b[1;32m    841\u001b[0m     \u001b[39mif\u001b[39;00m node \u001b[39mis\u001b[39;00m \u001b[39mNone\u001b[39;00m \u001b[39mor\u001b[39;00m \u001b[39misinstance\u001b[39m(e, (\u001b[39mNotImplementedError\u001b[39;00m, CompilationError)):\n\u001b[1;32m    842\u001b[0m         \u001b[39mraise\u001b[39;00m e\n\u001b[0;32m--> 843\u001b[0m     \u001b[39mraise\u001b[39;00m CompilationError(fn\u001b[39m.\u001b[39msrc, node) \u001b[39mfrom\u001b[39;00m \u001b[39me\u001b[39;00m\n\u001b[1;32m    844\u001b[0m ret \u001b[39m=\u001b[39m generator\u001b[39m.\u001b[39mmodule\n\u001b[1;32m    845\u001b[0m \u001b[39m# module takes ownership of the context\u001b[39;00m\n\n\u001b[0;31mCompilationError\u001b[0m: at 18:17:\ndef _clip_embed_fused(\n\tA,         # *Pointer* to first input vector (257xN).\n\tOut,       # *Pointer* to output vector (257xN).\n\tPosition,  # *Pointer* position vector (257xN).\n\tWeight,    # *Pointer* weight vector (N).\n\tBias,      # *Pointer* bias vector (N).\n\teps: tl.constexpr,\n):\n\t# position of elements processed by this program\n\trow = tl.program_id(0)\n\tOut += row * 1024\n\tA += row * 1024\n\tPosition += row * 1024\n\n\t# loads\n\tcols = tl.arange(0, 1024)\n\tmask = cols < 1024\n\ta = tl.load(A + cols, other=0.).to(tl.float32)\n                 ^"
}