arsenm added inline comments.
================ Comment at: clang/lib/CodeGen/CGBuiltin.cpp:13428 + return Call; + return CGF.Builder.CreateAddrSpaceCast(Call, RetTy); +} ---------------- Why is this necessary? The builtin always has the same return type? ================ Comment at: clang/lib/CodeGen/CGBuiltin.cpp:13435 + auto *DP = EmitAMDGPUDispatchPtr(CGF); + auto *Offset = llvm::ConstantInt::get(CGF.Int32Ty, XOffset + Index * 2); + auto *GEP = CGF.Builder.CreateGEP(DP, Offset); ---------------- Comment that this is indexing the hsa_kernel_dispatch_packet sstruct? ================ Comment at: clang/lib/CodeGen/CGBuiltin.cpp:13442 + llvm::MDBuilder MDHelper(CGF.getLLVMContext()); + llvm::MDNode *RNode = MDHelper.createRange(APInt(16, 1), APInt(16, 1025)); + LD->setMetadata(llvm::LLVMContext::MD_range, RNode); ---------------- I thought I had a patch to include the maximum group size in AMDGPUTargetInfo to avoid hardcoding it, but I guess it was never committed ================ Comment at: clang/lib/CodeGen/CGBuiltin.cpp:13443 + llvm::MDNode *RNode = MDHelper.createRange(APInt(16, 1), APInt(16, 1025)); + LD->setMetadata(llvm::LLVMContext::MD_range, RNode); + return LD; ---------------- Also set it's invariant ================ Comment at: clang/test/CodeGenOpenCL/builtins-amdgcn.cl:539 +void test_get_workgroup_size(int d, global int *out) +{ + switch (d) { ---------------- Also run in a hip test, or some case where the addrspacecast is needed? CHANGES SINCE LAST ACTION https://reviews.llvm.org/D76772/new/ https://reviews.llvm.org/D76772 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits