<div dir="ltr"><div dir="auto">Thanks for the reply! Indeed declaring the function myself somewhere is probably the easiest workaround.<div dir="auto"><br></div><div>I asked in the <a href="https://forums.developer.nvidia.com/t/cuda-11-4-cooperative-groups-no-longer-supported-on-sm-7-0/194001/2">Nvidia forums</a> and it seems they forgot to guard the problematic code inside an #ifdef that is only active for SM 7.x or later. I can see they've fixed it in CUDA 11.5, and indeed using CUDA 11.5 with Clang solves the problem (modulo Clang warning due to "too new CUDA version").</div><div><br></div><div>Therefore my next question is: would it be possible to bump the supported CUDA version in Clang to 11.5? I could do it myself following the steps done for 11.4, but I'm not sure what things I should be looking out for or how to test it. Let me know if I should open a new thread for this new question.</div><div><br></div><div>/Carlos</div></div></div><br><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Thu, Nov 4, 2021, 18:52 Johannes Doerfert <<a href="mailto:johannesdoerfert@gmail.com" target="_blank">johannesdoerfert@gmail.com</a>> wrote:<br></div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">It's not a CUDA issue per se (IMHO) and not related to the builtin at all.<br>
<br>
C++ says this ain't alright, nvcc is just very forgiving apparently:<br>
<a href="https://godbolt.org/z/14xYfbKdz" rel="noreferrer noreferrer" target="_blank">https://godbolt.org/z/14xYfbKdz</a><br>
<br>
I'd recommend to either declare the builtin yourself in you header or use<br>
a CUDA version that comes with a declaration.<br>
<br>
~ Johannes<br>
<br>
<br>
On 11/4/21 02:44, Carlos Galvez via cfe-dev wrote:<br>
> Hi!<br>
><br>
> I wanted to ask what's the expected support for CUDA 11.4 in Clang? We want<br>
> to update from 10.2 to 11.4, and came across the following issue when<br>
> compiling:<br>
><br>
> // <a href="http://foo.cu" rel="noreferrer noreferrer" target="_blank">foo.cu</a><br>
> #include <cooperative_groups.h><br>
><br>
> Just by including this header, and compiling for "--cuda-gpu-arch=sm_61",<br>
> we get the following error building with Clang:<br>
><br>
> /usr/local/cuda-11.4/targets/x86_64-linux/include/cooperative_groups/details/partitioning.h:85:32:<br>
> error: use of undeclared identifier '__match_any_sync'<br>
>          unsigned int subMask = __match_any_sync(thisMask, pred);<br>
>                                 ^<br>
><br>
> The error goes away if we use sm_70 or later, but we need sm_61 in my<br>
> application.<br>
><br>
> Cooperative groups have been working just fine since they were introduced<br>
> in CUDA 9, so I wonder why this would break now. NVCC has no problems with<br>
> the above code. Clang can also build it just fine on CUDA 10.2.<br>
><br>
> Digging a bit deeper, I can narrow it down to this:<br>
><br>
> // <a href="http://foo.cu" rel="noreferrer noreferrer" target="_blank">foo.cu</a><br>
><br>
> template <typename T><br>
> __device__ void foo()<br>
> {<br>
>      __match_any_sync(0,0);<br>
> }<br>
><br>
> The above code compiles fine with NVCC + sm_61. However, *if I remove the<br>
> template*, then NVCC complains that __match_any_sync is not defined. Clang<br>
> complains regardless.<br>
> Does this make any sense? Is NVCC doing some magic that allows it to ignore<br>
> undeclared functions if they are used in a non-instantiated template<br>
> function?<br>
><br>
> PS: from what I understand, __match_any_sync is indeed only available on SM<br>
> 7.0 or later. I still wonder how NVCC can handle this, is it a bug or a<br>
> feature? Should Clang be updated to match NVCC?<br>
><br>
> PS2: The problem goes away in CUDA 11.5, even though it's not officially<br>
> supported in Clang. I get a warning but the code compiles. I can still see<br>
> the __match_any_sync function in the CUDA header so I don't understand<br>
> what's happening. Is Clang really compiling, or what does it do when using<br>
> a "too new" CUDA version?<br>
><br>
> Thanks!<br>
><br>
><br>
> _______________________________________________<br>
> cfe-dev mailing list<br>
> <a href="mailto:cfe-dev@lists.llvm.org" rel="noreferrer" target="_blank">cfe-dev@lists.llvm.org</a><br>
> <a href="https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-dev" rel="noreferrer noreferrer" target="_blank">https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-dev</a><br>
</blockquote></div>