Skip to content

[mlir] Change default NVVM compilation to fatbin from bin #70052

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Oct 24, 2023

Conversation

grypp
Copy link
Member

@grypp grypp commented Oct 24, 2023

Change the NVVM assembly to fatbin as it is executable for multiple architectures. Using bin caused test errors at runtime in the test systems.

Change the NVVM assembly to `fatbin` as it is executable for multiple architectures. Using `bin` caused test errors at runtime in the test systems.
@llvmbot
Copy link
Member

llvmbot commented Oct 24, 2023

@llvm/pr-subscribers-mlir

@llvm/pr-subscribers-mlir-gpu

Author: Guray Ozen (grypp)

Changes

Change the NVVM assembly to fatbin as it is executable for multiple architectures. Using bin caused test errors at runtime in the test systems.


Full diff: https://github.com/llvm/llvm-project/pull/70052.diff

1 Files Affected:

  • (modified) mlir/test/lib/Dialect/GPU/TestLowerToNVVM.cpp (+2-2)
diff --git a/mlir/test/lib/Dialect/GPU/TestLowerToNVVM.cpp b/mlir/test/lib/Dialect/GPU/TestLowerToNVVM.cpp
index c4cc0d5ae38d9be..ed7634fbecf49fd 100644
--- a/mlir/test/lib/Dialect/GPU/TestLowerToNVVM.cpp
+++ b/mlir/test/lib/Dialect/GPU/TestLowerToNVVM.cpp
@@ -60,7 +60,7 @@ struct TestLowerToNVVMOptions
   PassOptions::Option<std::string> cubinFormat{
       *this, "cubin-format",
       llvm::cl::desc("Compilation format to use to serialize to cubin."),
-      llvm::cl::init("bin")};
+      llvm::cl::init("fatbin")};
   PassOptions::Option<int> optLevel{
       *this, "opt-level",
       llvm::cl::desc("Optimization level for NVVM compilation"),
@@ -156,4 +156,4 @@ void registerTestLowerToNVVM() {
 }
 } // namespace test
 } // namespace mlir
-#endif // MLIR_CUDA_CONVERSIONS_ENABLED
+#endif // MLIR_CUDA_CONVERSIONS_ENABLED
\ No newline at end of file

@grypp
Copy link
Member Author

grypp commented Oct 24, 2023

Submitting this as it should fix the integration test problem below:

******************** TEST 'MLIR :: Integration/GPU/CUDA/all-reduce-and.mlir' FAILED ********************
Exit Code: 1

Command Output (stdout):
--
# RUN: at line 1
/vol/worker/mlir-nvidia/mlir-nvidia/llvm.obj/bin/mlir-opt /vol/worker/mlir-nvidia/mlir-nvidia/llvm.src/mlir/test/Integration/GPU/CUDA/all-reduce-and.mlir  | /vol/worker/mlir-nvidia/mlir-nvidia/llvm.obj/bin/mlir-opt -test-lower-to-nvvm  | /vol/worker/mlir-nvidia/mlir-nvidia/llvm.obj/bin/mlir-cpu-runner    --shared-libs=/vol/worker/mlir-nvidia/mlir-nvidia/llvm.obj/lib/libmlir_cuda_runtime.so    --shared-libs=/vol/worker/mlir-nvidia/mlir-nvidia/llvm.obj/lib/libmlir_runner_utils.so    --entry-point-result=void  | /vol/worker/mlir-nvidia/mlir-nvidia/llvm.obj/bin/FileCheck /vol/worker/mlir-nvidia/mlir-nvidia/llvm.src/mlir/test/Integration/GPU/CUDA/all-reduce-and.mlir
# executed command: /vol/worker/mlir-nvidia/mlir-nvidia/llvm.obj/bin/mlir-opt /vol/worker/mlir-nvidia/mlir-nvidia/llvm.src/mlir/test/Integration/GPU/CUDA/all-reduce-and.mlir
# executed command: /vol/worker/mlir-nvidia/mlir-nvidia/llvm.obj/bin/mlir-opt -test-lower-to-nvvm
# executed command: /vol/worker/mlir-nvidia/mlir-nvidia/llvm.obj/bin/mlir-cpu-runner --shared-libs=/vol/worker/mlir-nvidia/mlir-nvidia/llvm.obj/lib/libmlir_cuda_runtime.so --shared-libs=/vol/worker/mlir-nvidia/mlir-nvidia/llvm.obj/lib/libmlir_runner_utils.so --entry-point-result=void
# .---command stderr------------
# | 'cuModuleLoadData(&module, data)' failed with 'CUDA_ERROR_INVALID_SOURCE'
# | 'cuModuleGetFunction(&function, module, name)' failed with 'CUDA_ERROR_INVALID_HANDLE'
# | 'cuFuncSetAttribute( function, CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, smem)' 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'

@grypp grypp merged commit 8875f78 into llvm:main Oct 24, 2023
@grypp grypp deleted the fix-nvvm-bin branch October 24, 2023 14:56
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Projects
None yet
Development

Successfully merging this pull request may close these issues.

2 participants