Skip to content

Commit 80972ce

Browse files
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 67f9b75 commit 80972ce

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());
@@ -4631,7 +4633,7 @@ bool AMDGPUAsmParser::validateOpSel(const MCInst &Inst) {
46314633

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

4634-
if (isGFX942() && (TSFlags & SIInstrFlags::IsDOT)) {
4636+
if (isGFX940() && (TSFlags & SIInstrFlags::IsDOT)) {
46354637
int OpSelIdx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::op_sel);
46364638
if (OpSelIdx != -1) {
46374639
if (Inst.getOperand(OpSelIdx).getImm() != 0)
@@ -4940,12 +4942,12 @@ bool AMDGPUAsmParser::validateBLGP(const MCInst &Inst,
49404942
bool IsNeg = StringRef(BLGPLoc.getPointer()).starts_with("neg:");
49414943
auto FB = getFeatureBits();
49424944
bool UsesNeg = false;
4943-
if (FB[AMDGPU::FeatureGFX942Insts]) {
4945+
if (FB[AMDGPU::FeatureGFX940Insts]) {
49444946
switch (Opc) {
4945-
case AMDGPU::V_MFMA_F64_16X16X4F64_gfx942_acd:
4946-
case AMDGPU::V_MFMA_F64_16X16X4F64_gfx942_vcd:
4947-
case AMDGPU::V_MFMA_F64_4X4X4F64_gfx942_acd:
4948-
case AMDGPU::V_MFMA_F64_4X4X4F64_gfx942_vcd:
4947+
case AMDGPU::V_MFMA_F64_16X16X4F64_gfx940_acd:
4948+
case AMDGPU::V_MFMA_F64_16X16X4F64_gfx940_vcd:
4949+
case AMDGPU::V_MFMA_F64_4X4X4F64_gfx940_acd:
4950+
case AMDGPU::V_MFMA_F64_4X4X4F64_gfx940_vcd:
49494951
UsesNeg = true;
49504952
}
49514953
}
@@ -5060,7 +5062,7 @@ bool AMDGPUAsmParser::validateCoherencyBits(const MCInst &Inst,
50605062
}
50615063
}
50625064

5063-
if (isGFX90A() && !isGFX942() && (CPol & CPol::SCC)) {
5065+
if (isGFX90A() && !isGFX940() && (CPol & CPol::SCC)) {
50645066
const uint64_t AllowSCCModifier = SIInstrFlags::MUBUF |
50655067
SIInstrFlags::MTBUF | SIInstrFlags::MIMG |
50665068
SIInstrFlags::FLAT;
@@ -5079,7 +5081,7 @@ bool AMDGPUAsmParser::validateCoherencyBits(const MCInst &Inst,
50795081

50805082
if (TSFlags & SIInstrFlags::IsAtomicRet) {
50815083
if (!(TSFlags & SIInstrFlags::MIMG) && !(CPol & CPol::GLC)) {
5082-
Error(IDLoc, isGFX942() ? "instruction must use sc0"
5084+
Error(IDLoc, isGFX940() ? "instruction must use sc0"
50835085
: "instruction must use glc");
50845086
return false;
50855087
}
@@ -5088,8 +5090,8 @@ bool AMDGPUAsmParser::validateCoherencyBits(const MCInst &Inst,
50885090
SMLoc S = getImmLoc(AMDGPUOperand::ImmTyCPol, Operands);
50895091
StringRef CStr(S.getPointer());
50905092
S = SMLoc::getFromPointer(
5091-
&CStr.data()[CStr.find(isGFX942() ? "sc0" : "glc")]);
5092-
Error(S, isGFX942() ? "instruction must not use sc0"
5093+
&CStr.data()[CStr.find(isGFX940() ? "sc0" : "glc")]);
5094+
Error(S, isGFX940() ? "instruction must not use sc0"
50935095
: "instruction must not use glc");
50945096
return false;
50955097
}
@@ -6655,7 +6657,7 @@ unsigned AMDGPUAsmParser::getCPolKind(StringRef Id, StringRef Mnemo,
66556657
bool &Disabling) const {
66566658
Disabling = Id.consume_front("no");
66576659

6658-
if (isGFX942() && !Mnemo.starts_with("s_")) {
6660+
if (isGFX940() && !Mnemo.starts_with("s_")) {
66596661
return StringSwitch<unsigned>(Id)
66606662
.Case("nt", AMDGPU::CPol::NT)
66616663
.Case("sc0", AMDGPU::CPol::SC0)

0 commit comments

Comments
 (0)