[clang] [flang] [llvm] [AMDGPU] Use a target feature to enable __builtin_amdgcn_global_load_lds on gfx9/10 (PR #133055)

Matt Arsenault via llvm-commits llvm-commits at lists.llvm.org
Mon Mar 31 02:28:36 PDT 2025


Juan Manuel Martinez =?utf-8?q?Caamaño?= <juamarti at amd.com>,
Juan Manuel Martinez =?utf-8?q?Caamaño?= <juamarti at amd.com>,
Juan Manuel Martinez =?utf-8?q?Caamaño?= <juamarti at amd.com>
Message-ID:
In-Reply-To: <llvm.org/llvm/llvm-project/pull/133055 at github.com>


================
@@ -260,7 +260,7 @@ AMDGPUTargetInfo::AMDGPUTargetInfo(const llvm::Triple &Triple,
 
   MaxAtomicPromoteWidth = MaxAtomicInlineWidth = 64;
   CUMode = !(GPUFeatures & llvm::AMDGPU::FEATURE_WGP);
-  for (auto F : {"image-insts", "gws"})
+  for (auto F : {"image-insts", "gws", "mem-to-lds-load-insts"})
----------------
arsenm wrote:

```suggestion
  for (auto F : {"image-insts", "gws", "vmem-to-lds-load-insts"})
```

Not sure about the naming. The manual refers to this as "LDS DMA", but was new in gfx8 (or maybe, the store case was new? I thought the load case always existed, we just never made use of it).

https://github.com/llvm/llvm-project/pull/133055


More information about the llvm-commits mailing list