Merge branch 'main' into cmake-msvc-no-shortfile

Author: RKSimon

clang/include/clang/Basic/BuiltinsAMDGPU.def

@@ -651,6 +651,16 @@ TARGET_BUILTIN(__builtin_amdgcn_global_load_monitor_b128, "V4iV4i*1Ii", "nc", "g
 TARGET_BUILTIN(__builtin_amdgcn_flat_load_monitor_b32, "ii*0Ii", "nc", "gfx1250-insts")
 TARGET_BUILTIN(__builtin_amdgcn_flat_load_monitor_b64, "V2iV2i*0Ii", "nc", "gfx1250-insts")
 TARGET_BUILTIN(__builtin_amdgcn_flat_load_monitor_b128, "V4iV4i*0Ii", "nc", "gfx1250-insts")
+TARGET_BUILTIN(__builtin_amdgcn_global_load_async_to_lds_b8, "vc*1c*3IiIi", "nc", "gfx1250-insts")
+TARGET_BUILTIN(__builtin_amdgcn_global_load_async_to_lds_b32, "vi*1i*3IiIi", "nc", "gfx1250-insts")
+TARGET_BUILTIN(__builtin_amdgcn_global_load_async_to_lds_b64, "vV2i*1V2i*3IiIi", "nc", "gfx1250-insts")
+TARGET_BUILTIN(__builtin_amdgcn_global_load_async_to_lds_b128, "vV4i*1V4i*3IiIi", "nc", "gfx1250-insts")
+TARGET_BUILTIN(__builtin_amdgcn_global_store_async_from_lds_b8, "vc*1c*3IiIi", "nc", "gfx1250-insts")
+TARGET_BUILTIN(__builtin_amdgcn_global_store_async_from_lds_b32, "vi*1i*3IiIi", "nc", "gfx1250-insts")
+TARGET_BUILTIN(__builtin_amdgcn_global_store_async_from_lds_b64, "vV2i*1V2i*3IiIi", "nc", "gfx1250-insts")
+TARGET_BUILTIN(__builtin_amdgcn_global_store_async_from_lds_b128, "vV4i*1V4i*3IiIi", "nc", "gfx1250-insts")
+TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_async_barrier_arrive_b64, "vLi*3", "nc", "gfx1250-insts")
+TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_barrier_arrive_rtn_b64, "LiLi*3Li", "nc", "gfx1250-insts")
 
 TARGET_BUILTIN(__builtin_amdgcn_tensor_load_to_lds, "vV4iV8iV4iV4iIi", "nc", "gfx1250-insts")
 TARGET_BUILTIN(__builtin_amdgcn_tensor_load_to_lds_d2, "vV4iV8iIi", "nc", "gfx1250-insts")
@@ -670,9 +680,6 @@ TARGET_BUILTIN(__builtin_amdgcn_ds_load_tr16_b128_v8i16, "V8sV8s*3", "nc", "gfx1
 TARGET_BUILTIN(__builtin_amdgcn_ds_load_tr16_b128_v8f16, "V8hV8h*3", "nc", "gfx1250-insts,wavefrontsize32")
 TARGET_BUILTIN(__builtin_amdgcn_ds_load_tr16_b128_v8bf16, "V8yV8y*3", "nc", "gfx1250-insts,wavefrontsize32")
 
-TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_async_barrier_arrive_b64, "vLi*3", "nc", "gfx1250-insts")
-TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_barrier_arrive_rtn_b64, "LiLi*3Li", "nc", "gfx1250-insts")
-
 TARGET_BUILTIN(__builtin_amdgcn_s_setprio_inc_wg, "vIs", "n", "setprio-inc-wg-inst")
 TARGET_BUILTIN(__builtin_amdgcn_s_monitor_sleep,  "vIs", "n", "gfx1250-insts")
 

clang/include/clang/Basic/FileManager.h

@@ -237,11 +237,6 @@ class FileManager : public RefCountedBase<FileManager> {
   FileEntryRef getVirtualFileRef(StringRef Filename, off_t Size,
                                  time_t ModificationTime);
 
-  LLVM_DEPRECATED("Functions returning FileEntry are deprecated.",
-                  "getVirtualFileRef()")
-  const FileEntry *getVirtualFile(StringRef Filename, off_t Size,
-                                  time_t ModificationTime);
-
   /// Retrieve a FileEntry that bypasses VFE, which is expected to be a virtual
   /// file entry, to access the real file.  The returned FileEntry will have
   /// the same filename as FE but a different identity and its own stat.

clang/include/clang/CIR/MissingFeatures.h

@@ -224,7 +224,6 @@ struct MissingFeatures {
   static bool moduleNameHash() { return false; }
   static bool msabi() { return false; }
   static bool needsGlobalCtorDtor() { return false; }
-  static bool nonFineGrainedBitfields() { return false; }
   static bool objCBlocks() { return false; }
   static bool objCGC() { return false; }
   static bool objCLifetime() { return false; }

clang/lib/Basic/FileManager.cpp

@@ -368,11 +368,6 @@ void FileManager::trackVFSUsage(bool Active) {
   });
 }
 
-const FileEntry *FileManager::getVirtualFile(StringRef Filename, off_t Size,
-                                             time_t ModificationTime) {
-  return &getVirtualFileRef(Filename, Size, ModificationTime).getFileEntry();
-}
-
 FileEntryRef FileManager::getVirtualFileRef(StringRef Filename, off_t Size,
                                             time_t ModificationTime) {
   ++NumFileLookups;

clang/lib/Basic/Targets/AMDGPU.cpp

@@ -266,8 +266,11 @@ AMDGPUTargetInfo::AMDGPUTargetInfo(const llvm::Triple &Triple,
 
   MaxAtomicPromoteWidth = MaxAtomicInlineWidth = 64;
   CUMode = !(GPUFeatures & llvm::AMDGPU::FEATURE_WGP);
-  for (auto F : {"image-insts", "gws", "vmem-to-lds-load-insts"})
-    ReadOnlyFeatures.insert(F);
+
+  for (auto F : {"image-insts", "gws", "vmem-to-lds-load-insts"}) {
+    if (GPUKind != llvm::AMDGPU::GK_NONE)
+      ReadOnlyFeatures.insert(F);
+  }
   HalfArgsAndReturns = true;
 }
 

clang/lib/CIR/CodeGen/CIRGenRecordLayoutBuilder.cpp

@@ -438,9 +438,7 @@ CIRRecordLowering::accumulateBitFields(RecordDecl::field_iterator field,
           } else if (cirGenTypes.getCGModule()
                          .getCodeGenOpts()
                          .FineGrainedBitfieldAccesses) {
-            assert(!cir::MissingFeatures::nonFineGrainedBitfields());
-            cirGenTypes.getCGModule().errorNYI(field->getSourceRange(),
-                                               "NYI FineGrainedBitfield");
+            installBest = true;
           } else {
             // Otherwise, we're not installing. Update the bit size
             // of the current span to go all the way to limitOffset, which is

clang/test/CIR/CodeGen/finegrain-bitfield-access.cpp

@@ -0,0 +1,271 @@
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -emit-cir -ffine-grained-bitfield-accesses %s -o %t.cir
+// RUN: FileCheck --input-file=%t.cir %s --check-prefix=CIR
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -emit-llvm -ffine-grained-bitfield-accesses %s -o %t-cir.ll
+// RUN: FileCheck --input-file=%t-cir.ll %s --check-prefix=LLVM
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -ffine-grained-bitfield-accesses %s -o %t.ll
+// RUN: FileCheck --input-file=%t.ll %s --check-prefix=OGCG
+
+struct S1 {
+  unsigned f1:2;
+  unsigned f2:6;
+  unsigned f3:8;
+  unsigned f4:4;
+  unsigned f5:8;
+};
+
+// CIR-DAG: !rec_S1 = !cir.record<struct "S1" {!u8i, !u8i, !u16i}>
+// LLVM-DAG: %struct.S1 = type { i8, i8, i16 }
+// OGCG-DAG: %struct.S1 = type { i8, i8, i16 }
+
+struct S2 {
+  unsigned long f1:16;
+  unsigned long f2:16;
+  unsigned long f3:6;
+};
+
+// CIR-DAG: !rec_S2 = !cir.record<struct "S2" padded {!u16i, !u16i, !u8i, !cir.array<!u8i x 3>}>
+// LLVM-DAG: %struct.S2 = type { i16, i16, i8, [3 x i8] }
+// OGCG-DAG: %struct.S2 = type { i16, i16, i8, [3 x i8] }
+
+struct S3 {
+  unsigned long f1:14;
+  unsigned long f2:18;
+  unsigned long f3:32;
+};
+
+// CIR-DAG: !rec_S3 = !cir.record<struct "S3" {!u32i, !u32i}>
+// LLVM-DAG: %struct.S3 = type { i32, i32 }
+// OGCG-DAG: %struct.S3 = type { i32, i32 }
+
+S1 a1;
+S2 a2;
+S3 a3;
+
+unsigned read8_1() {
+  return a1.f3;
+}
+
+// CIR-LABEL: @_Z7read8_1v
+// CIR: [[MEMBER:%.*]] = cir.get_member %1[1] {name = "f3"} : !cir.ptr<!rec_S1> -> !cir.ptr<!u8i>
+// CIR: [[BITFI:%.*]] = cir.get_bitfield align(1) (#bfi_f3, [[MEMBER]] : !cir.ptr<!u8i>) -> !u32i
+// CIR: cir.store [[BITFI]], {{.*}} : !u32i, !cir.ptr<!u32i>
+// CIR: [[RET:%.*]] = cir.load {{.*}} : !cir.ptr<!u32i>, !u32i
+// CIR: cir.return [[RET]] : !u32i
+
+// LLVM-LABEL: @_Z7read8_1v
+// LLVM:  [[MEMBER:%.*]] = load i8, ptr getelementptr inbounds nuw (i8, ptr {{.*}}, i64 1), align 1
+// LLVM:  [[BFCAST:%.*]] = zext i8 [[MEMBER]] to i32
+// LLVM:  store i32 [[BFCAST]], ptr {{.*}}, align 4
+// LLVM:  [[RET:%.*]] = load i32, ptr {{.*}}, align 4
+// LLVM:  ret i32 [[RET]]
+
+// OGCG-LABEL: @_Z7read8_1v
+// OGCG: [[BFLOAD:%.*]] = load i8, ptr getelementptr inbounds nuw (%struct.S1, ptr {{.*}}, i32 0, i32 1), align 1
+// OGCG-NEXT: [[BFCAST:%.*]] = zext i8 [[BFLOAD]] to i32
+// OGCG-NEXT: ret i32 [[BFCAST]]
+
+void write8_1() {
+  a1.f3 = 3;
+}
+
+// CIR-LABEL: @_Z8write8_1v
+// CIR: [[CONST3:%.*]] = cir.const #cir.int<3> : !s32i
+// CIR: [[INT3:%.*]] = cir.cast(integral, [[CONST3]] : !s32i), !u32i
+// CIR: [[MEMBER:%.*]] = cir.get_member {{.*}}[1] {name = "f3"} : !cir.ptr<!rec_S1> -> !cir.ptr<!u8i>
+// CIR: cir.set_bitfield align(1) (#bfi_f3, [[MEMBER]] : !cir.ptr<!u8i>, [[INT3]] : !u32i) -> !u32i
+
+// LLVM-LABEL: @_Z8write8_1v
+// LLVM:  store i8 3, ptr getelementptr inbounds nuw (i8, ptr {{.*}}, i64 1), align 1
+// LLVM:  ret void
+
+// OGCG-LABEL: @_Z8write8_1v
+// OGCG: store i8 3, ptr getelementptr inbounds nuw (%struct.S1, ptr {{.*}}, i32 0, i32 1), align 1
+// OGCG-NEXT: ret void
+
+unsigned read8_2() {
+
+  return a1.f5;
+}
+
+// CIR-LABEL: @_Z7read8_2v
+// CIR: [[MEMBER:%.*]] = cir.get_member {{.*}}[2] {name = "f5"} : !cir.ptr<!rec_S1> -> !cir.ptr<!u16i>
+// CIR: [[BITFI:%.*]] = cir.get_bitfield align(2) (#bfi_f5, [[MEMBER]] : !cir.ptr<!u16i>) -> !u32i
+// CIR: cir.store [[BITFI]], {{.*}} : !u32i, !cir.ptr<!u32i>
+// CIR: [[RET:%.*]] = cir.load {{.*}} : !cir.ptr<!u32i>, !u32i
+// CIR: cir.return [[RET]] : !u32i
+
+// LLVM-LABEL: @_Z7read8_2v
+// LLVM:  [[BFLOAD:%.*]] = load i16, ptr getelementptr inbounds nuw (i8, ptr {{.*}}, i64 2), align 2
+// LLVM:  [[BFLSHR:%.*]] = lshr i16 [[BFLOAD]], 4
+// LLVM:  [[BFCLEAR:%.*]] = and i16 [[BFLSHR]], 255
+// LLVM:  [[BFCAST:%.*]] = zext i16 [[BFCLEAR]] to i32
+// LLVM:  store i32 [[BFCAST]], ptr {{.*}}, align 4
+// LLVM:  [[RET:%.*]] = load i32, ptr {{.*}}, align 4
+// LLVM:  ret i32 [[RET]]
+
+// OGCG-LABEL: @_Z7read8_2v
+// OGCG: [[BFLOAD:%.*]] = load i16, ptr getelementptr inbounds nuw (%struct.S1, ptr {{.*}}, i32 0, i32 2), align 2
+// OGCG-NEXT: [[BFLSHR:%.*]] = lshr i16 [[BFLOAD]], 4
+// OGCG-NEXT: [[BFCLEAR:%.*]] = and i16 [[BFLSHR]], 255
+// OGCG-NEXT: [[BFCAST:%.*]] = zext i16 [[BFCLEAR]] to i32
+// OGCG-NEXT: ret i32 [[BFCAST]]
+
+void write8_2() {
+  a1.f5 = 3;
+}
+
+// CIR-LABEL: @_Z8write8_2v
+// CIR: [[CONST3:%.*]] = cir.const #cir.int<3> : !s32i
+// CIR: [[INT3:%.*]] = cir.cast(integral, [[CONST3]] : !s32i), !u32i
+// CIR: [[MEMBER:%.*]] = cir.get_member {{.*}}[2] {name = "f5"} : !cir.ptr<!rec_S1> -> !cir.ptr<!u16i>
+// CIR: cir.set_bitfield align(2) (#bfi_f5, %3 : !cir.ptr<!u16i>, {{.*}} : !u32i) -> !u32i
+
+// LLVM-LABEL: @_Z8write8_2v
+// LLVM:  [[BFLOAD:%.*]] = load i16, ptr getelementptr inbounds nuw (i8, ptr {{.*}}, i64 2), align 2
+// LLVM:  [[BFCLEAR:%.*]] = and i16 [[BFLOAD]], -4081
+// LLVM:  [[BFSET:%.*]] = or i16 [[BFCLEAR]], 48
+// LLVM:  store i16 [[BFSET]], ptr getelementptr inbounds nuw (i8, ptr {{.*}}, i64 2), align 2
+// LLVM:  ret void
+
+// OGCG-LABEL: @_Z8write8_2v
+// OGCG: [[BFLOAD:%.*]] = load i16, ptr getelementptr inbounds nuw (%struct.S1, ptr {{.*}}, i32 0, i32 2), align 2
+// OGCG-NEXT: [[BFCLEAR:%.*]] = and i16 [[BFLOAD]], -4081
+// OGCG-NEXT: [[BFSET:%.*]] = or i16 [[BFCLEAR]], 48
+// OGCG-NEXT: store i16 [[BFSET]], ptr getelementptr inbounds nuw (%struct.S1, ptr {{.*}}, i32 0, i32 2), align 2
+// OGCG-NEXT: ret void
+
+unsigned read16_1() {
+  return a2.f1;
+}
+
+// CIR-LABEL: @_Z8read16_1v
+// CIR: [[MEMBER:%.*]] = cir.get_member {{.*}}[0] {name = "f1"} : !cir.ptr<!rec_S2> -> !cir.ptr<!u16i>
+// CIR: [[BITFI:%.*]] = cir.get_bitfield align(8) (#bfi_f1, [[MEMBER]] : !cir.ptr<!u16i>) -> !u64i
+// CIR: [[BFCAST:%.*]] = cir.cast(integral, [[BITFI]] : !u64i), !u32i
+// CIR: cir.store [[BFCAST]], {{.*}} : !u32i, !cir.ptr<!u32i>
+// CIR: [[RET:%.*]] = cir.load {{.*}} : !cir.ptr<!u32i>, !u32i
+// CIR: cir.return [[RET]] : !u32i
+
+// LLVM-LABEL: @_Z8read16_1v
+// LLVM:  [[BFLOAD:%.*]] = load i16, ptr {{.*}}, align 8
+// LLVM:  [[BFCAST:%.*]] = zext i16 [[BFLOAD]] to i64
+// LLVM:  [[BF:%.*]] = trunc i64 [[BFCAST]] to i32
+// LLVM:  store i32 [[BF]], ptr {{.*}}, align 4
+// LLVM:  [[RET:%.*]] = load i32, ptr {{.*}}, align 4
+// LLVM:  ret i32 [[RET]]
+
+// OGCG-LABEL: @_Z8read16_1v
+// OGCG: [[BFLOAD:%.*]] = load i16, ptr {{.*}}, align 8
+// OGCG-NEXT: [[BFCAST:%.*]] = zext i16 [[BFLOAD]] to i64
+// OGCG-NEXT: [[RET:%.*]] = trunc i64 [[BFCAST]] to i32
+// OGCG-NEXT: ret i32 [[RET]]
+
+unsigned read16_2() {
+  return a2.f2;
+}
+
+// CIR-LABEL: @_Z8read16_2v
+// CIR: [[MEMBER:%.*]] = cir.get_member {{.*}}[1] {name = "f2"} : !cir.ptr<!rec_S2> -> !cir.ptr<!u16i>
+// CIR: [[BITFI:%.*]] = cir.get_bitfield align(2) (#bfi_f2, [[MEMBER]] : !cir.ptr<!u16i>) -> !u64i
+// CIR: [[BFCAST:%.*]] = cir.cast(integral, [[BITFI]] : !u64i), !u32i
+// CIR: cir.store [[BFCAST]], {{.*}} : !u32i, !cir.ptr<!u32i>
+// CIR: [[RET:%.*]] = cir.load {{.*}} : !cir.ptr<!u32i>, !u32i
+// CIR: cir.return [[RET]] : !u32i
+
+// LLVM-LABEL: @_Z8read16_2v
+// LLVM:  [[BFLOAD:%.*]] = load i16, ptr getelementptr inbounds nuw (i8, ptr {{.*}}, i64 2), align 2
+// LLVM:  [[BFCAST:%.*]] = zext i16 [[BFLOAD]] to i64
+// LLVM:  [[BF:%.*]] = trunc i64 [[BFCAST]] to i32
+// LLVM:  store i32 [[BF]], ptr {{.*}}, align 4
+// LLVM:  [[RET:%.*]] = load i32, ptr {{.*}}, align 4
+// LLVM:  ret i32 [[RET]]
+
+// OGCG-LABEL: @_Z8read16_2v
+// OGCG: [[BFLOAD:%.*]] = load i16, ptr getelementptr inbounds nuw (%struct.S2, ptr {{.*}}, i32 0, i32 1), align 2
+// OGCG-NEXT: [[BFCAST:%.*]] = zext i16 [[BFLOAD]] to i64
+// OGCG-NEXT: [[RET:%.*]] = trunc i64 [[BFCAST]] to i32
+// OGCG-NEXT: ret i32 [[RET]]
+
+void write16_1() {
+  a2.f1 = 5;
+}
+
+// CIR-LABEL: @_Z9write16_1v
+// CIR: [[CONST5:%.*]] = cir.const #cir.int<5> : !s32i
+// CIR: [[INT5:%.*]] = cir.cast(integral, [[CONST5]] : !s32i), !u64i
+// CIR: [[MEMBER:%.*]]  = cir.get_member {{.*}}[0] {name = "f1"} : !cir.ptr<!rec_S2> -> !cir.ptr<!u16i>
+// CIR: cir.set_bitfield align(8) (#bfi_f1, [[MEMBER]] : !cir.ptr<!u16i>, [[INT5]] : !u64i) -> !u64i
+// CIR: cir.return
+
+// LLVM-LABEL: @_Z9write16_1v
+// LLVM:  store i16 5, ptr {{.*}}, align 8
+// LLVM:  ret void
+
+// OGCG-LABEL: @_Z9write16_1v
+// OGCG: store i16 5, ptr {{.*}}, align 8
+// OGCG-NEXT: ret void
+
+void write16_2() {
+
+  a2.f2 = 5;
+}
+
+// CIR-LABEL: @_Z9write16_2v
+// CIR: [[CONST5:%.*]] = cir.const #cir.int<5> : !s32i
+// CIR: [[INT5:%.*]] = cir.cast(integral, [[CONST5]] : !s32i), !u64i
+// CIR: [[MEMBER:%.*]] = cir.get_member {{.*}}[1] {name = "f2"} : !cir.ptr<!rec_S2> -> !cir.ptr<!u16i>
+// CIR: cir.set_bitfield align(2) (#bfi_f2, [[MEMBER]] : !cir.ptr<!u16i>, {{.*}} : !u64i) -> !u64i
+// CIR: cir.return
+
+// LLVM-LABEL: @_Z9write16_2v
+// LLVM: store i16 5, ptr getelementptr inbounds nuw (i8, ptr {{.*}}, i64 2), align 2
+// LLVM: ret void
+
+// OGCG-LABEL: @_Z9write16_2v
+// OGCG: store i16 5, ptr getelementptr inbounds nuw (%struct.S2, ptr {{.*}}, i32 0, i32 1), align 2
+// OGCG-NEXT: ret void
+
+unsigned read32_1() {
+
+  return a3.f3;
+}
+// CIR-LABEL: @_Z8read32_1v
+// CIR: [[MEMBER:%.*]] = cir.get_member {{.*}}[1] {name = "f3"} : !cir.ptr<!rec_S3> -> !cir.ptr<!u32i>
+// CIR: [[BITFI:%.*]] = cir.get_bitfield align(4) (#bfi_f3_1, [[MEMBER]] : !cir.ptr<!u32i>) -> !u64i
+// CIR: [[BFCAST:%.*]] = cir.cast(integral, [[BITFI]] : !u64i), !u32i
+// CIR: cir.store [[BFCAST]], {{.*}} : !u32i, !cir.ptr<!u32i>
+// CIR: [[RET:%.*]] = cir.load {{.*}} : !cir.ptr<!u32i>, !u32i
+// CIR: cir.return [[RET]] : !u32i
+
+// LLVM-LABEL: @_Z8read32_1v
+// LLVM: [[BFLOAD:%.*]] = load i32, ptr getelementptr inbounds nuw (i8, ptr {{.*}}, i64 4), align 4
+// LLVM: [[BFCAST:%.*]] = zext i32 [[BFLOAD]] to i64
+// LLVM: [[BF:%.*]] = trunc i64 [[BFCAST]] to i32
+// LLVM: store i32 [[BF]], ptr {{.*}}, align 4
+// LLVM: [[RET:%.*]] = load i32, ptr {{.*}}, align 4
+// LLVM: ret i32 [[RET]]
+
+// OGCG-LABEL: @_Z8read32_1v
+// OGCG: [[BFLOAD:%.*]] = load i32, ptr getelementptr inbounds nuw (%struct.S3, ptr {{.*}}, i32 0, i32 1), align 4
+// OGCG-NEXT: [[BFCAST:%.*]] = zext i32 %bf.load to i64
+// OGCG-NEXT: [[RET:%.*]] = trunc i64 %bf.cast to i32
+// OGCG-NEXT: ret i32 [[RET]]
+
+void write32_1() {
+  a3.f3 = 5;
+}
+
+// CIR-LABEL: @_Z9write32_1v
+// CIR: [[CONST5:%.*]] = cir.const #cir.int<5> : !s32i
+// CIR: [[INT5:%.*]] = cir.cast(integral, [[CONST5]] : !s32i), !u64i
+// CIR: [[MEMBER:%.*]] = cir.get_member {{.*}}[1] {name = "f3"} : !cir.ptr<!rec_S3> -> !cir.ptr<!u32i>
+// CIR: cir.set_bitfield align(4) (#bfi_f3_1, [[MEMBER]] : !cir.ptr<!u32i>, [[INT5]] : !u64i) -> !u64i
+// CIR: cir.return
+
+// LLVM-LABEL: @_Z9write32_1v
+// LLVM:  store i32 5, ptr getelementptr inbounds nuw (i8, ptr {{.*}}, i64 4), align 4
+// LLVM:  ret void
+
+// OGCG-LABEL: @_Z9write32_1v
+// OGCG: store i32 5, ptr getelementptr inbounds nuw (%struct.S3, ptr {{.*}}, i32 0, i32 1), align 4
+// OGCG-NEXT: ret void

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

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -triple amdgcn -target-feature +gws -o /dev/null %s 2>&1 \
+// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx942 -target-feature +gws -o /dev/null %s 2>&1 \
 // RUN:   | FileCheck --check-prefix=GWS %s
 
 // GWS: warning: feature flag '+gws' is ignored since the feature is read only [-Winvalid-command-line-argument]

clang/test/CodeGenOpenCL/amdgpu-readonly-features-written-with-no-target.cl

@@ -0,0 +1,16 @@
+// REQUIRES: amdgpu-registered-target
+
+// Check the readonly feature will can be written to the IR
+// if there is no target specified.
+
+// RUN: %clang_cc1 -triple amdgcn -emit-llvm -o - %s | FileCheck --check-prefix=NOCPU %s
+// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx942 -emit-llvm -o - %s | FileCheck --check-prefix=GFX942 %s
+// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx1100 -emit-llvm -o - %s | FileCheck --check-prefix=GFX1100 %s
+// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx1200 -emit-llvm -o - %s | FileCheck --check-prefix=GFX1200 %s
+
+__attribute__((target("gws,image-insts,vmem-to-lds-load-insts"))) void test() {}
+
+// NOCPU: "target-features"="+gws,+image-insts,+vmem-to-lds-load-insts"
+// 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"
+// GFX1100: "target-features"="+16-bit-insts,+atomic-fadd-rtn-insts,+ci-insts,+dl-insts,+dot10-insts,+dot12-insts,+dot5-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize32"
+// GFX1200: "target-features"="+16-bit-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,+ci-insts,+dl-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+fp8-conversion-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize32"