[AMDGPU] Split dot-insts feature

[kzhuravl]

LGTM

include/clang/Basic/BuiltinsAMDGPU.def

 //===----------------------------------------------------------------------===//
 
 TARGET_BUILTIN(__builtin_amdgcn_fmed3h, "hhhh", "nc", "gfx9-insts")
 
 //===----------------------------------------------------------------------===//
 // Deep learning builtins.
 //===----------------------------------------------------------------------===//
 
-TARGET_BUILTIN(__builtin_amdgcn_fdot2, "fV2hV2hfIb", "nc", "dot-insts")
-TARGET_BUILTIN(__builtin_amdgcn_sdot2, "SiV2SsV2SsSiIb", "nc", "dot-insts")
-TARGET_BUILTIN(__builtin_amdgcn_udot2, "UiV2UsV2UsUiIb", "nc", "dot-insts")
-TARGET_BUILTIN(__builtin_amdgcn_sdot4, "SiSiSiSiIb", "nc", "dot-insts")
-TARGET_BUILTIN(__builtin_amdgcn_udot4, "UiUiUiUiIb", "nc", "dot-insts")
-TARGET_BUILTIN(__builtin_amdgcn_sdot8, "SiSiSiSiIb", "nc", "dot-insts")
-TARGET_BUILTIN(__builtin_amdgcn_udot8, "UiUiUiUiIb", "nc", "dot-insts")
+TARGET_BUILTIN(__builtin_amdgcn_fdot2, "fV2hV2hfIb", "nc", "dot2-insts")
+TARGET_BUILTIN(__builtin_amdgcn_sdot2, "SiV2SsV2SsSiIb", "nc", "dot2-insts")
+TARGET_BUILTIN(__builtin_amdgcn_udot2, "UiV2UsV2UsUiIb", "nc", "dot2-insts")
+TARGET_BUILTIN(__builtin_amdgcn_sdot4, "SiSiSiSiIb", "nc", "dot1-insts")
+TARGET_BUILTIN(__builtin_amdgcn_udot4, "UiUiUiUiIb", "nc", "dot2-insts")
+TARGET_BUILTIN(__builtin_amdgcn_sdot8, "SiSiSiSiIb", "nc", "dot1-insts")
+TARGET_BUILTIN(__builtin_amdgcn_udot8, "UiUiUiUiIb", "nc", "dot2-insts")
 
 //===----------------------------------------------------------------------===//
 // Special builtins.
 //===----------------------------------------------------------------------===//
 BUILTIN(__builtin_amdgcn_read_exec, "LUi", "nc")
 BUILTIN(__builtin_amdgcn_read_exec_lo, "Ui", "nc")
 BUILTIN(__builtin_amdgcn_read_exec_hi, "Ui", "nc")
 

lib/Basic/Targets/AMDGPU.cpp

bool AMDGPUTargetInfo::initFeatureMap(
   // XXX - What does the member GPU mean if device name string passed here?
   if (isAMDGCN(getTriple())) {
     if (CPU.empty())
       CPU = "gfx600";
 
     switch (llvm::AMDGPU::parseArchAMDGCN(CPU)) {
     case GK_GFX906:
       Features["dl-insts"] = true;
-      Features["dot-insts"] = true;
+      Features["dot1-insts"] = true;
+      Features["dot2-insts"] = true;
       LLVM_FALLTHROUGH;
     case GK_GFX909:
     case GK_GFX904:
     case GK_GFX902:
     case GK_GFX900:
       Features["gfx9-insts"] = true;
       LLVM_FALLTHROUGH;
     case GK_GFX810:

test/CodeGenOpenCL/amdgpu-features.cl

 // REQUIRES: amdgpu-registered-target
 
 // Check that appropriate features are defined for every supported AMDGPU
 // "-target" and "-mcpu" options.
 
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx904 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX904 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX906 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx801 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX801 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx700 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX700 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx600 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX600 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx601 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX601 %s
 
 // GFX904: "target-features"="+16-bit-insts,+ci-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx9-insts,+s-memrealtime,+vi-insts"
-// GFX906: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx9-insts,+s-memrealtime,+vi-insts"
+// GFX906: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx9-insts,+s-memrealtime,+vi-insts"
 // GFX801: "target-features"="+16-bit-insts,+ci-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+s-memrealtime,+vi-insts"
 // GFX700: "target-features"="+ci-insts,+fp64-fp16-denormals,-fp32-denormals"
 // GFX600: "target-features"="+fp64-fp16-denormals,-fp32-denormals"
 // GFX601: "target-features"="+fp64-fp16-denormals,-fp32-denormals"
 
 kernel void test() {}

test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl

 // REQUIRES: amdgpu-registered-target
 
 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -verify -S -emit-llvm -o - %s
 
 typedef unsigned int uint;
 typedef half __attribute__((ext_vector_type(2))) half2;
 typedef short __attribute__((ext_vector_type(2))) short2;
 typedef unsigned short __attribute__((ext_vector_type(2))) ushort2;
 
 kernel void builtins_amdgcn_dl_insts_err(
     global float *fOut, global int *siOut, global uint *uiOut,
     half2 v2hA, half2 v2hB, float fC,
     short2 v2ssA, short2 v2ssB, int siA, int siB, int siC,
     ushort2 v2usA, ushort2 v2usB, uint uiA, uint uiB, uint uiC) {
-  fOut[0] = __builtin_amdgcn_fdot2(v2hA, v2hB, fC, false);     // expected-error {{'__builtin_amdgcn_fdot2' needs target feature dot-insts}}
-  fOut[1] = __builtin_amdgcn_fdot2(v2hA, v2hB, fC, true);      // expected-error {{'__builtin_amdgcn_fdot2' needs target feature dot-insts}}
+  fOut[0] = __builtin_amdgcn_fdot2(v2hA, v2hB, fC, false);     // expected-error {{'__builtin_amdgcn_fdot2' needs target feature dot2-insts}}
+  fOut[1] = __builtin_amdgcn_fdot2(v2hA, v2hB, fC, true);      // expected-error {{'__builtin_amdgcn_fdot2' needs target feature dot2-insts}}
 
-  siOut[0] = __builtin_amdgcn_sdot2(v2ssA, v2ssB, siC, false); // expected-error {{'__builtin_amdgcn_sdot2' needs target feature dot-insts}}
-  siOut[1] = __builtin_amdgcn_sdot2(v2ssA, v2ssB, siC, true);  // expected-error {{'__builtin_amdgcn_sdot2' needs target feature dot-insts}}
+  siOut[0] = __builtin_amdgcn_sdot2(v2ssA, v2ssB, siC, false); // expected-error {{'__builtin_amdgcn_sdot2' needs target feature dot2-insts}}
+  siOut[1] = __builtin_amdgcn_sdot2(v2ssA, v2ssB, siC, true);  // expected-error {{'__builtin_amdgcn_sdot2' needs target feature dot2-insts}}
 
-  uiOut[0] = __builtin_amdgcn_udot2(v2usA, v2usB, uiC, false); // expected-error {{'__builtin_amdgcn_udot2' needs target feature dot-insts}}
-  uiOut[1] = __builtin_amdgcn_udot2(v2usA, v2usB, uiC, true);  // expected-error {{'__builtin_amdgcn_udot2' needs target feature dot-insts}}
+  uiOut[0] = __builtin_amdgcn_udot2(v2usA, v2usB, uiC, false); // expected-error {{'__builtin_amdgcn_udot2' needs target feature dot2-insts}}
+  uiOut[1] = __builtin_amdgcn_udot2(v2usA, v2usB, uiC, true);  // expected-error {{'__builtin_amdgcn_udot2' needs target feature dot2-insts}}
 
-  siOut[2] = __builtin_amdgcn_sdot4(siA, siB, siC, false);     // expected-error {{'__builtin_amdgcn_sdot4' needs target feature dot-insts}}
-  siOut[3] = __builtin_amdgcn_sdot4(siA, siB, siC, true);      // expected-error {{'__builtin_amdgcn_sdot4' needs target feature dot-insts}}
+  siOut[2] = __builtin_amdgcn_sdot4(siA, siB, siC, false);     // expected-error {{'__builtin_amdgcn_sdot4' needs target feature dot1-insts}}
+  siOut[3] = __builtin_amdgcn_sdot4(siA, siB, siC, true);      // expected-error {{'__builtin_amdgcn_sdot4' needs target feature dot1-insts}}
 
-  uiOut[2] = __builtin_amdgcn_udot4(uiA, uiB, uiC, false);     // expected-error {{'__builtin_amdgcn_udot4' needs target feature dot-insts}}
-  uiOut[3] = __builtin_amdgcn_udot4(uiA, uiB, uiC, true);      // expected-error {{'__builtin_amdgcn_udot4' needs target feature dot-insts}}
+  uiOut[2] = __builtin_amdgcn_udot4(uiA, uiB, uiC, false);     // expected-error {{'__builtin_amdgcn_udot4' needs target feature dot2-insts}}
+  uiOut[3] = __builtin_amdgcn_udot4(uiA, uiB, uiC, true);      // expected-error {{'__builtin_amdgcn_udot4' needs target feature dot2-insts}}
 
-  siOut[4] = __builtin_amdgcn_sdot8(siA, siB, siC, false);     // expected-error {{'__builtin_amdgcn_sdot8' needs target feature dot-insts}}
-  siOut[5] = __builtin_amdgcn_sdot8(siA, siB, siC, true);      // expected-error {{'__builtin_amdgcn_sdot8' needs target feature dot-insts}}
+  siOut[4] = __builtin_amdgcn_sdot8(siA, siB, siC, false);     // expected-error {{'__builtin_amdgcn_sdot8' needs target feature dot1-insts}}
+  siOut[5] = __builtin_amdgcn_sdot8(siA, siB, siC, true);      // expected-error {{'__builtin_amdgcn_sdot8' needs target feature dot1-insts}}
 
-  uiOut[4] = __builtin_amdgcn_udot8(uiA, uiB, uiC, false);     // expected-error {{'__builtin_amdgcn_udot8' needs target feature dot-insts}}
-  uiOut[5] = __builtin_amdgcn_udot8(uiA, uiB, uiC, true);      // expected-error {{'__builtin_amdgcn_udot8' needs target feature dot-insts}}
+  uiOut[4] = __builtin_amdgcn_udot8(uiA, uiB, uiC, false);     // expected-error {{'__builtin_amdgcn_udot8' needs target feature dot2-insts}}
+  uiOut[5] = __builtin_amdgcn_udot8(uiA, uiB, uiC, true);      // expected-error {{'__builtin_amdgcn_udot8' needs target feature dot2-insts}}
 }