Author: Arseniy Obolenskiy Date: 2026-10-13T14:20:37+02:00 New Revision: 56aa8ae4e5fc2653c1b74e87f6bfd0e771ff70ee
URL: https://github.com/llvm/llvm-project/commit/56aa8ae4e5fc2653c1b74e87f6bfd0e771ff70ee DIFF: https://github.com/llvm/llvm-project/commit/56aa8ae4e5fc2653c1b74e87f6bfd0e771ff70ee.diff LOG: [CIR][AMDGPU] Add support for AMDGCN workgroup_size builtins (#231156) Port EmitAMDGPUWorkGroupSize from classic codegen Added: clang/test/CIR/CodeGenHIP/amdgpu-workgroup-size.hip Modified: clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip Removed: ################################################################################ diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 4b28d8c402a16..f28bfc67d9cc5 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -22,6 +22,16 @@ using namespace clang; using namespace clang::CIRGen; using namespace cir; +// Emit a no-operand intrinsic returning a pointer in the constant AS. +static mlir::Value emitAMDGPUConstantPtrIntrinsic(CIRGenFunction &cgf, + mlir::Location loc, + llvm::StringRef name) { + CIRGenBuilderTy &builder = cgf.getBuilder(); + mlir::Type retTy = builder.getVoidPtrTy(cir::TargetAddressSpaceAttr::get( + builder.getContext(), llvm::AMDGPUAS::CONSTANT_ADDRESS)); + return builder.emitIntrinsicCallOp(loc, name, retTy, mlir::ValueRange{}); +} + /// Map a constant integral to memory order. static cir::MemOrder decodeAtomicOrder(const Expr *arg, ASTContext &ctx) { Expr::EvalResult orderRes; @@ -108,13 +118,8 @@ static mlir::Value emitAMDGPUDispatchPtr(CIRGenFunction &cgf, mlir::Location loc, const CallExpr *e = nullptr) { CIRGenBuilderTy &builder = cgf.getBuilder(); - // The intrinsic always returns a pointer in the constant AS. - mlir::Type retTy = cir::PointerType::get( - cir::VoidType::get(builder.getContext()), - cir::TargetAddressSpaceAttr::get(builder.getContext(), - llvm::AMDGPUAS::CONSTANT_ADDRESS)); - mlir::Value call = builder.emitIntrinsicCallOp(loc, "amdgcn.dispatch.ptr", - retTy, mlir::ValueRange{}); + mlir::Value call = + emitAMDGPUConstantPtrIntrinsic(cgf, loc, "amdgcn.dispatch.ptr"); if (!e) return call; // Only cast when the caller-visible AS diff ers from the intrinsic's AS; @@ -126,6 +131,23 @@ static mlir::Value emitAMDGPUDispatchPtr(CIRGenFunction &cgf, return builder.createAddrSpaceCast(loc, call, expectedPtrTy); } +static mlir::Value emitAMDGPUImplicitArgPtr(CIRGenFunction &cgf, + mlir::Location loc) { + // OG also marks the result align 8 dereferenceable(256). + assert(!cir::MissingFeatures::opCallAttrs()); + return emitAMDGPUConstantPtrIntrinsic(cgf, loc, "amdgcn.implicitarg.ptr"); +} + +static mlir::Value emitAMDGPUWorkGroupID(CIRGenFunction &cgf, + mlir::Location loc, unsigned index) { + static constexpr llvm::StringLiteral names[] = {"amdgcn.workgroup.id.x", + "amdgcn.workgroup.id.y", + "amdgcn.workgroup.id.z"}; + CIRGenBuilderTy &builder = cgf.getBuilder(); + return builder.emitIntrinsicCallOp(loc, names[index], builder.getUInt32Ty(), + mlir::ValueRange{}); +} + // Load a \p ty value at byte \p offset from \p base as an invariant load. static mlir::Value emitAMDGPUInvariantLoad(CIRGenFunction &cgf, mlir::Location loc, mlir::Value base, @@ -150,6 +172,112 @@ static mlir::Value emitAMDGPUGridSize(CIRGenFunction &cgf, mlir::Location loc, builder.getUInt32Ty(), CharUnits::fromQuantity(4)); } +static mlir::Value emitAMDGPUWorkGroupSizeV5(CIRGenFunction &cgf, + mlir::Location loc, + unsigned index) { + CIRGenBuilderTy &builder = cgf.getBuilder(); + mlir::Value implicitArgPtr = emitAMDGPUImplicitArgPtr(cgf, loc); + + // offsetof(amdhsa_implicit_kernarg_v5, block_count[index]) + unsigned blockCountOffset = 0 + index * 4; + // offsetof(amdhsa_implicit_kernarg_v5, group_size[index]) + unsigned groupSizeOffset = 12 + index * 2; + // offsetof(amdhsa_implicit_kernarg_v5, remainder[index]) + unsigned remainderOffset = 18 + index * 2; + + mlir::Value sizeOffset; + if (cgf.getLangOpts().OffloadUniformBlock) { + sizeOffset = builder.getUInt64(groupSizeOffset, loc); + } else { + mlir::Value blockCount = emitAMDGPUInvariantLoad( + cgf, loc, implicitArgPtr, builder.getUInt64(blockCountOffset, loc), + builder.getUInt32Ty(), CharUnits::fromQuantity(4)); + mlir::Value workgroupID = emitAMDGPUWorkGroupID(cgf, loc, index); + mlir::Value isFull = + builder.createCompare(loc, cir::CmpOpKind::lt, workgroupID, blockCount); + sizeOffset = builder.createSelect(loc, isFull, + builder.getUInt64(groupSizeOffset, loc), + builder.getUInt64(remainderOffset, loc)); + } + + mlir::Value size = emitAMDGPUInvariantLoad(cgf, loc, implicitArgPtr, + sizeOffset, builder.getUInt16Ty(), + CharUnits::fromQuantity(2)); + return builder.createIntCast(size, builder.getUInt32Ty()); +} + +static mlir::Value emitAMDGPUWorkGroupSizeV4(CIRGenFunction &cgf, + mlir::Location loc, + unsigned index) { + CIRGenBuilderTy &builder = cgf.getBuilder(); + mlir::Value dispatchPtr = emitAMDGPUDispatchPtr(cgf, loc); + + // Indexing the HSA kernel_dispatch_packet struct. + mlir::Value groupSizeLd = emitAMDGPUInvariantLoad( + cgf, loc, dispatchPtr, builder.getUInt64(4 + index * 2, loc), + builder.getUInt16Ty(), CharUnits::fromQuantity(2)); + mlir::Value groupSize = + builder.createIntCast(groupSizeLd, builder.getUInt32Ty()); + + if (cgf.getLangOpts().OffloadUniformBlock) + return groupSize; + + mlir::Value workgroupID = emitAMDGPUWorkGroupID(cgf, loc, index); + mlir::Value gridSize = emitAMDGPUInvariantLoad( + cgf, loc, dispatchPtr, builder.getUInt64(12 + index * 4, loc), + builder.getUInt32Ty(), CharUnits::fromQuantity(4)); + mlir::Value mul = builder.createMul(loc, workgroupID, groupSize); + mlir::Value remainder = builder.createSub(loc, gridSize, mul); + mlir::Value isPartial = + builder.createCompare(loc, cir::CmpOpKind::lt, remainder, groupSize); + return builder.createSelect(loc, isPartial, remainder, groupSize); +} + +// \p index is 0, 1, and 2 for x, y, and z dimension, respectively. +// Emit code based on Code Object ABI version. +// COV_4 : Emit code to use dispatch ptr +// COV_5+ : Emit code to use implicitarg ptr +// COV_NONE : Emit code to load a global variable "__oclc_ABI_version" +// and use its value for COV_4 or COV_5+ approach. It is used for +// compiling device libraries in an ABI-agnostic way. +static mlir::Value emitAMDGPUWorkGroupSize(CIRGenFunction &cgf, + mlir::Location loc, unsigned index) { + llvm::CodeObjectVersionKind cov = + cgf.getTarget().getTargetOpts().CodeObjectVersion; + + // Do not emit __oclc_ABI_version references with non-empty environment. + if (cov == llvm::CodeObjectVersionKind::COV_None && + cgf.getTarget().getTriple().hasEnvironment()) + cov = llvm::CodeObjectVersionKind::COV_6; + + if (cov != llvm::CodeObjectVersionKind::COV_None) + return cov >= llvm::CodeObjectVersionKind::COV_5 + ? emitAMDGPUWorkGroupSizeV5(cgf, loc, index) + : emitAMDGPUWorkGroupSizeV4(cgf, loc, index); + + CIRGenBuilderTy &builder = cgf.getBuilder(); + CIRGenModule &cgm = cgf.cgm; + StringRef name = "__oclc_ABI_version"; + mlir::Type int32Ty = builder.getSInt32Ty(); + auto abiVersionGV = + mlir::dyn_cast_or_null<cir::GlobalOp>(cgm.getGlobalValue(name)); + if (!abiVersionGV) + abiVersionGV = + cgm.createGlobalOp(loc, name, int32Ty, /*isConstant=*/false, + cir::toCIRAddressSpaceAttr(cgm.getMLIRContext(), + LangAS::opencl_constant)); + mlir::Value abiVersionPtr = builder.createGetGlobal(loc, abiVersionGV); + mlir::Value abiVersion = builder.createAlignedLoad( + loc, int32Ty, abiVersionPtr, + cgf.getContext().getTypeAlignInChars(cgf.getContext().IntTy)); + mlir::Value isCOV5 = builder.createCompare( + loc, cir::CmpOpKind::ge, abiVersion, + builder.getConstantInt(loc, int32Ty, llvm::CodeObjectVersionKind::COV_5)); + mlir::Value v5Impl = emitAMDGPUWorkGroupSizeV5(cgf, loc, index); + mlir::Value v4Impl = emitAMDGPUWorkGroupSizeV4(cgf, loc, index); + return builder.createSelect(loc, isCOV5, v5Impl, v4Impl); +} + static mlir::Value emitBinaryExpMaybeConstrainedFPBuiltin( CIRGenFunction &cgf, const CallExpr *e, llvm::StringRef intrinsicName, llvm::StringRef constrainedIntrinsicName) { @@ -1090,13 +1218,11 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, } // amdgcn workgroup size case AMDGPU::BI__builtin_amdgcn_workgroup_size_x: + return emitAMDGPUWorkGroupSize(*this, getLoc(expr->getExprLoc()), 0); case AMDGPU::BI__builtin_amdgcn_workgroup_size_y: - case AMDGPU::BI__builtin_amdgcn_workgroup_size_z: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; - } + return emitAMDGPUWorkGroupSize(*this, getLoc(expr->getExprLoc()), 1); + case AMDGPU::BI__builtin_amdgcn_workgroup_size_z: + return emitAMDGPUWorkGroupSize(*this, getLoc(expr->getExprLoc()), 2); // amdgcn grid size case AMDGPU::BI__builtin_amdgcn_grid_size_x: return emitAMDGPUGridSize(*this, getLoc(expr->getExprLoc()), 0); diff --git a/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip b/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip index 25aa1c4a7d187..99e2ba1c8e237 100644 --- a/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip +++ b/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip @@ -138,3 +138,66 @@ __device__ unsigned test_grid_size_z() { // OGCG: [[P:%.*]] = call{{.*}} ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() // OGCG: [[GEP:%.*]] = getelementptr i8, ptr addrspace(4) [[P]], i32 20 // OGCG: load i32, ptr addrspace(4) [[GEP]], align 4, !range !{{.*}}, !invariant.load + +__device__ unsigned test_workgroup_size_x() { + return __builtin_amdgcn_workgroup_size_x(); +} + +// CIR-LABEL: cir.func no_inline @_Z21test_workgroup_size_xv +// CIR: %[[P:.*]] = cir.call_llvm_intrinsic "amdgcn.implicitarg.ptr" : () -> !cir.ptr<!void, target_address_space(4)> +// CIR: %[[GEP:.*]] = cir.ptr_stride %[[P]], %{{.*}} : (!cir.ptr<!void, target_address_space(4)>, !u64i) +// CIR: cir.load invariant align(2) %{{.*}} : !cir.ptr<!u16i, target_address_space(4)>, !u16i + +// LLVM-LABEL: define spir_func noundef i32 @_Z21test_workgroup_size_xv +// LLVM: [[P:%.*]] = call{{.*}} ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() +// LLVM: [[GEP:%.*]] = getelementptr i8, ptr addrspace(4) [[P]], i64 12 +// LLVM: [[SIZE:%.*]] = load i16, ptr addrspace(4) [[GEP]], align 2, !invariant.load +// LLVM: zext i16 [[SIZE]] to i32 + +// OGCG-LABEL: define spir_func noundef i32 @_Z21test_workgroup_size_xv +// OGCG: [[P:%.*]] = call{{.*}} ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() +// OGCG: [[GEP:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[P]], i64 12 +// OGCG: [[SIZE:%.*]] = load i16, ptr addrspace(4) [[GEP]], align 2, !range !{{.*}}, !invariant.load !{{.*}}, !noundef +// OGCG: zext i16 [[SIZE]] to i32 + +__device__ unsigned test_workgroup_size_y() { + return __builtin_amdgcn_workgroup_size_y(); +} + +// CIR-LABEL: cir.func no_inline @_Z21test_workgroup_size_yv +// CIR: %[[P:.*]] = cir.call_llvm_intrinsic "amdgcn.implicitarg.ptr" : () -> !cir.ptr<!void, target_address_space(4)> +// CIR: %[[GEP:.*]] = cir.ptr_stride %[[P]], %{{.*}} : (!cir.ptr<!void, target_address_space(4)>, !u64i) +// CIR: cir.load invariant align(2) %{{.*}} : !cir.ptr<!u16i, target_address_space(4)>, !u16i + +// LLVM-LABEL: define spir_func noundef i32 @_Z21test_workgroup_size_yv +// LLVM: [[P:%.*]] = call{{.*}} ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() +// LLVM: [[GEP:%.*]] = getelementptr i8, ptr addrspace(4) [[P]], i64 14 +// LLVM: [[SIZE:%.*]] = load i16, ptr addrspace(4) [[GEP]], align 2, !invariant.load +// LLVM: zext i16 [[SIZE]] to i32 + +// OGCG-LABEL: define spir_func noundef i32 @_Z21test_workgroup_size_yv +// OGCG: [[P:%.*]] = call{{.*}} ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() +// OGCG: [[GEP:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[P]], i64 14 +// OGCG: [[SIZE:%.*]] = load i16, ptr addrspace(4) [[GEP]], align 2, !range !{{.*}}, !invariant.load !{{.*}}, !noundef +// OGCG: zext i16 [[SIZE]] to i32 + +__device__ unsigned test_workgroup_size_z() { + return __builtin_amdgcn_workgroup_size_z(); +} + +// CIR-LABEL: cir.func no_inline @_Z21test_workgroup_size_zv +// CIR: %[[P:.*]] = cir.call_llvm_intrinsic "amdgcn.implicitarg.ptr" : () -> !cir.ptr<!void, target_address_space(4)> +// CIR: %[[GEP:.*]] = cir.ptr_stride %[[P]], %{{.*}} : (!cir.ptr<!void, target_address_space(4)>, !u64i) +// CIR: cir.load invariant align(2) %{{.*}} : !cir.ptr<!u16i, target_address_space(4)>, !u16i + +// LLVM-LABEL: define spir_func noundef i32 @_Z21test_workgroup_size_zv +// LLVM: [[P:%.*]] = call{{.*}} ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() +// LLVM: [[GEP:%.*]] = getelementptr i8, ptr addrspace(4) [[P]], i64 16 +// LLVM: [[SIZE:%.*]] = load i16, ptr addrspace(4) [[GEP]], align 2, !invariant.load +// LLVM: zext i16 [[SIZE]] to i32 + +// OGCG-LABEL: define spir_func noundef i32 @_Z21test_workgroup_size_zv +// OGCG: [[P:%.*]] = call{{.*}} ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() +// OGCG: [[GEP:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[P]], i64 16 +// OGCG: [[SIZE:%.*]] = load i16, ptr addrspace(4) [[GEP]], align 2, !range !{{.*}}, !invariant.load !{{.*}}, !noundef +// OGCG: zext i16 [[SIZE]] to i32 diff --git a/clang/test/CIR/CodeGenHIP/amdgpu-workgroup-size.hip b/clang/test/CIR/CodeGenHIP/amdgpu-workgroup-size.hip new file mode 100644 index 0000000000000..e5b3f28efc3be --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/amdgpu-workgroup-size.hip @@ -0,0 +1,57 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fclangir -fcuda-is-device \ +// RUN: -mcode-object-version=4 -fno-offload-uniform-block -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=COV4 %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \ +// RUN: -mcode-object-version=4 -fno-offload-uniform-block -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=COV4 %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fclangir -fcuda-is-device \ +// RUN: -mcode-object-version=5 -fno-offload-uniform-block -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=COV5 %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \ +// RUN: -mcode-object-version=5 -fno-offload-uniform-block -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=COV5 %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fclangir -fcuda-is-device \ +// RUN: -mcode-object-version=none -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=NONE %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \ +// RUN: -mcode-object-version=none -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=NONE %s + +#define __device__ __attribute__((device)) + +// NONE: @__oclc_ABI_version = external addrspace(4) global i32 + +// COV4-LABEL: @_Z4testv +// COV4: [[DP:%.*]] = call{{.*}} ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() +// COV4: [[GEP:%.*]] = getelementptr{{.*}} i8, ptr addrspace(4) [[DP]], i64 6 +// COV4: [[LD:%.*]] = load i16, ptr addrspace(4) [[GEP]], align 2 +// COV4: [[GS:%.*]] = zext i16 [[LD]] to i32 +// COV4: [[ID:%.*]] = call i32 @llvm.amdgcn.workgroup.id.y() +// COV4: [[GGEP:%.*]] = getelementptr{{.*}} i8, ptr addrspace(4) [[DP]], i64 16 +// COV4: [[GRID:%.*]] = load i32, ptr addrspace(4) [[GGEP]], align 4 +// COV4: [[MUL:%.*]] = mul i32 [[ID]], [[GS]] +// COV4: [[REM:%.*]] = sub i32 [[GRID]], [[MUL]] +// COV4: [[CMP:%.*]] = icmp ult i32 [[REM]], [[GS]] +// COV4: select i1 [[CMP]], i32 [[REM]], i32 [[GS]] + +// COV5-LABEL: @_Z4testv +// COV5: [[IP:%.*]] = call{{.*}} ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() +// COV5: [[BGEP:%.*]] = getelementptr{{.*}} i8, ptr addrspace(4) [[IP]], i64 4 +// COV5: [[BC:%.*]] = load i32, ptr addrspace(4) [[BGEP]], align 4 +// COV5: [[ID:%.*]] = call i32 @llvm.amdgcn.workgroup.id.y() +// COV5: [[FULL:%.*]] = icmp ult i32 [[ID]], [[BC]] +// COV5: [[OFF:%.*]] = select i1 [[FULL]], i{{32|64}} 14, i{{32|64}} 20 +// COV5: [[SGEP:%.*]] = getelementptr{{.*}} i8, ptr addrspace(4) [[IP]], i{{32|64}} [[OFF]] +// COV5: [[SZ:%.*]] = load i16, ptr addrspace(4) [[SGEP]], align 2 +// COV5: zext i16 [[SZ]] to i32 + +// NONE-LABEL: @_Z4testv +// NONE: [[ABI:%.*]] = load i32, ptr addrspace(4) @__oclc_ABI_version, align 4 +// NONE: [[IS5:%.*]] = icmp sge i32 [[ABI]], 500 +// NONE: call{{.*}} ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() +// NONE: [[V5:%.*]] = zext i16 %{{.*}} to i32 +// NONE: call{{.*}} ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() +// NONE: [[V4:%.*]] = zext i16 %{{.*}} to i32 +// NONE: select i1 [[IS5]], i32 [[V5]], i32 [[V4]] +__device__ unsigned test() { return __builtin_amdgcn_workgroup_size_y(); } diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip index 768936bbcb8d3..180f632b5050b 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip @@ -132,6 +132,30 @@ __device__ void test_grid_size(unsigned *out) { out[2] = __builtin_amdgcn_grid_size_z(); } +// CIR-LABEL: @_Z19test_workgroup_sizePj +// CIR: %[[P:.*]] = cir.call_llvm_intrinsic "amdgcn.implicitarg.ptr" : () -> !cir.ptr<!void, target_address_space(4)> +// CIR: %[[OFF:.*]] = cir.const #cir.int<12> : !u64i +// CIR: %[[GEP:.*]] = cir.ptr_stride %[[P]], %[[OFF]] +// CIR: %[[PTR:.*]] = cir.cast bitcast %[[GEP]] : !cir.ptr<!void, target_address_space(4)> -> !cir.ptr<!u16i, target_address_space(4)> +// CIR: %[[SIZE:.*]] = cir.load invariant align(2) %[[PTR]] +// CIR: cir.cast integral %[[SIZE]] : !u16i -> !u32i +// CIR: cir.const #cir.int<14> : !u64i +// CIR: cir.const #cir.int<16> : !u64i +// LLVM-LABEL: @_Z19test_workgroup_sizePj +// LLVM: [[P:%.*]] = call{{.*}} ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() +// LLVM: [[GEP:%.*]] = getelementptr{{.*}} i8, ptr addrspace(4) [[P]], i64 12 +// LLVM: [[SIZE:%.*]] = load i16, ptr addrspace(4) [[GEP]], align 2, {{.*}}!invariant.load +// LLVM: zext i16 [[SIZE]] to i32 +// LLVM: getelementptr{{.*}} i8, ptr addrspace(4) %{{.*}}, i64 14 +// LLVM: getelementptr{{.*}} i8, ptr addrspace(4) %{{.*}}, i64 16 +// LLVM-NOT: @llvm.amdgcn.workgroup.id +// LLVM: ret void +__device__ void test_workgroup_size(unsigned *out) { + out[0] = __builtin_amdgcn_workgroup_size_x(); + out[1] = __builtin_amdgcn_workgroup_size_y(); + out[2] = __builtin_amdgcn_workgroup_size_z(); +} + // CIR-LABEL: @_Z12test_rcp_f32Pff // CIR: cir.call_llvm_intrinsic "amdgcn.rcp" {{.*}} : (!cir.float) -> !cir.float // LLVM: define{{.*}} void @_Z12test_rcp_f32Pff _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
