arsenm added inline comments.
================ Comment at: clang/include/clang/Basic/LangOptions.def:274 LANGOPT(OffloadingNewDriver, 1, 0, "use the new driver for generating offloading code.") +ENUM_LANGOPT(AMDGPUPrintfKindVal, AMDGPUPrintfKind, 2, AMDGPUPrintfKind::Buffered, "printf lowering scheme to be used, hostcall or buffer based") ---------------- yaxunl wrote: > This should be a target option like https://reviews.llvm.org/D91546 instead > of a language option since it is target specific. Should be -m option ================ Comment at: clang/lib/Driver/ToolChains/Clang.cpp:4673 + + // unconditionally claim the pritnf option now to avoid unused diagnostic. + // TODO: OpenCL targets will should use this option to switch between ---------------- Capitalize, typo 'pritnf' ================ Comment at: clang/lib/Driver/ToolChains/Clang.cpp:4674 + // unconditionally claim the pritnf option now to avoid unused diagnostic. + // TODO: OpenCL targets will should use this option to switch between + // hostcall and buffered printf schemes. ---------------- Typo 'will should' Don't really understand the TODO, this should trigger for OpenCL as it is ================ Comment at: clang/test/CodeGenHIP/printf_nonhostcall.cpp:5 +// RUN: -o - %s | FileCheck --enable-var-scope %s + +#define __device__ __attribute__((device)) ---------------- Do we need a test with -fno-builtin? ================ Comment at: clang/test/CodeGenHIP/printf_nonhostcall.cpp:69 + const char *s = "hello world"; + return printf("%.*f %*.*s %p\n", 8, 3.14159, 8, 4, s, s); +} ---------------- Doesn't cover the full range of printable types. need some other non-string pointers and different address spaces, some FP promotions, 16 and 64 bit integers ================ Comment at: llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp:259-261 + // This is a tradeoff. we might end up taking more compile + // time to calculate string contents if possible, but the generated + // code would be better runtime wise. ---------------- Don't understand the point of the comment, I would assume anything involving analysis of a constant string has ignorable compile time ================ Comment at: llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp:371-384 + Function *TheFn = Intrinsic::getDeclaration( + Builder.GetInsertBlock()->getModule(), Intrinsic::memcpy, Tys); + SmallVector<Value *, 1> BuffOffset; + + Value *FnArgs[] = { + PtrToStore, Args[i], val, + ConstantInt::get(Type::getInt1Ty(Builder.getContext()), false)}; ---------------- Builder.CreateMemCpy ================ Comment at: llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp:402 + SmallVector<Value *, 1> BuffOffset; + uint offsetVal = toStore->getType()->getIntegerBitWidth() == 32 ? 4 : 8; + BuffOffset.push_back(ConstantInt::get(Builder.getInt32Ty(), offsetVal)); ---------------- Just use the type store size ================ Comment at: llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp:403 + uint offsetVal = toStore->getType()->getIntegerBitWidth() == 32 ? 4 : 8; + BuffOffset.push_back(ConstantInt::get(Builder.getInt32Ty(), offsetVal)); + ---------------- You don't need a SmallVector to push back a single entry ================ Comment at: llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp:428 + SmallVector<StringData, 8> StringContents; + llvm::Module *M = Builder.GetInsertBlock()->getModule(); + LLVMContext &Ctx = Builder.getContext(); ---------------- Don't need llvm:: ================ Comment at: llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp:430 + LLVMContext &Ctx = Builder.getContext(); + auto Int1Ty = Builder.getInt1Ty(); + auto Int8Ty = Builder.getInt8Ty(); ---------------- No auto ================ Comment at: llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp:458 + auto CreateControlDWord = M->getOrInsertFunction( + StringRef("__ockl_create_control_dword"), Builder.getInt32Ty(), + Builder.getInt32Ty(), Int1Ty, Int1Ty); ---------------- Do we really need another ockl control variable for this? Why isn't it a parameter? printf=stdout always ================ Comment at: llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp:473 + if (!FmtStr.empty()) { + llvm::MD5 Hasher; + llvm::MD5::MD5Result Hash; ---------------- Don't need llvm:: ================ Comment at: llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp:498 + NamedMDNode *metaD = M->getOrInsertNamedMetadata("llvm.printf.fmts"); + if(0 == metaD->getNumOperands()) { + MDString *fmtStrArray = MDString::get(Ctx, "0:0:deadbeef,\"\""); ---------------- Backwards conditional Repository: rG LLVM Github Monorepo CHANGES SINCE LAST ACTION https://reviews.llvm.org/D150427/new/ https://reviews.llvm.org/D150427 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits