[Relay][TE] Use Relay parameter name to generated TE tensor name by Lunderberg · Pull Request #10516 · apache/tvm

Lunderberg · 2022-03-07T15:41:53Z

Previously, the TE placeholders representing relay function parameters were all named "placeholder", which could be difficult to follow when debugging larger functions.

Lunderberg · 2022-03-09T15:13:25Z

Rebase onto main following #10535

Lunderberg · 2022-03-14T14:26:36Z

Rebase onto main following merge conflict with #10577.

Lunderberg · 2022-03-14T20:34:37Z

Current CI failures due to checks for the name "placeholder" in auto_scheduler (link). Will investigate further as time allows, as I'm not sure the history of this check.

tkonolige

This is a very helpful change. Thanks @Lunderberg!

areusch · 2022-04-11T23:31:07Z

@Lunderberg what's the status on this one?

Lunderberg · 2022-04-12T14:35:31Z

Currently there's one CI failure that I haven't had time to track down yet. The error message occurs during memory verification, that a buffer is undefined, but likely is caused by an earlier pass failing to update a buffer.

Lunderberg · 2022-04-19T14:21:50Z

Rebased onto main, since the failing test case was disabled in #10717.

Lunderberg · 2022-09-13T15:50:24Z

I let this one languish long enough that I no longer trust the CI checks to be up-to-date. Re-running CI.

Previously, the TE placeholders representing relay function parameters were all named `"placeholder"`, which could be difficult to follow when debugging larger functions.

The tensor name "ethos-u", once passed through from relay to TE, resulted in invalid C++ codegen.

The tensor name is part of the hash used for these results, so the previous hashes are no longer valid.

#10516 used the Relay parameter name when lowering to TE. However, this creates an issue when the parameter name is empty. This is legal in Relay, but results in errors during code generation. For example, this is the generated CUDA kernel for bias add: ``` extern "C" __global__ void __launch_bounds__(1024) fused_raf_op_tvm_add_kernel0( float* __restrict__ T_add, float* __restrict__ , /* Name is missing and it results in compile errors. */ float* __restrict__ _1) { T_add[((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] = ([((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] + _1[((((((int)blockIdx.x) * 16) + (((int)threadIdx.x) >> 6)) % 54) / 9)]); } ``` This PR adds "placeholder" back as a default to make sure no empty string will be passed when lowering to TE.

…che#10516) * [Relay][TE] Use Relay parameter name to generated TE tensor name Previously, the TE placeholders representing relay function parameters were all named `"placeholder"`, which could be difficult to follow when debugging larger functions.

apache#10516 used the Relay parameter name when lowering to TE. However, this creates an issue when the parameter name is empty. This is legal in Relay, but results in errors during code generation. For example, this is the generated CUDA kernel for bias add: ``` extern "C" __global__ void __launch_bounds__(1024) fused_raf_op_tvm_add_kernel0( float* __restrict__ T_add, float* __restrict__ , /* Name is missing and it results in compile errors. */ float* __restrict__ _1) { T_add[((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] = ([((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] + _1[((((((int)blockIdx.x) * 16) + (((int)threadIdx.x) >> 6)) % 54) / 9)]); } ``` This PR adds "placeholder" back as a default to make sure no empty string will be passed when lowering to TE.

Lunderberg · 2022-03-07T15:41:53Z

Previously, the TE placeholders representing relay function parameters were all named "placeholder", which could be difficult to follow when debugging larger functions.

Lunderberg · 2022-03-09T15:13:25Z

Rebase onto main following #10535

Lunderberg · 2022-03-14T14:26:36Z

Rebase onto main following merge conflict with #10577.

Lunderberg · 2022-03-14T20:34:37Z

Current CI failures due to checks for the name "placeholder" in auto_scheduler (link). Will investigate further as time allows, as I'm not sure the history of this check.

tkonolige

This is a very helpful change. Thanks @Lunderberg!

areusch · 2022-04-11T23:31:07Z

@Lunderberg what's the status on this one?

Lunderberg · 2022-04-12T14:35:31Z

Currently there's one CI failure that I haven't had time to track down yet. The error message occurs during memory verification, that a buffer is undefined, but likely is caused by an earlier pass failing to update a buffer.

Lunderberg · 2022-04-19T14:21:50Z

Rebased onto main, since the failing test case was disabled in #10717.

Lunderberg · 2022-09-13T15:50:24Z

I let this one languish long enough that I no longer trust the CI checks to be up-to-date. Re-running CI.

Previously, the TE placeholders representing relay function parameters were all named `"placeholder"`, which could be difficult to follow when debugging larger functions.

The tensor name "ethos-u", once passed through from relay to TE, resulted in invalid C++ codegen.

The tensor name is part of the hash used for these results, so the previous hashes are no longer valid.

#10516 used the Relay parameter name when lowering to TE. However, this creates an issue when the parameter name is empty. This is legal in Relay, but results in errors during code generation. For example, this is the generated CUDA kernel for bias add: ``` extern "C" __global__ void __launch_bounds__(1024) fused_raf_op_tvm_add_kernel0( float* __restrict__ T_add, float* __restrict__ , /* Name is missing and it results in compile errors. */ float* __restrict__ _1) { T_add[((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] = ([((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] + _1[((((((int)blockIdx.x) * 16) + (((int)threadIdx.x) >> 6)) % 54) / 9)]); } ``` This PR adds "placeholder" back as a default to make sure no empty string will be passed when lowering to TE.

…che#10516) * [Relay][TE] Use Relay parameter name to generated TE tensor name Previously, the TE placeholders representing relay function parameters were all named `"placeholder"`, which could be difficult to follow when debugging larger functions.

apache#10516 used the Relay parameter name when lowering to TE. However, this creates an issue when the parameter name is empty. This is legal in Relay, but results in errors during code generation. For example, this is the generated CUDA kernel for bias add: ``` extern "C" __global__ void __launch_bounds__(1024) fused_raf_op_tvm_add_kernel0( float* __restrict__ T_add, float* __restrict__ , /* Name is missing and it results in compile errors. */ float* __restrict__ _1) { T_add[((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] = ([((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] + _1[((((((int)blockIdx.x) * 16) + (((int)threadIdx.x) >> 6)) % 54) / 9)]); } ``` This PR adds "placeholder" back as a default to make sure no empty string will be passed when lowering to TE.

Lunderberg · 2022-03-07T15:41:53Z

Previously, the TE placeholders representing relay function parameters were all named "placeholder", which could be difficult to follow when debugging larger functions.

Lunderberg · 2022-03-09T15:13:25Z

Rebase onto main following #10535

Lunderberg · 2022-03-14T14:26:36Z

Rebase onto main following merge conflict with #10577.

Lunderberg · 2022-03-14T20:34:37Z

Current CI failures due to checks for the name "placeholder" in auto_scheduler (link). Will investigate further as time allows, as I'm not sure the history of this check.

tkonolige

This is a very helpful change. Thanks @Lunderberg!

areusch · 2022-04-11T23:31:07Z

@Lunderberg what's the status on this one?

Lunderberg · 2022-04-12T14:35:31Z

Currently there's one CI failure that I haven't had time to track down yet. The error message occurs during memory verification, that a buffer is undefined, but likely is caused by an earlier pass failing to update a buffer.

Lunderberg · 2022-04-19T14:21:50Z

Rebased onto main, since the failing test case was disabled in #10717.

Lunderberg · 2022-09-13T15:50:24Z

I let this one languish long enough that I no longer trust the CI checks to be up-to-date. Re-running CI.

Previously, the TE placeholders representing relay function parameters were all named `"placeholder"`, which could be difficult to follow when debugging larger functions.

The tensor name "ethos-u", once passed through from relay to TE, resulted in invalid C++ codegen.

The tensor name is part of the hash used for these results, so the previous hashes are no longer valid.

#10516 used the Relay parameter name when lowering to TE. However, this creates an issue when the parameter name is empty. This is legal in Relay, but results in errors during code generation. For example, this is the generated CUDA kernel for bias add: ``` extern "C" __global__ void __launch_bounds__(1024) fused_raf_op_tvm_add_kernel0( float* __restrict__ T_add, float* __restrict__ , /* Name is missing and it results in compile errors. */ float* __restrict__ _1) { T_add[((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] = ([((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] + _1[((((((int)blockIdx.x) * 16) + (((int)threadIdx.x) >> 6)) % 54) / 9)]); } ``` This PR adds "placeholder" back as a default to make sure no empty string will be passed when lowering to TE.

…che#10516) * [Relay][TE] Use Relay parameter name to generated TE tensor name Previously, the TE placeholders representing relay function parameters were all named `"placeholder"`, which could be difficult to follow when debugging larger functions.

apache#10516 used the Relay parameter name when lowering to TE. However, this creates an issue when the parameter name is empty. This is legal in Relay, but results in errors during code generation. For example, this is the generated CUDA kernel for bias add: ``` extern "C" __global__ void __launch_bounds__(1024) fused_raf_op_tvm_add_kernel0( float* __restrict__ T_add, float* __restrict__ , /* Name is missing and it results in compile errors. */ float* __restrict__ _1) { T_add[((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] = ([((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] + _1[((((((int)blockIdx.x) * 16) + (((int)threadIdx.x) >> 6)) % 54) / 9)]); } ``` This PR adds "placeholder" back as a default to make sure no empty string will be passed when lowering to TE.

Lunderberg · 2022-03-07T15:41:53Z

Previously, the TE placeholders representing relay function parameters were all named "placeholder", which could be difficult to follow when debugging larger functions.

Lunderberg · 2022-03-09T15:13:25Z

Rebase onto main following #10535

Lunderberg · 2022-03-14T14:26:36Z

Rebase onto main following merge conflict with #10577.

Lunderberg · 2022-03-14T20:34:37Z

Current CI failures due to checks for the name "placeholder" in auto_scheduler (link). Will investigate further as time allows, as I'm not sure the history of this check.

tkonolige

This is a very helpful change. Thanks @Lunderberg!

areusch · 2022-04-11T23:31:07Z

@Lunderberg what's the status on this one?

Lunderberg · 2022-04-12T14:35:31Z

Currently there's one CI failure that I haven't had time to track down yet. The error message occurs during memory verification, that a buffer is undefined, but likely is caused by an earlier pass failing to update a buffer.

Lunderberg · 2022-04-19T14:21:50Z

Rebased onto main, since the failing test case was disabled in #10717.

Lunderberg · 2022-09-13T15:50:24Z

I let this one languish long enough that I no longer trust the CI checks to be up-to-date. Re-running CI.

Previously, the TE placeholders representing relay function parameters were all named `"placeholder"`, which could be difficult to follow when debugging larger functions.

The tensor name "ethos-u", once passed through from relay to TE, resulted in invalid C++ codegen.

The tensor name is part of the hash used for these results, so the previous hashes are no longer valid.

#10516 used the Relay parameter name when lowering to TE. However, this creates an issue when the parameter name is empty. This is legal in Relay, but results in errors during code generation. For example, this is the generated CUDA kernel for bias add: ``` extern "C" __global__ void __launch_bounds__(1024) fused_raf_op_tvm_add_kernel0( float* __restrict__ T_add, float* __restrict__ , /* Name is missing and it results in compile errors. */ float* __restrict__ _1) { T_add[((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] = ([((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] + _1[((((((int)blockIdx.x) * 16) + (((int)threadIdx.x) >> 6)) % 54) / 9)]); } ``` This PR adds "placeholder" back as a default to make sure no empty string will be passed when lowering to TE.

…che#10516) * [Relay][TE] Use Relay parameter name to generated TE tensor name Previously, the TE placeholders representing relay function parameters were all named `"placeholder"`, which could be difficult to follow when debugging larger functions.

apache#10516 used the Relay parameter name when lowering to TE. However, this creates an issue when the parameter name is empty. This is legal in Relay, but results in errors during code generation. For example, this is the generated CUDA kernel for bias add: ``` extern "C" __global__ void __launch_bounds__(1024) fused_raf_op_tvm_add_kernel0( float* __restrict__ T_add, float* __restrict__ , /* Name is missing and it results in compile errors. */ float* __restrict__ _1) { T_add[((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] = ([((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] + _1[((((((int)blockIdx.x) * 16) + (((int)threadIdx.x) >> 6)) % 54) / 9)]); } ``` This PR adds "placeholder" back as a default to make sure no empty string will be passed when lowering to TE.

Lunderberg · 2022-03-07T15:41:53Z

Previously, the TE placeholders representing relay function parameters were all named "placeholder", which could be difficult to follow when debugging larger functions.

Lunderberg · 2022-03-09T15:13:25Z

Rebase onto main following #10535

Lunderberg · 2022-03-14T14:26:36Z

Rebase onto main following merge conflict with #10577.

Lunderberg · 2022-03-14T20:34:37Z

Current CI failures due to checks for the name "placeholder" in auto_scheduler (link). Will investigate further as time allows, as I'm not sure the history of this check.

tkonolige

This is a very helpful change. Thanks @Lunderberg!

areusch · 2022-04-11T23:31:07Z

@Lunderberg what's the status on this one?

Lunderberg · 2022-04-12T14:35:31Z

Currently there's one CI failure that I haven't had time to track down yet. The error message occurs during memory verification, that a buffer is undefined, but likely is caused by an earlier pass failing to update a buffer.

Lunderberg · 2022-04-19T14:21:50Z

Rebased onto main, since the failing test case was disabled in #10717.

Lunderberg · 2022-09-13T15:50:24Z

I let this one languish long enough that I no longer trust the CI checks to be up-to-date. Re-running CI.

Previously, the TE placeholders representing relay function parameters were all named `"placeholder"`, which could be difficult to follow when debugging larger functions.

The tensor name "ethos-u", once passed through from relay to TE, resulted in invalid C++ codegen.

The tensor name is part of the hash used for these results, so the previous hashes are no longer valid.

#10516 used the Relay parameter name when lowering to TE. However, this creates an issue when the parameter name is empty. This is legal in Relay, but results in errors during code generation. For example, this is the generated CUDA kernel for bias add: ``` extern "C" __global__ void __launch_bounds__(1024) fused_raf_op_tvm_add_kernel0( float* __restrict__ T_add, float* __restrict__ , /* Name is missing and it results in compile errors. */ float* __restrict__ _1) { T_add[((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] = ([((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] + _1[((((((int)blockIdx.x) * 16) + (((int)threadIdx.x) >> 6)) % 54) / 9)]); } ``` This PR adds "placeholder" back as a default to make sure no empty string will be passed when lowering to TE.

…che#10516) * [Relay][TE] Use Relay parameter name to generated TE tensor name Previously, the TE placeholders representing relay function parameters were all named `"placeholder"`, which could be difficult to follow when debugging larger functions.

apache#10516 used the Relay parameter name when lowering to TE. However, this creates an issue when the parameter name is empty. This is legal in Relay, but results in errors during code generation. For example, this is the generated CUDA kernel for bias add: ``` extern "C" __global__ void __launch_bounds__(1024) fused_raf_op_tvm_add_kernel0( float* __restrict__ T_add, float* __restrict__ , /* Name is missing and it results in compile errors. */ float* __restrict__ _1) { T_add[((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] = ([((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] + _1[((((((int)blockIdx.x) * 16) + (((int)threadIdx.x) >> 6)) % 54) / 9)]); } ``` This PR adds "placeholder" back as a default to make sure no empty string will be passed when lowering to TE.

Lunderberg · 2022-03-07T15:41:53Z

Previously, the TE placeholders representing relay function parameters were all named "placeholder", which could be difficult to follow when debugging larger functions.

Lunderberg · 2022-03-09T15:13:25Z

Rebase onto main following #10535

Lunderberg · 2022-03-14T14:26:36Z

Rebase onto main following merge conflict with #10577.

Lunderberg · 2022-03-14T20:34:37Z

Current CI failures due to checks for the name "placeholder" in auto_scheduler (link). Will investigate further as time allows, as I'm not sure the history of this check.

tkonolige

This is a very helpful change. Thanks @Lunderberg!

areusch · 2022-04-11T23:31:07Z

@Lunderberg what's the status on this one?

Lunderberg · 2022-04-12T14:35:31Z

Currently there's one CI failure that I haven't had time to track down yet. The error message occurs during memory verification, that a buffer is undefined, but likely is caused by an earlier pass failing to update a buffer.

Lunderberg · 2022-04-19T14:21:50Z

Rebased onto main, since the failing test case was disabled in #10717.

Lunderberg · 2022-09-13T15:50:24Z

I let this one languish long enough that I no longer trust the CI checks to be up-to-date. Re-running CI.

Previously, the TE placeholders representing relay function parameters were all named `"placeholder"`, which could be difficult to follow when debugging larger functions.

The tensor name "ethos-u", once passed through from relay to TE, resulted in invalid C++ codegen.

The tensor name is part of the hash used for these results, so the previous hashes are no longer valid.

#10516 used the Relay parameter name when lowering to TE. However, this creates an issue when the parameter name is empty. This is legal in Relay, but results in errors during code generation. For example, this is the generated CUDA kernel for bias add: ``` extern "C" __global__ void __launch_bounds__(1024) fused_raf_op_tvm_add_kernel0( float* __restrict__ T_add, float* __restrict__ , /* Name is missing and it results in compile errors. */ float* __restrict__ _1) { T_add[((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] = ([((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] + _1[((((((int)blockIdx.x) * 16) + (((int)threadIdx.x) >> 6)) % 54) / 9)]); } ``` This PR adds "placeholder" back as a default to make sure no empty string will be passed when lowering to TE.

…che#10516) * [Relay][TE] Use Relay parameter name to generated TE tensor name Previously, the TE placeholders representing relay function parameters were all named `"placeholder"`, which could be difficult to follow when debugging larger functions.

apache#10516 used the Relay parameter name when lowering to TE. However, this creates an issue when the parameter name is empty. This is legal in Relay, but results in errors during code generation. For example, this is the generated CUDA kernel for bias add: ``` extern "C" __global__ void __launch_bounds__(1024) fused_raf_op_tvm_add_kernel0( float* __restrict__ T_add, float* __restrict__ , /* Name is missing and it results in compile errors. */ float* __restrict__ _1) { T_add[((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] = ([((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] + _1[((((((int)blockIdx.x) * 16) + (((int)threadIdx.x) >> 6)) % 54) / 9)]); } ``` This PR adds "placeholder" back as a default to make sure no empty string will be passed when lowering to TE.

Lunderberg · 2022-03-07T15:41:53Z

Previously, the TE placeholders representing relay function parameters were all named "placeholder", which could be difficult to follow when debugging larger functions.

Lunderberg · 2022-03-09T15:13:25Z

Rebase onto main following #10535

Lunderberg · 2022-03-14T14:26:36Z

Rebase onto main following merge conflict with #10577.

Lunderberg · 2022-03-14T20:34:37Z

Current CI failures due to checks for the name "placeholder" in auto_scheduler (link). Will investigate further as time allows, as I'm not sure the history of this check.

tkonolige

This is a very helpful change. Thanks @Lunderberg!

areusch · 2022-04-11T23:31:07Z

@Lunderberg what's the status on this one?

Lunderberg · 2022-04-12T14:35:31Z

Currently there's one CI failure that I haven't had time to track down yet. The error message occurs during memory verification, that a buffer is undefined, but likely is caused by an earlier pass failing to update a buffer.

Lunderberg · 2022-04-19T14:21:50Z

Rebased onto main, since the failing test case was disabled in #10717.

Lunderberg · 2022-09-13T15:50:24Z

I let this one languish long enough that I no longer trust the CI checks to be up-to-date. Re-running CI.

Previously, the TE placeholders representing relay function parameters were all named `"placeholder"`, which could be difficult to follow when debugging larger functions.

The tensor name "ethos-u", once passed through from relay to TE, resulted in invalid C++ codegen.

The tensor name is part of the hash used for these results, so the previous hashes are no longer valid.

#10516 used the Relay parameter name when lowering to TE. However, this creates an issue when the parameter name is empty. This is legal in Relay, but results in errors during code generation. For example, this is the generated CUDA kernel for bias add: ``` extern "C" __global__ void __launch_bounds__(1024) fused_raf_op_tvm_add_kernel0( float* __restrict__ T_add, float* __restrict__ , /* Name is missing and it results in compile errors. */ float* __restrict__ _1) { T_add[((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] = ([((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] + _1[((((((int)blockIdx.x) * 16) + (((int)threadIdx.x) >> 6)) % 54) / 9)]); } ``` This PR adds "placeholder" back as a default to make sure no empty string will be passed when lowering to TE.

…che#10516) * [Relay][TE] Use Relay parameter name to generated TE tensor name Previously, the TE placeholders representing relay function parameters were all named `"placeholder"`, which could be difficult to follow when debugging larger functions.

apache#10516 used the Relay parameter name when lowering to TE. However, this creates an issue when the parameter name is empty. This is legal in Relay, but results in errors during code generation. For example, this is the generated CUDA kernel for bias add: ``` extern "C" __global__ void __launch_bounds__(1024) fused_raf_op_tvm_add_kernel0( float* __restrict__ T_add, float* __restrict__ , /* Name is missing and it results in compile errors. */ float* __restrict__ _1) { T_add[((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] = ([((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] + _1[((((((int)blockIdx.x) * 16) + (((int)threadIdx.x) >> 6)) % 54) / 9)]); } ``` This PR adds "placeholder" back as a default to make sure no empty string will be passed when lowering to TE.

Lunderberg · 2022-03-07T15:41:53Z

Previously, the TE placeholders representing relay function parameters were all named "placeholder", which could be difficult to follow when debugging larger functions.

Lunderberg · 2022-03-09T15:13:25Z

Rebase onto main following #10535

Lunderberg · 2022-03-14T14:26:36Z

Rebase onto main following merge conflict with #10577.

Lunderberg · 2022-03-14T20:34:37Z

Current CI failures due to checks for the name "placeholder" in auto_scheduler (link). Will investigate further as time allows, as I'm not sure the history of this check.

tkonolige

This is a very helpful change. Thanks @Lunderberg!

areusch · 2022-04-11T23:31:07Z

@Lunderberg what's the status on this one?

Lunderberg · 2022-04-12T14:35:31Z

Currently there's one CI failure that I haven't had time to track down yet. The error message occurs during memory verification, that a buffer is undefined, but likely is caused by an earlier pass failing to update a buffer.

Lunderberg · 2022-04-19T14:21:50Z

Rebased onto main, since the failing test case was disabled in #10717.

Lunderberg · 2022-09-13T15:50:24Z

I let this one languish long enough that I no longer trust the CI checks to be up-to-date. Re-running CI.

Previously, the TE placeholders representing relay function parameters were all named `"placeholder"`, which could be difficult to follow when debugging larger functions.

The tensor name "ethos-u", once passed through from relay to TE, resulted in invalid C++ codegen.

The tensor name is part of the hash used for these results, so the previous hashes are no longer valid.

#10516 used the Relay parameter name when lowering to TE. However, this creates an issue when the parameter name is empty. This is legal in Relay, but results in errors during code generation. For example, this is the generated CUDA kernel for bias add: ``` extern "C" __global__ void __launch_bounds__(1024) fused_raf_op_tvm_add_kernel0( float* __restrict__ T_add, float* __restrict__ , /* Name is missing and it results in compile errors. */ float* __restrict__ _1) { T_add[((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] = ([((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] + _1[((((((int)blockIdx.x) * 16) + (((int)threadIdx.x) >> 6)) % 54) / 9)]); } ``` This PR adds "placeholder" back as a default to make sure no empty string will be passed when lowering to TE.

…che#10516) * [Relay][TE] Use Relay parameter name to generated TE tensor name Previously, the TE placeholders representing relay function parameters were all named `"placeholder"`, which could be difficult to follow when debugging larger functions.

apache#10516 used the Relay parameter name when lowering to TE. However, this creates an issue when the parameter name is empty. This is legal in Relay, but results in errors during code generation. For example, this is the generated CUDA kernel for bias add: ``` extern "C" __global__ void __launch_bounds__(1024) fused_raf_op_tvm_add_kernel0( float* __restrict__ T_add, float* __restrict__ , /* Name is missing and it results in compile errors. */ float* __restrict__ _1) { T_add[((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] = ([((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] + _1[((((((int)blockIdx.x) * 16) + (((int)threadIdx.x) >> 6)) % 54) / 9)]); } ``` This PR adds "placeholder" back as a default to make sure no empty string will be passed when lowering to TE.

Lunderberg requested review from MarisaKirisame, ZihengJiang, anijain2305, icemelon, jroesch, junrushao, jwfromm, mbrookhart, slyubomirsky, vinx13, wweic, yzhliu and zhiics as code owners March 7, 2022 15:41

Lunderberg requested review from kparzysz-quic, masahi and tqchen as code owners March 8, 2022 22:57

Lunderberg force-pushed the relay_name_to_te_placeholder branch from c37278f to 6d7b337Compare March 9, 2022 15:12

Lunderberg force-pushed the relay_name_to_te_placeholder branch from 564bf5e to c5f19dcCompare March 14, 2022 14:26

tkonolige approved these changes Mar 15, 2022
View reviewed changes

Lunderberg force-pushed the relay_name_to_te_placeholder branch from 68d2c5b to 9ac6636Compare March 22, 2022 19:10

Lunderberg force-pushed the relay_name_to_te_placeholder branch from 9ac6636 to f163bddCompare April 19, 2022 14:21

Lunderberg force-pushed the relay_name_to_te_placeholder branch 2 times, most recently from 47e9550 to 30c4c8bCompare April 26, 2022 14:14

Lunderberg added 4 commits September 13, 2022 10:54

[Relay][TE] Use Relay parameter name to generated TE tensor name
df6a6dc
Previously, the TE placeholders representing relay function parameters were all named `"placeholder"`, which could be difficult to follow when debugging larger functions.

Avoid identifiers with dash signs.
bcba05d
The tensor name "ethos-u", once passed through from relay to TE, resulted in invalid C++ codegen.

Avoid checking for an explicit "placeholder" name.
21b1764

Updated cached workload results
c7dc0ff
The tensor name is part of the hash used for these results, so the previous hashes are no longer valid.

Lunderberg force-pushed the relay_name_to_te_placeholder branch from 8a70448 to c7dc0ffCompare September 13, 2022 15:55

Lunderberg merged commit a408493 into apache:mainSep 14, 2022

Lunderberg deleted the relay_name_to_te_placeholder branch September 14, 2022 14:23

comaniac mentioned this pull request Sep 26, 2022
[Relay][TE] Add default param name if needed #12912
Merged