[AMDGPU] Add missing __builtin_amdgcn_wavefrontsize builtin (#80741)

Summary:
The backend supports the wavefrontsize intrinsic, and suggests that it
is tied to a corresponding clang builtin, but it is not actually
present. This simply adds it in so it can be used from clang. This
attribute likely isn't the best to rely on, but for the `libc` use-case
we will need to detect a struct's differing size in a way that will
depend on the wavefront size.

Author: jhuber6

clang/include/clang/Basic/BuiltinsAMDGPU.def

@@ -69,6 +69,7 @@ BUILTIN(__builtin_amdgcn_s_dcache_inv, "v", "n")
 BUILTIN(__builtin_amdgcn_buffer_wbinvl1, "v", "n")
 BUILTIN(__builtin_amdgcn_fence, "vUicC*", "n")
 BUILTIN(__builtin_amdgcn_groupstaticsize, "Ui", "n")
+BUILTIN(__builtin_amdgcn_wavefrontsize, "Ui", "nc")
 
 BUILTIN(__builtin_amdgcn_atomic_inc32, "UZiUZiD*UZiUicC*", "n")
 BUILTIN(__builtin_amdgcn_atomic_inc64, "UWiUWiD*UWiUicC*", "n")

clang/test/CodeGenOpenCL/builtins-amdgcn.cl

@@ -832,6 +832,13 @@ void test_atomic_inc_dec(local uint *lptr, global uint *gptr, uint val) {
   res = __builtin_amdgcn_atomic_dec32((volatile global uint*)gptr, val, __ATOMIC_SEQ_CST, "");
 }
 
+// CHECK-LABEL test_wavefrontsize(
+unsigned test_wavefrontsize() {
+
+  // CHECK: call i32 @llvm.amdgcn.wavefrontsize()
+  return __builtin_amdgcn_wavefrontsize();
+}
+
 // CHECK-DAG: [[$WI_RANGE]] = !{i32 0, i32 1024}
 // CHECK-DAG: [[$WS_RANGE]] = !{i16 1, i16 1025}
 // CHECK-DAG: attributes #[[$NOUNWIND_READONLY]] = { convergent mustprogress nocallback nofree nounwind willreturn memory(none) }