[BugFix] fix nvptx not supported by device_enabled error

[ZQPei]

Fix Issue #9513.

I think this bug was mistakenly introduced by #8032.

[junrushao]

LGTM. Thanks!

[ZQPei]

@junrushao1994
I think thir pr needs a further review, because another error occurs in ci https://ci.tlcpack.ai/blue/organizations/jenkins/tvm/detail/PR-9585/1/pipeline/267.

Quoted error is

E File "/workspace/src/target/llvm/codegen_nvptx.cc", line 146
E TVMError: Do not support sync shared.dyn

This error is raised because of shared.dyn not supported by "nvptx" target.
And it was previously covered by device_enabled("nvptx") returning False in the last few months.

tvm/tests/python/unittest/test_tir_ir_builder.py

Lines 499 to 552 in 0a4cc89

	@tvm.testing.requires_gpu
	def test_dyn_shared():
	n = te.size_var("n")
	dtype = "float32"
	A = te.placeholder((n,), name="A")
	def test_device_ir(A, B):
	n = A.shape[0]
	ib = tvm.tir.ir_builder.create()
	tx = te.thread_axis("threadIdx.x")
	ib.scope_attr(tx, "thread_extent", n)
	temp = ib.allocate(dtype, (n,), scope="shared.dyn") # n is symbolic size
	Aptr = ib.buffer_ptr(A)
	Bptr = ib.buffer_ptr(B)
	temp[tx] = Aptr[tx]
	depth = tvm.tir.log2(cast(n, "float32"))
	with ib.for_range(0, depth) as i:
	ib.emit(tvm.tir.Call(None, "tir.tvm_storage_sync", tvm.runtime.convert(["shared"])))
	d = n >> (i + 1)
	with ib.if_scope(tx < d):
	temp[tx] += temp[tx + d]
	Bptr[0] = temp[0]
	return ib.get()
	B = te.extern(
	(1,),
	[A],
	lambda ins, outs: test_device_ir(ins[0], outs[0]),
	name="reduce",
	dtype=dtype,
	)
	s = te.create_schedule(B.op)
	def check_target(target):
	if not tvm.testing.device_enabled(target):
	return
	freduce = tvm.build(s, [A, B], target)
	dev = tvm.device(target, 0)
	for n in [512, 1024]:
	a = tvm.nd.array(np.random.uniform(size=n).astype(A.dtype), dev)
	b = tvm.nd.array(np.zeros(1, dtype=B.dtype), dev)
	freduce(a, b)
	tvm.testing.assert_allclose(b.numpy()[0], np.sum(a.numpy()), 1e-4, 1e-4)
	for target in ["cuda", "nvptx"]:
	check_target(target)

I have tried to fix it with this commit a8d14b1, and it works fine on my 2080ti gpu card.
But I am not sure if it is really OK.

[vinx13]

The codegen fix looks good

[junrushao]

Yeah I think it's reasonable to handle dynamic shared memory this way for now

[junrushao]

Thank you so much for the fix! @ZQPei