AMDGPU: Add __builtin_amdgcn_permlane64

[yaxunl]

LGTM. Thanks.

[arsenm]

f59f116bd5c357b1cb4b04693c88d41484e168d5

clang/include/clang/Basic/BuiltinsAMDGPU.def

 TARGET_BUILTIN(__builtin_amdgcn_image_bvh_intersect_ray_l, "V4UiWUifV4fV4fV4fV4Ui", "nc", "gfx10-insts")
 TARGET_BUILTIN(__builtin_amdgcn_image_bvh_intersect_ray_lh, "V4UiWUifV4fV4hV4hV4Ui", "nc", "gfx10-insts")
 
 
 //===----------------------------------------------------------------------===//
 // GFX11+ only builtins.
 //===----------------------------------------------------------------------===//
 
+// TODO: This is a no-op in wave32. Should the builtin require wavefrontsize64?
+TARGET_BUILTIN(__builtin_amdgcn_permlane64, "UiUi", "nc", "gfx11-insts")
+
 //===----------------------------------------------------------------------===//
 // WMMA builtins.
 // Postfix w32 indicates the builtin requires wavefront size of 32.
 // Postfix w64 indicates the builtin requires wavefront size of 64.
 //===----------------------------------------------------------------------===//
 TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_f16_w32, "V8fV16hV16hV8f", "nc", "gfx11-insts")
 TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32, "V8fV16sV16sV8f", "nc", "gfx11-insts")
 TARGET_BUILTIN(__builtin_amdgcn_wmma_f16_16x16x16_f16_w32, "V16hV16hV16hV16hIb", "nc", "gfx11-insts")

clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11.cl

 // CHECK: %1 = extractvalue { i32, i32 } %0, 0
 // CHECK: %2 = extractvalue { i32, i32 } %0, 1
 // CHECK: %3 = insertelement <2 x i32> poison, i32 %1, i64 0
 // CHECK: %4 = insertelement <2 x i32> %3, i32 %2, i64 1
 void test_ds_bvh_stack_rtn(global uint2* out, uint addr, uint data, uint4 data1)
 {
   *out = __builtin_amdgcn_ds_bvh_stack_rtn(addr, data, data1, 128);
 }
+
+// CHECK-LABEL: @test_permlane64(
+// CHECK: call i32 @llvm.amdgcn.permlane64(i32 %a)
+void test_permlane64(global uint* out, uint a) {
+  *out = __builtin_amdgcn_permlane64(a);
+}

clang/test/SemaOpenCL/builtins-amdgcn-gfx11.cl

 // RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1030 -verify=GFX10 -S -o - %s
 // RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1100 -verify=GFX11 -S -o - %s
 
 typedef unsigned int uint;
 typedef unsigned long ulong;
 
 void test(global uint* out1, global ulong* out2, int x) {
   *out1 = __builtin_amdgcn_s_sendmsg_rtn(0); // GFX10-error {{'__builtin_amdgcn_s_sendmsg_rtn' needs target feature gfx11-insts}}
   *out2 = __builtin_amdgcn_s_sendmsg_rtnl(0); // GFX10-error {{'__builtin_amdgcn_s_sendmsg_rtnl' needs target feature gfx11-insts}}
 #if __has_builtin(__builtin_amdgcn_s_sendmsg_rtn)
   *out1 = __builtin_amdgcn_s_sendmsg_rtn(x); // GFX11-error {{argument to '__builtin_amdgcn_s_sendmsg_rtn' must be a constant integer}}
 #endif
 #if __has_builtin(__builtin_amdgcn_s_sendmsg_rtnl)
   *out2 = __builtin_amdgcn_s_sendmsg_rtnl(x); // GFX11-error {{argument to '__builtin_amdgcn_s_sendmsg_rtnl' must be a constant integer}}
 #endif
+
+  *out1 = __builtin_amdgcn_permlane64(x); // GFX10-error {{'__builtin_amdgcn_permlane64' needs target feature gfx11-insts}}
 }

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

Intrinsic<[llvm_v4i32_ty],
             [IntrReadMem, IntrWillReturn]>;
 
 //===----------------------------------------------------------------------===//
 // GFX11 Intrinsics
 //===----------------------------------------------------------------------===//
 
 // llvm.amdgcn.permlane64 <src0>
 def int_amdgcn_permlane64 :
+  ClangBuiltin<"__builtin_amdgcn_permlane64">,
   Intrinsic<[llvm_i32_ty], [llvm_i32_ty],
             [IntrNoMem, IntrConvergent, IntrWillReturn]>;
 
 def int_amdgcn_ds_add_gs_reg_rtn :
   ClangBuiltin<"__builtin_amdgcn_ds_add_gs_reg_rtn">,
   Intrinsic<[llvm_anyint_ty], [llvm_i32_ty, llvm_i32_ty],
             [ImmArg<ArgIndex<1>>, IntrHasSideEffects, IntrWillReturn]>;