Skip to content

Commit 8e01569

Browse files
ritter-x2aritter-x2a
committed
Revert part of the initial PR to leave the subtarget feature alone.
Also revert the FeatureForceStoreSC0SC1 removal so that it can be handled in a separate PR.
1 parent 149bb69 commit 8e01569

24 files changed

+328
-297
lines changed

‎clang/include/clang/Basic/BuiltinsAMDGPU.def

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -248,13 +248,13 @@ TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fmax_f64, "dd*0d", "t", "gfx90a-inst
248248
TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_f64, "dd*3d", "t", "gfx90a-insts")
249249
TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_f32, "ff*3f", "t", "gfx8-insts")
250250

251-
TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_f32, "ff*0f", "t", "gfx942-insts")
251+
TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_f32, "ff*0f", "t", "gfx940-insts")
252252
TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_v2f16, "V2hV2h*0V2h", "t", "atomic-flat-pk-add-16-insts")
253253
TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_v2bf16, "V2sV2s*0V2s", "t", "atomic-flat-pk-add-16-insts")
254254
TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_v2bf16, "V2sV2s*1V2s", "t", "atomic-global-pk-add-bf16-inst")
255255
TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2bf16, "V2sV2s*3V2s", "t", "atomic-ds-pk-add-16-insts")
256256
TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2f16, "V2hV2h*3V2h", "t", "atomic-ds-pk-add-16-insts")
257-
TARGET_BUILTIN(__builtin_amdgcn_global_load_lds, "vv*1v*3IUiIiIUi", "t", "gfx942-insts")
257+
TARGET_BUILTIN(__builtin_amdgcn_global_load_lds, "vv*1v*3IUiIiIUi", "t", "gfx940-insts")
258258

259259
//===----------------------------------------------------------------------===//
260260
// Deep learning builtins.

‎clang/test/CodeGenCXX/dynamic-cast-address-space.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -112,9 +112,9 @@ const B& f(A *a) {
112112
// CHECK: attributes #[[ATTR3]] = { nounwind }
113113
// CHECK: attributes #[[ATTR4]] = { noreturn }
114114
//.
115-
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR0]] = { mustprogress noinline optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+bf8-cvt-scale-insts,+bitop3-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx942-insts,+gfx950-insts,+gws,+image-insts,+mai-insts,+permlane16-swap,+permlane32-swap,+prng-inst,+s-memrealtime,+s-memtime-inst,+wavefrontsize32,+wavefrontsize64" }
115+
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR0]] = { mustprogress noinline optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+bf8-cvt-scale-insts,+bitop3-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+mai-insts,+permlane16-swap,+permlane32-swap,+prng-inst,+s-memrealtime,+s-memtime-inst,+wavefrontsize32,+wavefrontsize64" }
116116
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR1:[0-9]+]] = { nounwind willreturn memory(read) }
117-
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR2:[0-9]+]] = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+bf8-cvt-scale-insts,+bitop3-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx942-insts,+gfx950-insts,+gws,+image-insts,+mai-insts,+permlane16-swap,+permlane32-swap,+prng-inst,+s-memrealtime,+s-memtime-inst,+wavefrontsize32,+wavefrontsize64" }
117+
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR2:[0-9]+]] = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+bf8-cvt-scale-insts,+bitop3-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+mai-insts,+permlane16-swap,+permlane32-swap,+prng-inst,+s-memrealtime,+s-memtime-inst,+wavefrontsize32,+wavefrontsize64" }
118118
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR3]] = { nounwind }
119119
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR4]] = { noreturn }
120120
//.

‎clang/test/CodeGenOpenCL/amdgpu-features.cl

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -83,9 +83,9 @@
8383
// GFX909: "target-features"="+16-bit-insts,+ci-insts,+dpp,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
8484
// GFX90A: "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
8585
// GFX90C: "target-features"="+16-bit-insts,+ci-insts,+dpp,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
86-
// GFX942: "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+fp8-conversion-insts,+fp8-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx942-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64,+xf32-insts"
87-
// GFX9_4_Generic: "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx942-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
88-
// GFX950: "target-features"="+16-bit-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+bf8-cvt-scale-insts,+bitop3-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx942-insts,+gfx950-insts,+mai-insts,+permlane16-swap,+permlane32-swap,+prng-inst,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
86+
// GFX942: "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+fp8-conversion-insts,+fp8-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64,+xf32-insts"
87+
// GFX9_4_Generic: "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
88+
// GFX950: "target-features"="+16-bit-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+bf8-cvt-scale-insts,+bitop3-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+mai-insts,+permlane16-swap,+permlane32-swap,+prng-inst,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
8989
// GFX1010: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"
9090
// GFX1011: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"
9191
// GFX1012: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"

‎clang/test/CodeGenOpenCL/builtins-amdgcn-fp-atomics-gfx90a-err.cl

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -9,7 +9,7 @@ typedef short __attribute__((ext_vector_type(2))) short2;
99
void test_atomic_fadd(__global half2 *addrh2, __local half2 *addrh2l, half2 xh2,
1010
__global short2 *addrs2, __local short2 *addrs2l, short2 xs2,
1111
__global float *addrf, float xf) {
12-
__builtin_amdgcn_flat_atomic_fadd_f32(addrf, xf); // expected-error{{'__builtin_amdgcn_flat_atomic_fadd_f32' needs target feature gfx942-insts}}
12+
__builtin_amdgcn_flat_atomic_fadd_f32(addrf, xf); // expected-error{{'__builtin_amdgcn_flat_atomic_fadd_f32' needs target feature gfx940-insts}}
1313
__builtin_amdgcn_flat_atomic_fadd_v2f16(addrh2, xh2); // expected-error{{'__builtin_amdgcn_flat_atomic_fadd_v2f16' needs target feature atomic-flat-pk-add-16-insts}}
1414
__builtin_amdgcn_flat_atomic_fadd_v2bf16(addrs2, xs2); // expected-error{{'__builtin_amdgcn_flat_atomic_fadd_v2bf16' needs target feature atomic-flat-pk-add-16-insts}}
1515
__builtin_amdgcn_global_atomic_fadd_v2bf16(addrs2, xs2); // expected-error{{'__builtin_amdgcn_global_atomic_fadd_v2bf16' needs target feature atomic-global-pk-add-bf16-inst}}

‎llvm/lib/Target/AMDGPU/AMDGPU.td

Lines changed: 26 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -372,10 +372,10 @@ def FeatureGFX90AInsts : SubtargetFeature<"gfx90a-insts",
372372
// [HasAtomicFMinFMaxF64GlobalInsts, HasAtomicFMinFMaxF64FlatInsts] // TODO
373373
>;
374374

375-
def FeatureGFX942Insts : SubtargetFeature<"gfx942-insts",
376-
"GFX942Insts",
375+
def FeatureGFX940Insts : SubtargetFeature<"gfx940-insts",
376+
"GFX940Insts",
377377
"true",
378-
"Additional instructions for GFX942+"
378+
"Additional instructions for GFX940+"
379379
>;
380380

381381
def FeaturePermlane16Swap : SubtargetFeature<"permlane16-swap",
@@ -1040,6 +1040,12 @@ def FeatureVALUTransUseHazard : SubtargetFeature<"valu-trans-use-hazard",
10401040
"Hazard when TRANS instructions are closely followed by a use of the result"
10411041
>;
10421042

1043+
def FeatureForceStoreSC0SC1 : SubtargetFeature<"force-store-sc0-sc1",
1044+
"HasForceStoreSC0SC1",
1045+
"true",
1046+
"Has SC0 and SC1 on stores"
1047+
>;
1048+
10431049
def FeatureSALUFloatInsts : SubtargetFeature<"salu-float",
10441050
"HasSALUFloatInsts",
10451051
"true",
@@ -1558,7 +1564,7 @@ def FeatureISAVersion9_0_C : FeatureSet<
15581564
def FeatureISAVersion9_4_Common : FeatureSet<
15591565
[FeatureGFX9,
15601566
FeatureGFX90AInsts,
1561-
FeatureGFX942Insts,
1567+
FeatureGFX940Insts,
15621568
FeatureFmaMixInsts,
15631569
FeatureLDSBankCount32,
15641570
FeatureDLInsts,
@@ -2047,20 +2053,20 @@ def isGFX8GFX9NotGFX90A :
20472053
AssemblerPredicate<(all_of FeatureGFX8Insts, FeatureGCN3Encoding, (not FeatureGFX90AInsts))>;
20482054

20492055
def isGFX90AOnly :
2050-
Predicate<"Subtarget->hasGFX90AInsts() && !Subtarget->hasGFX942Insts()">,
2051-
AssemblerPredicate<(all_of FeatureGFX90AInsts, (not FeatureGFX942Insts))>;
2056+
Predicate<"Subtarget->hasGFX90AInsts() && !Subtarget->hasGFX940Insts()">,
2057+
AssemblerPredicate<(all_of FeatureGFX90AInsts, (not FeatureGFX940Insts))>;
20522058

20532059
def isGFX908orGFX90A :
2054-
Predicate<"Subtarget->hasMAIInsts() && !Subtarget->hasGFX942Insts()">,
2055-
AssemblerPredicate<(all_of FeatureMAIInsts, (not FeatureGFX942Insts))>;
2060+
Predicate<"Subtarget->hasMAIInsts() && !Subtarget->hasGFX940Insts()">,
2061+
AssemblerPredicate<(all_of FeatureMAIInsts, (not FeatureGFX940Insts))>;
20562062

2057-
def isGFX942Plus :
2058-
Predicate<"Subtarget->hasGFX942Insts()">,
2059-
AssemblerPredicate<(all_of FeatureGFX942Insts)>;
2063+
def isGFX940Plus :
2064+
Predicate<"Subtarget->hasGFX940Insts()">,
2065+
AssemblerPredicate<(all_of FeatureGFX940Insts)>;
20602066

2061-
def isNotGFX942Plus :
2062-
Predicate<"!Subtarget->hasGFX942Insts()">,
2063-
AssemblerPredicate<(all_of (not FeatureGFX942Insts))>;
2067+
def isNotGFX940Plus :
2068+
Predicate<"!Subtarget->hasGFX940Insts()">,
2069+
AssemblerPredicate<(all_of (not FeatureGFX940Insts))>;
20642070

20652071
def HasGFX950Insts :
20662072
Predicate<"Subtarget->hasGFX950Insts()">,
@@ -2074,11 +2080,11 @@ def HasPermlane32Swap :
20742080
Predicate<"Subtarget->hasPermlane32Swap()">,
20752081
AssemblerPredicate<(all_of FeaturePermlane32Swap)>;
20762082

2077-
def isGFX8GFX9NotGFX942 :
2078-
Predicate<"!Subtarget->hasGFX942Insts() &&"
2083+
def isGFX8GFX9NotGFX940 :
2084+
Predicate<"!Subtarget->hasGFX940Insts() &&"
20792085
"(Subtarget->getGeneration() == AMDGPUSubtarget::VOLCANIC_ISLANDS ||"
20802086
" Subtarget->getGeneration() == AMDGPUSubtarget::GFX9)">,
2081-
AssemblerPredicate<(all_of FeatureGFX8Insts, FeatureGCN3Encoding, (not FeatureGFX942Insts))>;
2087+
AssemblerPredicate<(all_of FeatureGFX8Insts, FeatureGCN3Encoding, (not FeatureGFX940Insts))>;
20822088

20832089
def isGFX8GFX9 :
20842090
Predicate<"Subtarget->getGeneration() == AMDGPUSubtarget::VOLCANIC_ISLANDS ||"
@@ -2185,9 +2191,9 @@ def HasD16LoadStore : Predicate<"Subtarget->hasD16LoadStore()">,
21852191
AssemblerPredicate<(all_of FeatureGFX9Insts)>;
21862192

21872193
def HasFlatScratchSTMode : Predicate<"Subtarget->hasFlatScratchSTMode()">,
2188-
AssemblerPredicate<(any_of FeatureGFX10_3Insts, FeatureGFX942Insts)>;
2194+
AssemblerPredicate<(any_of FeatureGFX10_3Insts, FeatureGFX940Insts)>;
21892195
def HasFlatScratchSVSMode : Predicate<"Subtarget->hasFlatScratchSVSMode()">,
2190-
AssemblerPredicate<(any_of FeatureGFX942Insts, FeatureGFX11Insts)>;
2196+
AssemblerPredicate<(any_of FeatureGFX940Insts, FeatureGFX11Insts)>;
21912197

21922198
def HasGFX10_AEncoding : Predicate<"Subtarget->hasGFX10_AEncoding()">,
21932199
AssemblerPredicate<(all_of FeatureGFX10_AEncoding)>;
@@ -2295,7 +2301,7 @@ def HasPkMovB32 : Predicate<"Subtarget->hasPkMovB32()">,
22952301

22962302
def HasFmaakFmamkF32Insts :
22972303
Predicate<"Subtarget->hasFmaakFmamkF32Insts()">,
2298-
AssemblerPredicate<(any_of FeatureGFX10Insts, FeatureGFX942Insts)>;
2304+
AssemblerPredicate<(any_of FeatureGFX10Insts, FeatureGFX940Insts)>;
22992305

23002306
def HasImageInsts : Predicate<"Subtarget->hasImageInsts()">,
23012307
AssemblerPredicate<(all_of FeatureImageInsts)>;

‎llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp

Lines changed: 15 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -1492,12 +1492,14 @@ class AMDGPUAsmParser : public MCTargetAsmParser {
14921492
return AMDGPU::isGFX9(getSTI());
14931493
}
14941494

1495-
// TODO: isGFX90A is also true for GFX942. We need to clean it.
1495+
// TODO: isGFX90A is also true for GFX940. We need to clean it.
14961496
bool isGFX90A() const {
14971497
return AMDGPU::isGFX90A(getSTI());
14981498
}
14991499

1500-
bool isGFX942() const { return AMDGPU::isGFX942(getSTI()); }
1500+
bool isGFX940() const {
1501+
return AMDGPU::isGFX940(getSTI());
1502+
}
15011503

15021504
bool isGFX9Plus() const {
15031505
return AMDGPU::isGFX9Plus(getSTI());
@@ -4630,7 +4632,7 @@ bool AMDGPUAsmParser::validateOpSel(const MCInst &Inst) {
46304632

46314633
uint64_t TSFlags = MII.get(Opc).TSFlags;
46324634

4633-
if (isGFX942() && (TSFlags & SIInstrFlags::IsDOT)) {
4635+
if (isGFX940() && (TSFlags & SIInstrFlags::IsDOT)) {
46344636
int OpSelIdx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::op_sel);
46354637
if (OpSelIdx != -1) {
46364638
if (Inst.getOperand(OpSelIdx).getImm() != 0)
@@ -4938,12 +4940,12 @@ bool AMDGPUAsmParser::validateBLGP(const MCInst &Inst,
49384940
bool IsNeg = StringRef(BLGPLoc.getPointer()).starts_with("neg:");
49394941
auto FB = getFeatureBits();
49404942
bool UsesNeg = false;
4941-
if (FB[AMDGPU::FeatureGFX942Insts]) {
4943+
if (FB[AMDGPU::FeatureGFX940Insts]) {
49424944
switch (Opc) {
4943-
case AMDGPU::V_MFMA_F64_16X16X4F64_gfx942_acd:
4944-
case AMDGPU::V_MFMA_F64_16X16X4F64_gfx942_vcd:
4945-
case AMDGPU::V_MFMA_F64_4X4X4F64_gfx942_acd:
4946-
case AMDGPU::V_MFMA_F64_4X4X4F64_gfx942_vcd:
4945+
case AMDGPU::V_MFMA_F64_16X16X4F64_gfx940_acd:
4946+
case AMDGPU::V_MFMA_F64_16X16X4F64_gfx940_vcd:
4947+
case AMDGPU::V_MFMA_F64_4X4X4F64_gfx940_acd:
4948+
case AMDGPU::V_MFMA_F64_4X4X4F64_gfx940_vcd:
49474949
UsesNeg = true;
49484950
}
49494951
}
@@ -5058,7 +5060,7 @@ bool AMDGPUAsmParser::validateCoherencyBits(const MCInst &Inst,
50585060
}
50595061
}
50605062

5061-
if (isGFX90A() && !isGFX942() && (CPol & CPol::SCC)) {
5063+
if (isGFX90A() && !isGFX940() && (CPol & CPol::SCC)) {
50625064
const uint64_t AllowSCCModifier = SIInstrFlags::MUBUF |
50635065
SIInstrFlags::MTBUF | SIInstrFlags::MIMG |
50645066
SIInstrFlags::FLAT;
@@ -5077,7 +5079,7 @@ bool AMDGPUAsmParser::validateCoherencyBits(const MCInst &Inst,
50775079

50785080
if (TSFlags & SIInstrFlags::IsAtomicRet) {
50795081
if (!(TSFlags & SIInstrFlags::MIMG) && !(CPol & CPol::GLC)) {
5080-
Error(IDLoc, isGFX942() ? "instruction must use sc0"
5082+
Error(IDLoc, isGFX940() ? "instruction must use sc0"
50815083
: "instruction must use glc");
50825084
return false;
50835085
}
@@ -5086,8 +5088,8 @@ bool AMDGPUAsmParser::validateCoherencyBits(const MCInst &Inst,
50865088
SMLoc S = getImmLoc(AMDGPUOperand::ImmTyCPol, Operands);
50875089
StringRef CStr(S.getPointer());
50885090
S = SMLoc::getFromPointer(
5089-
&CStr.data()[CStr.find(isGFX942() ? "sc0" : "glc")]);
5090-
Error(S, isGFX942() ? "instruction must not use sc0"
5091+
&CStr.data()[CStr.find(isGFX940() ? "sc0" : "glc")]);
5092+
Error(S, isGFX940() ? "instruction must not use sc0"
50915093
: "instruction must not use glc");
50925094
return false;
50935095
}
@@ -6641,7 +6643,7 @@ unsigned AMDGPUAsmParser::getCPolKind(StringRef Id, StringRef Mnemo,
66416643
bool &Disabling) const {
66426644
Disabling = Id.consume_front("no");
66436645

6644-
if (isGFX942() && !Mnemo.starts_with("s_")) {
6646+
if (isGFX940() && !Mnemo.starts_with("s_")) {
66456647
return StringSwitch<unsigned>(Id)
66466648
.Case("nt", AMDGPU::CPol::NT)
66476649
.Case("sc0", AMDGPU::CPol::SC0)

0 commit comments

Comments
 (0)