[PATCH] D150427: [AMDGPU] Non hostcall printf support for HIP

[Vikram Hegde via Phabricator via llvm-commits]

vikramRH marked 4 inline comments as done.
vikramRH added inline comments.


================
Comment at: clang/test/CodeGenHIP/printf_nonhostcall.cpp:137
+
+__device__ float f1 = 3.14f;
+__device__ double f2 = 2.71828;
----------------
arsenm wrote:
> Also half 
C++ default arg promotions does not seem to list float16 case and consequently clang does not handle it. I have handled this case by promoting to double currenlty. Clang however still emits a warning when used with a %f conv specifier


================
Comment at: llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp:386-387
+    } else {
+      auto IntTy = dyn_cast<IntegerType>(Args[i]->getType());
+      if (IntTy && IntTy->getBitWidth() == 32)
+        WhatToStore.push_back(
----------------
arsenm wrote:
> isIntegerTy(32).
> 
> I also do not understand why only 32-bit integers would be promoted to 64-bit (or why this would be zext). This doesn't match varargs ABI handling, where everything smaller would be promoted to i32.
Right, But default promotions would have happened already, this additional promotion is due to runtime processing requirement


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D150427/new/

https://reviews.llvm.org/D150427