Fix compile issue for Marin qqq on sm<8.0 #1651
Open
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Closes #1648
Test on Google Colab T4 (sm75): https://colab.research.google.com/drive/107-fKXymnK-QNCvnfXggTK8SSqy9oz8o?usp=sharing
The main issue is that
TORCH_CHECK_NOT_IMPLEMENTED()
should not be called inside__global__
function. However, I also believe the aggressive use of__CUDA_ARCH__
guard will lead to undefined behavior, as outlined in here (https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#cuda-arch). Note that__CUDA_ARCH__
is not defined in host code.Following the practice from @alexsamardzic, I simply let the kernel empty when
__CUDA_ARCH__ < 800
, and add a runtime check in the torch function.at::cuda::getCurrentDeviceProperties()
call.Run
python benchmarks/benchmark_marlin_qqq.py
on 4070Ti SUPERmain
this PR