<div dir="ltr">Justin, <div><br></div><div>I am confused. I didn't find any link that says "<span style="font-size:12.8000001907349px">CUDA interface expects kernel parameters to be in the generic address space". Instead, I found </span><a href="https://urldefense.proofpoint.com/v2/url?u=http-3A__docs.nvidia.com_cuda_parallel-2Dthread-2Dexecution_-23kernel-2Dfunction-2Dparameters&d=AwMFaQ&c=8hUWFZcy2Z-Za5rBPlktOQ&r=mQ4LZ2PUj9hpadE3cDHZnIdEwhEBrbAstXeMaFoB9tg&m=QMR7k4yt-zpNBNSir9MWrsksLrGgzGl8A-A8jypW4Co&s=KD9xqa-6r4rEXLqoam48vM3WO7N39PVaDZjM5h2Q3AY&e=">PTX ISA</a> says "... kernel function parameters ... may hold address to objects in constant, global, local, or shared state spaces". PTX can even attach .ptr metadata to pointer parameters to indicate which address space. When the .ptr annotation is missing, PTX assumes the pointer parameter is generic. </div><div><br></div><div>Am I missing something? </div><div><br></div><div>Jingyue</div></div><div class="gmail_extra"><br><div class="gmail_quote">On Tue, Jun 30, 2015 at 9:26 AM, Justin Holewinski <span dir="ltr"><<a href="mailto:jholewinski@nvidia.com" target="_blank">jholewinski@nvidia.com</a>></span> wrote:<br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">Hi Tobias,<br>
<br>
One thing to keep in mind is that the CUDA interface expects kernel parameters to be in the generic address space.  If you write kernels that pass arguments as global memory pointers, you are technically violating the contract between host and device.  Now, this just so happens to work since the global->generic and generic->global pointer conversions are no-ops, but as far as I know this is not actually guaranteed and may not be true in the future.<br>
<span class=""><br>
> On Jun 26, 2015, at 6:36 PM, Jingyue Wu <<a href="mailto:jingyue@google.com">jingyue@google.com</a>> wrote:<br>
><br>
> REPOSITORY<br>
>  rL LLVM<br>
><br>
> <a href="https://urldefense.proofpoint.com/v2/url?u=http-3A__reviews.llvm.org_D10779&d=AwMFaQ&c=8hUWFZcy2Z-Za5rBPlktOQ&r=mQ4LZ2PUj9hpadE3cDHZnIdEwhEBrbAstXeMaFoB9tg&m=QMR7k4yt-zpNBNSir9MWrsksLrGgzGl8A-A8jypW4Co&s=jwL6J6DYl0hZ0b5OPBfZCiyGOL5_cgXOD0i3u2yQigE&e=" rel="noreferrer" target="_blank">http://reviews.llvm.org/D10779</a><br>
><br>
> Files:<br>
>  llvm/trunk/lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp<br>
>  llvm/trunk/test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll<br>
><br>
> Index: llvm/trunk/lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp<br>
> ===================================================================<br>
> --- llvm/trunk/lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp<br>
> +++ llvm/trunk/lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp<br>
</span><span class="">> @@ -132,6 +132,10 @@<br>
>   assert(!Arg->hasByValAttr() &&<br>
>          "byval params should be handled by handleByValParam");<br>
><br>
> +  // Do nothing if the argument already points to the global address space.<br>
> +  if (Arg->getType()->getPointerAddressSpace() == ADDRESS_SPACE_GLOBAL)<br>
> +    return;<br>
> +<br>
>   Instruction *FirstInst = Arg->getParent()->getEntryBlock().begin();<br>
>   Instruction *ArgInGlobal = new AddrSpaceCastInst(<br>
>       Arg, PointerType::get(Arg->getType()->getPointerElementType(),<br>
</span><span class="">> Index: llvm/trunk/test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll<br>
> ===================================================================<br>
> --- llvm/trunk/test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll<br>
> +++ llvm/trunk/test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll<br>
</span><span class="">> @@ -16,5 +16,16 @@<br>
>   ret void<br>
> }<br>
><br>
> -!nvvm.annotations = !{!0}<br>
> +define void @kernel2(float addrspace(1)* %input, float addrspace(1)* %output) {<br>
> +; CHECK-LABEL: .visible .entry kernel2(<br>
> +; CHECK-NOT: cvta.to.global.u64<br>
> +  %1 = load float, float addrspace(1)* %input, align 4<br>
> +; CHECK: ld.global.f32<br>
> +  store float %1, float addrspace(1)* %output, align 4<br>
> +; CHECK: st.global.f32<br>
> +  ret void<br>
> +}<br>
> +<br>
> +!nvvm.annotations = !{!0, !1}<br>
> !0 = !{void (float*, float*)* @kernel, !"kernel", i32 1}<br>
> +!1 = !{void (float addrspace(1)*, float addrspace(1)*)* @kernel2, !"kernel", i32 1}<br>
><br>
> EMAIL PREFERENCES<br>
>  <a href="https://urldefense.proofpoint.com/v2/url?u=http-3A__reviews.llvm.org_settings_panel_emailpreferences_&d=AwMFaQ&c=8hUWFZcy2Z-Za5rBPlktOQ&r=mQ4LZ2PUj9hpadE3cDHZnIdEwhEBrbAstXeMaFoB9tg&m=QMR7k4yt-zpNBNSir9MWrsksLrGgzGl8A-A8jypW4Co&s=XbXDOkVsIaxR09V83F1M89x7ztJjqfLcw06ZDw302Y4&e=" rel="noreferrer" target="_blank">http://reviews.llvm.org/settings/panel/emailpreferences/</a><br>
</span>> <D10779.28605.patch><br>
<br>
-----------------------------------------------------------------------------------<br>
This email message is for the sole use of the intended recipient(s) and may contain<br>
confidential information.  Any unauthorized review, use, disclosure or distribution<br>
is prohibited.  If you are not the intended recipient, please contact the sender by<br>
reply email and destroy all copies of the original message.<br>
-----------------------------------------------------------------------------------<br>
</blockquote></div><br></div>