Unverified Commit 857baf37 authored by David Rivera's avatar David Rivera Committed by GitHub
Browse files

[CIR][AMDGPU] Set amdgpu_kernel calling convention on HIP kernels (#195381)

parent 8c77d077

clang/lib/CIR/CodeGen/Targets/AMDGPU.cpp

+2 −1
Original line number Diff line number Diff line
@@ -240,7 +240,8 @@ void clang::CIRGen::setAMDGPUTargetFunctionAttributes(const Decl *decl, 
    const bool isHIPKernel =

        cgm.getLangOpts().HIP && fd->hasAttr<CUDAGlobalAttr>();



    // TODO(CIR): Set amdgpu_kernel calling convention for HIP kernels.

    if (isHIPKernel)

      func.setCallingConv(cir::CallingConv::AMDGPUKernel);



    handleAMDGPUFlatWorkGroupSizeAttr(fd, func, cgm, builder, isOpenCLKernel,

                                      isHIPKernel);

clang/test/CIR/CodeGenHIP/amdgpu-attrs.hip

+2 −2
Original line number Diff line number Diff line
@@ -25,8 +25,8 @@ 
// Test that AMDGPU-specific attributes are generated for HIP kernels



// Test: Default attributes for simple kernel

// CIR: cir.func{{.*}} @_Z13kernel_simplev(){{.*}}"cir.amdgpu-flat-work-group-size" = "1,1024"

// LLVM: define{{.*}} void @_Z13kernel_simplev(){{.*}} #[[SIMPLE_ATTR:[0-9]+]]

// CIR: cir.func{{.*}} @_Z13kernel_simplev() cc(amdgpu_kernel){{.*}}"cir.amdgpu-flat-work-group-size" = "1,1024"

// LLVM: define{{.*}} amdgpu_kernel void @_Z13kernel_simplev(){{.*}} #[[SIMPLE_ATTR:[0-9]+]]

__global__ void kernel_simple() {}



// Test: Explicit flat work group size attribute