[GFX13][Clang] Add w32 predicate to DDS Clang builtins (#3471)

DDS instructions should only we used in wave32. Add the predicates to
the builtins, so clang can give an error if this is attempted.

Author: Sisyph

clang/include/clang/Basic/BuiltinsAMDGPU.def

@@ -1216,11 +1216,11 @@ TARGET_BUILTIN(__builtin_amdgcn_load_mcast_b32, "vi*10iC*Iii", "n", "gfx13-insts
 TARGET_BUILTIN(__builtin_amdgcn_load_mcast_b64, "vV2i*10V2iC*Iii", "n", "gfx13-insts,wavefrontsize32")
 TARGET_BUILTIN(__builtin_amdgcn_load_mcast_b128, "vV4i*10V4iC*Iii", "n", "gfx13-insts,wavefrontsize32")
 
-TARGET_BUILTIN(__builtin_amdgcn_map_shared_rank, "v*11v*3i", "n", "gfx13-insts")
-TARGET_BUILTIN(__builtin_amdgcn_query_shared_rank, "ivC*11", "n", "gfx13-insts")
-TARGET_BUILTIN(__builtin_amdgcn_dds_load_async_to_lds_b32, "viC*11i*3IiIi", "n", "gfx13-insts")
-TARGET_BUILTIN(__builtin_amdgcn_dds_load_async_to_lds_b64, "vV2iC*11V2i*3IiIi", "n", "gfx13-insts")
-TARGET_BUILTIN(__builtin_amdgcn_dds_load_async_to_lds_b128, "vV4iC*11V4i*3IiIi", "n", "gfx13-insts")
+TARGET_BUILTIN(__builtin_amdgcn_map_shared_rank, "v*11v*3i", "n", "gfx13-insts,wavefrontsize32")
+TARGET_BUILTIN(__builtin_amdgcn_query_shared_rank, "ivC*11", "n", "gfx13-insts,wavefrontsize32")
+TARGET_BUILTIN(__builtin_amdgcn_dds_load_async_to_lds_b32, "viC*11i*3IiIi", "n", "gfx13-insts,wavefrontsize32")
+TARGET_BUILTIN(__builtin_amdgcn_dds_load_async_to_lds_b64, "vV2iC*11V2i*3IiIi", "n", "gfx13-insts,wavefrontsize32")
+TARGET_BUILTIN(__builtin_amdgcn_dds_load_async_to_lds_b128, "vV4iC*11V4i*3IiIi", "n", "gfx13-insts,wavefrontsize32")
 TARGET_BUILTIN(__builtin_amdgcn_dds_load_async_mcast_to_lds_b32, "viC*11i*3IiIii", "n", "gfx13-insts,wavefrontsize32")
 TARGET_BUILTIN(__builtin_amdgcn_dds_load_async_mcast_to_lds_b64, "vV2iC*11V2i*3IiIii", "n", "gfx13-insts,wavefrontsize32")
 TARGET_BUILTIN(__builtin_amdgcn_dds_load_async_mcast_to_lds_b128, "vV4iC*11V4i*3IiIii", "n", "gfx13-insts,wavefrontsize32")

clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13-w64-err.cl

@@ -0,0 +1,19 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu gfx1300 -target-feature -wavefrontsize64 \
+// RUN: -emit-llvm -verify -o - %s
+
+typedef int    v2i   __attribute__((ext_vector_type(2)));
+typedef int    v4i   __attribute__((ext_vector_type(4)));
+
+__attribute__((address_space(11))) void*
+test_amdgcn_map_shared_rank(__local void* ptr, int rank, __attribute__((address_space(11))) const void* ptrsr,
+                            __attribute__((address_space(11))) const int* gaddr, __attribute__((address_space(11))) const v2i* gaddr2,
+                            __attribute__((address_space(11))) const v4i* gaddr4, local int* laddr, local v2i* laddr2, local v4i* laddr4)
+{
+  __builtin_amdgcn_query_shared_rank(ptrsr); // expected-error{{'__builtin_amdgcn_query_shared_rank' needs target feature gfx13-insts,wavefrontsize32}}
+  __builtin_amdgcn_dds_load_async_to_lds_b32(gaddr, laddr, 16, 0); // expected-error{{'__builtin_amdgcn_dds_load_async_to_lds_b32' needs target feature gfx13-insts,wavefrontsize32}}
+  __builtin_amdgcn_dds_load_async_to_lds_b64(gaddr2, laddr2, 16, 0); // expected-error{{'__builtin_amdgcn_dds_load_async_to_lds_b64' needs target feature gfx13-insts,wavefrontsize32}}
+  __builtin_amdgcn_dds_load_async_to_lds_b128(gaddr4, laddr4, 16, 0); // expected-error{{'__builtin_amdgcn_dds_load_async_to_lds_b128' needs target feature gfx13-insts,wavefrontsize32}}
+  return __builtin_amdgcn_map_shared_rank(ptr, rank); // expected-error{{'__builtin_amdgcn_map_shared_rank' needs target feature gfx13-insts,wavefrontsize32}}
+}
+