Revert "[mlir][gpu] Fallback to JIT compilation" "[mlir][gpu] Increase default SM...
authorMehdi Amini <joker.eph@gmail.com>
Mon, 24 Jul 2023 17:20:02 +0000 (10:20 -0700)
committerMehdi Amini <joker.eph@gmail.com>
Mon, 24 Jul 2023 17:23:15 +0000 (10:23 -0700)
commit5e8a1164f22757c19387abf79608e99782b8ec38
tree175dc5694137b043cdf0baafb2adb1bc2dbb1593
parent033efb3fe8c7be58803ecdb7edd5c509233d1850
Revert "[mlir][gpu] Fallback to JIT compilation" "[mlir][gpu] Increase default SM version from 35 to 50" and  "[mlir][gpu] Improving Cubin Serialization with ptxas Compiler"

This reverts commit 2e0e00ed841951e358a85a871647be9b3a622f51
and reverts commit a6eb40692c795a9cc29266779ceca2e304141114
and reverts commit 585cbe3f639783bf0307b47504acbd205f135310.

15 tests are broken on the mlir-nvidia buildbot:

'cuModuleLoadData(&module, data)' failed with 'CUDA_ERROR_INVALID_SOURCE'
'cuModuleGetFunction(&function, module, name)' failed with 'CUDA_ERROR_INVALID_HANDLE'
'cuLaunchKernel(function, gridX, gridY, gridZ, blockX, blockY, blockZ, smem, stream, params, extra)' failed with 'CUDA_ERROR_INVALID_HANDLE'
'cuModuleUnload(module)' failed with 'CUDA_ERROR_INVALID_HANDLE'
mlir/include/mlir/Dialect/GPU/Transforms/Passes.h
mlir/lib/Dialect/GPU/Transforms/SerializeToCubin.cpp
mlir/lib/Dialect/SparseTensor/Pipelines/SparseTensorPipelines.cpp
mlir/test/lib/Dialect/GPU/TestConvertGPUKernelToCubin.cpp
mlir/test/lib/Dialect/GPU/TestLowerToNVVM.cpp