[TRITON] support.conv3d.triton.kernel

[kxyk99]

Motivation

Technical Details

Test Plan

Test Result

Submission Checklist

• Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

[github-actions]

⚠️ [ruff] <E731> _{reported by reviewdog 🐶}
Do not assign a lambda expression, use a def

Suggested change

	padding_same = lambda i, w, stride, dilation: math.ceil(
	(stride * (i - 1) + 1 + dilation * (w - 1) - i) / 2
	)
	def padding_same(i, w, stride, dilation):
	return math.ceil(
	(stride * (i - 1) + 1 + dilation * (w - 1) - i) / 2
	)

[github-actions]

⚠️ [ruff] <E731> _{reported by reviewdog 🐶}
Do not assign a lambda expression, use a def

Suggested change

	out_padding = lambda i, w, padding, stride, dilation: int(
	(i + 2 * padding - dilation * (w - 1) - 1) / stride + 1
	)
	def out_padding(i, w, padding, stride, dilation):
	return int(
	(i + 2 * padding - dilation * (w - 1) - 1) / stride + 1
	)

[github-actions]

⚠️ [ruff] <E731> _{reported by reviewdog 🐶}
Do not assign a lambda expression, use a def

Suggested change

	grid_fn = lambda META: (
	triton.cdiv(N * od * oh * ow, META['BLOCK_NI_DO_HO_WO']),
	triton.cdiv(oc // groups, META['BLOCK_CO']),
	groups,
	)
	def grid_fn(META):
	return (
	triton.cdiv(N * od * oh * ow, META['BLOCK_NI_DO_HO_WO']),
	triton.cdiv(oc // groups, META['BLOCK_CO']),
	groups,
	)

[github-actions]

⚠️ [ruff] <E731> _{reported by reviewdog 🐶}
Do not assign a lambda expression, use a def

Suggested change

	padding_same = lambda i, w, stride, dilation: math.ceil(
	(stride * (i - 1) + 1 + dilation * (w - 1) - i) / 2
	)
	def padding_same(i, w, stride, dilation):
	return math.ceil(
	(stride * (i - 1) + 1 + dilation * (w - 1) - i) / 2
	)

[github-actions]

⚠️ [ruff] <E731> _{reported by reviewdog 🐶}
Do not assign a lambda expression, use a def

Suggested change

	out_padding = lambda i, w, padding, stride, dilation: int(
	(i + 2 * padding - dilation * (w - 1) - 1) / stride + 1
	)
	def out_padding(i, w, padding, stride, dilation):
	return int(
	(i + 2 * padding - dilation * (w - 1) - 1) / stride + 1
	)

[github-actions]

⚠️ [ruff] <E731> _{reported by reviewdog 🐶}
Do not assign a lambda expression, use a def

Suggested change

	grid_fn = lambda META: (
	triton.cdiv(N * od * oh * ow, META['BLOCK_NI_DO_HO_WO']),
	triton.cdiv(oc // groups, META['BLOCK_CO']),
	groups,
	)
	def grid_fn(META):
	return (
	triton.cdiv(N * od * oh * ow, META['BLOCK_NI_DO_HO_WO']),
	triton.cdiv(oc // groups, META['BLOCK_CO']),
	groups,
	)

[vgokhale]

This kernel is the same as the one here

https://github.com/flagos-ai/FlagGems/blob/18bd83adfead48c30d784e961ddad84d514131e4/src/flag_gems/ops/conv3d.py#L114

which has a different license. Also it is not good for our AI libs to have kernels from other companies.

I would suggest closing this PR and writing a conv 3d kernel from scratch. Thanks.