Skip to content

Conversation

@skc7
Copy link
Collaborator

@skc7 skc7 commented Dec 22, 2025

This PR implements AMDGPU target support for AST->CIR codgen.

  • Adds AMDGPUTargetCIRGenInfo with:
    • getOpenCLKernelCallingConv() returning AMDGPUKernel
    • setCUDAKernelCallingConvention() for HIP/CUDA kernel support

With changes in PR:

  • OpenCL kernels on AMDGPU get cc(amdgpu_kernel) calling convention
  • HIP __global__ kernels get cc(amdgpu_kernel) calling convention

@skc7 skc7 marked this pull request as ready for review December 22, 2025 16:14
Copy link
Member

@bcardosolopes bcardosolopes left a comment

Choose a reason for hiding this comment

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

This is not the right place for callconv lowering, there's a dedicated pass for this, you should teach that instead.

@skc7 skc7 force-pushed the skc7/cir_amdgpu_p1_abi branch from 0eb6b76 to ff60eaf Compare December 23, 2025 05:54
@skc7 skc7 force-pushed the skc7/cir_amdgpu_p1_abi branch from f8301fc to 8c5988b Compare December 23, 2025 06:06
@skc7 skc7 changed the title [CIR][AMDGPU] Add basic ABI implementation of AMDGPU for AST to CIR [CIR][AMDGPU] Add AMDGPU target support for AST to CIR Dec 23, 2025
@skc7
Copy link
Collaborator Author

skc7 commented Dec 23, 2025

This is not the right place for callconv lowering, there's a dedicated pass for this, you should teach that instead.

@bcardosolopes
Updated PR to add amdgpu_kernel CC enum and removed any CallingConv lowering in TargetInfo module.

@skc7 skc7 changed the title [CIR][AMDGPU] Add AMDGPU target support for AST to CIR [CIR][AMDGPU] Add AMDGPU target support for AST to CIR codegen Dec 23, 2025
@bcardosolopes bcardosolopes merged commit addef71 into llvm:main Dec 24, 2025
6 checks passed
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