[QST] How to avoid too many resources requested

[YSF-A]

What is your question?

I try to use the cutlass::conv::device::Convolution with the fixed ThreadblockShape, WarpShape and InstructionShape. There is internal error which is too many resources requested actually. It may be useful to modify the ThreadblockShape or WarpShape. Is there any other solutions? For example, __launch_bounds__ may be useful in such case of cuda kernel.

[hwu36]

what shapes do you want to use? we don't want to use too many resources which can hurt the performance.

[YSF-A]

@hwu36 hi, the ThreadblockShape is [128, 128, 32], the WarpShape is [32, 32, 32], the InstructionShape is [8, 8, 4].

[hwu36]

What is data type and architecture?

[YSF-A]

What is data type and architecture?

Data type is float16. Data type of internal accumulation is float32. GPU is 2080ti.

[hwu36]

your threadblock size is 128x128, your warp size is 32x32. so you need 128/32 x 128/32 = 16 warps. we usually use 4 or 8 warps. so you'd better use warp size 64x64. if i am not wrong, 2080ti is a turing card. it better use instruction shape 16x8x8.

so you could use threadblock shape 128x128x32, warp shape 64x64x32, and instruction shape 16x8x8. you could find more plausible tile sizes in our profilers https://github.com/NVIDIA/cutlass/blob/main/python/cutlass_library/generator.py#L1375-L1383

[wxthu]

how can I convert these tile descriptions to TVM tile size shape? thanks

[hwu36]

how can I convert these tile descriptions to TVM tile size shape? thanks

@masahi

[masahi]

What do you mean by "TVM tile size shape"?

[github-actions]

This issue has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d if there is no activity in the next 60 days.

[github-actions]

This issue has been labeled inactive-90d due to no recent activity in the past 90 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed.