Skip to content

For Metal project---adding drivers#92

Draft
gulugulubing wants to merge 5 commits into
libmir:masterfrom
gulugulubing:startOfMetal
Draft

For Metal project---adding drivers#92
gulugulubing wants to merge 5 commits into
libmir:masterfrom
gulugulubing:startOfMetal

Conversation

@gulugulubing
Copy link
Copy Markdown

Start a minimal Metal bindings and mimic cuda or ocl's style. I am not sure whether it's a right direction.
One major blocker was an Objective-C interop RTTI/linker problem in D:

  • Storing extern(Objective-C) interface types directly in wrapper structs sometimes(not always, but at most time, strange) caused LDC to emit RTTI/classinfo references such as __Interface / __Class.
  • Those symbols are not provided by the Objective-C runtime, which led to undefined-symbol linker errors.
  • The fix was to store Objective-C handles as raw void* internally and cast at the wrapper boundary, avoiding unwanted RTTI emission.

@gulugulubing gulugulubing marked this pull request as draft May 8, 2026 04:19
import ldc.dcompute;
pure: nothrow: @nogc:

pragma(LDC_intrinsic, "air.get_global_id.i32") uint get_global_id(uint dim);
Copy link
Copy Markdown
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@asindarov, is this style of index consistent with your compiler's side progress? And If you are working on how to inline these functions, I found MPM.addPass(llvm::ModuleInlinerWrapperPass()); works https://github.com/gulugulubing/ldc/tree/Metal-OS26-LLVM20.

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

yes, it is. I will see if I can run the module level pass during code gen for llvm bitcode


//suspends work-item execution until all work-items in the work-group have called the barrier
void barrier()
void barrier()()
Copy link
Copy Markdown
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

At first, I didn't wanna touch the codes related to cuda and ocl, but I found I had to use a zero-argument template to avoid compiling this file. This is related to #95

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants