https://github.com/jhuber6 updated https://github.com/llvm/llvm-project/pull/214865
>From 72bddda7810fe9db1b564f90c130d4ea83e76264 Mon Sep 17 00:00:00 2001 From: Joseph Huber <[email protected]> Date: Fri, 7 Aug 2026 17:16:44 -0500 Subject: [PATCH 1/3] [AMDGPU] Mark int64_t and intmax_t as LP64 Summary: The standalone amdgcn target currently inconsistently sets the bit widths for things like intptr_t but does not set int64_t and intmax_t. This diverges from the common Linux definitions, which HIP / OpenMP currently inherit for most cases, but doesn't correspond to other targets like Windows either. We should make this consistent with the Linux target, both so we are not in the in-between state and because that is what the vast majority of AMDGCN compilations already do through HIP. Note, this does not affect OpenCL, HIP, OpenMP, or SPIR-V compilations. It only affects the raw '--target=amdgcn-amd-amdhsa` without any other language mode set. --- clang/lib/Basic/Targets/AMDGPU.cpp | 2 ++ clang/test/CodeGen/AMDGPU/amdgcn-av-none-amdgpu-atomic.cpp | 2 +- clang/test/CodeGenCXX/builtin-amdgcn-atomic-inc-dec.cpp | 4 ++-- 3 files changed, 5 insertions(+), 3 deletions(-) diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp index 4109066ec910e..4fbc4e8a798c0 100644 --- a/clang/lib/Basic/Targets/AMDGPU.cpp +++ b/clang/lib/Basic/Targets/AMDGPU.cpp @@ -224,6 +224,8 @@ AMDGPUTargetInfo::AMDGPUTargetInfo(const llvm::Triple &Triple, SizeType = UnsignedLong; PtrDiffType = SignedLong; IntPtrType = SignedLong; + Int64Type = SignedLong; + IntMaxType = SignedLong; } MaxAtomicPromoteWidth = MaxAtomicInlineWidth = 64; diff --git a/clang/test/CodeGen/AMDGPU/amdgcn-av-none-amdgpu-atomic.cpp b/clang/test/CodeGen/AMDGPU/amdgcn-av-none-amdgpu-atomic.cpp index ada72942dcd4c..4ddd3f38dd896 100644 --- a/clang/test/CodeGen/AMDGPU/amdgcn-av-none-amdgpu-atomic.cpp +++ b/clang/test/CodeGen/AMDGPU/amdgcn-av-none-amdgpu-atomic.cpp @@ -10,7 +10,7 @@ void test_atomic_inc32_av(volatile unsigned *p, unsigned val) { [[clang::amdgpu_av("none")]] __builtin_amdgcn_atomic_inc32(p, val, __ATOMIC_SEQ_CST, "agent"); // expected-warning {{only applies to atomic expressions or Clang atomic builtins}} } -void test_atomic_dec64_av(volatile unsigned long long *p, unsigned long long val) { +void test_atomic_dec64_av(volatile unsigned long *p, unsigned long val) { [[clang::amdgpu_av("none")]] __builtin_amdgcn_atomic_dec64(p, val, __ATOMIC_ACQUIRE, "agent"); // expected-warning {{only applies to atomic expressions or Clang atomic builtins}} } diff --git a/clang/test/CodeGenCXX/builtin-amdgcn-atomic-inc-dec.cpp b/clang/test/CodeGenCXX/builtin-amdgcn-atomic-inc-dec.cpp index 09d5497a98228..d0375d177fe25 100644 --- a/clang/test/CodeGenCXX/builtin-amdgcn-atomic-inc-dec.cpp +++ b/clang/test/CodeGenCXX/builtin-amdgcn-atomic-inc-dec.cpp @@ -87,7 +87,7 @@ __attribute__((device)) void test_non_volatile_parameter32(__UINT32_TYPE__ *ptr) // CHECK-NEXT: [[TMP7:%.*]] = atomicrmw udec_wrap ptr [[TMP4]], i64 [[TMP6]] syncscope("workgroup") seq_cst, align 8, !amdgpu.no.fine.grained.memory [[META4]] // CHECK-NEXT: store i64 [[TMP7]], ptr [[RES_ASCAST]], align 8 // CHECK-NEXT: ret void -// GCN-LABEL: @_Z29test_non_volatile_parameter64Py( +// GCN-LABEL: @_Z29test_non_volatile_parameter64Pm( // GCN-NEXT: entry: // GCN-NEXT: [[PTR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5) // GCN-NEXT: [[RES:%.*]] = alloca i64, align 8, addrspace(5) @@ -213,7 +213,7 @@ __attribute__((device)) void test_volatile_parameter32(volatile __UINT32_TYPE__ // CHECK-NEXT: [[TMP7:%.*]] = atomicrmw volatile udec_wrap ptr [[TMP4]], i64 [[TMP6]] syncscope("workgroup") seq_cst, align 8, !amdgpu.no.fine.grained.memory [[META4]] // CHECK-NEXT: store i64 [[TMP7]], ptr [[RES_ASCAST]], align 8 // CHECK-NEXT: ret void -// GCN-LABEL: @_Z25test_volatile_parameter64PVy( +// GCN-LABEL: @_Z25test_volatile_parameter64PVm( // GCN-NEXT: entry: // GCN-NEXT: [[PTR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5) // GCN-NEXT: [[RES:%.*]] = alloca i64, align 8, addrspace(5) >From a5f27980dd3f9a489eedd4e8c85fff22ab23332c Mon Sep 17 00:00:00 2001 From: Joseph Huber <[email protected]> Date: Sat, 8 Aug 2026 07:34:29 -0500 Subject: [PATCH 2/3] test --- clang/test/Preprocessor/init-amdgpu.c | 60 +++++++++++++++++++++++++++ 1 file changed, 60 insertions(+) create mode 100644 clang/test/Preprocessor/init-amdgpu.c diff --git a/clang/test/Preprocessor/init-amdgpu.c b/clang/test/Preprocessor/init-amdgpu.c new file mode 100644 index 0000000000000..4ce767f0f5f46 --- /dev/null +++ b/clang/test/Preprocessor/init-amdgpu.c @@ -0,0 +1,60 @@ +// Standalone amdgcn is LP64, so 'long' is used for the 64-bit and maximal +// integer types. +// +// RUN: %clang_cc1 -E -dM -triple=amdgcn-amd-amdhsa < /dev/null \ +// RUN: | FileCheck -match-full-lines -check-prefix=LP64 %s +// RUN: %clang_cc1 -x c++ -E -dM -triple=amdgcn-amd-amdhsa < /dev/null \ +// RUN: | FileCheck -match-full-lines -check-prefix=LP64 %s +// RUN: %clang_cc1 -E -dM -triple=amdgcn-amd-amdhsa -target-cpu gfx90a < /dev/null \ +// RUN: | FileCheck -match-full-lines -check-prefix=LP64 %s +// +// r600 has 32-bit pointers and keeps the generic defaults. +// +// RUN: %clang_cc1 -E -dM -triple=r600 < /dev/null \ +// RUN: | FileCheck -match-full-lines -check-prefix=LLP64 %s +// +// OpenCL mandates its own widths regardless of the target. +// +// RUN: %clang_cc1 -x cl -cl-std=CL2.0 -E -dM -triple=amdgcn-amd-amdhsa < /dev/null \ +// RUN: | FileCheck -match-full-lines -check-prefix=OPENCL %s +// +// Offloading languages copy these from the auxiliary host target, so the +// standalone definitions above do not leak into them. +// +// RUN: %clang_cc1 -x hip -fcuda-is-device -E -dM -triple=amdgcn-amd-amdhsa \ +// RUN: -aux-triple x86_64-unknown-linux-gnu < /dev/null \ +// RUN: | FileCheck -match-full-lines -check-prefix=LP64 %s +// RUN: %clang_cc1 -x hip -fcuda-is-device -E -dM -triple=amdgcn-amd-amdhsa \ +// RUN: -aux-triple x86_64-pc-windows-msvc < /dev/null \ +// RUN: | FileCheck -match-full-lines -check-prefix=LLP64 %s +// RUN: %clang_cc1 -x hip -fcuda-is-device -E -dM -triple=amdgcn-amd-amdhsa \ +// RUN: -aux-triple i386-unknown-linux-gnu < /dev/null \ +// RUN: | FileCheck -match-full-lines -check-prefix=LLP64 %s +// RUN: %clang_cc1 -fopenmp -fopenmp-is-target-device -E -dM \ +// RUN: -triple=amdgcn-amd-amdhsa -aux-triple x86_64-pc-windows-msvc < /dev/null \ +// RUN: | FileCheck -match-full-lines -check-prefix=LLP64 %s + +// LP64-DAG: #define __INT64_TYPE__ long int +// LP64-DAG: #define __UINT64_TYPE__ long unsigned int +// LP64-DAG: #define __INT64_C_SUFFIX__ L +// LP64-DAG: #define __INTMAX_TYPE__ long int +// LP64-DAG: #define __UINTMAX_TYPE__ long unsigned int +// LP64-DAG: #define __INTMAX_C_SUFFIX__ L +// LP64-DAG: #define __INTMAX_WIDTH__ 64 +// LP64-DAG: #define __INTPTR_TYPE__ long int +// LP64-DAG: #define __SIZE_TYPE__ long unsigned int +// LP64-DAG: #define __PTRDIFF_TYPE__ long int + +// LLP64-DAG: #define __INT64_TYPE__ long long int +// LLP64-DAG: #define __UINT64_TYPE__ long long unsigned int +// LLP64-DAG: #define __INT64_C_SUFFIX__ LL +// LLP64-DAG: #define __INTMAX_TYPE__ long long int +// LLP64-DAG: #define __UINTMAX_TYPE__ long long unsigned int +// LLP64-DAG: #define __INTMAX_C_SUFFIX__ LL +// LLP64-DAG: #define __INTMAX_WIDTH__ 64 + +// OPENCL-DAG: #define __INT64_TYPE__ long int +// OPENCL-DAG: #define __INTMAX_TYPE__ long long int +// OPENCL-DAG: #define __UINTMAX_TYPE__ long long unsigned int +// OPENCL-DAG: #define __INTPTR_TYPE__ long int +// OPENCL-DAG: #define __SIZE_TYPE__ long unsigned int >From cc78402bd7a70715d6b4ba498f23a333e4908d18 Mon Sep 17 00:00:00 2001 From: Joseph Huber <[email protected]> Date: Sat, 8 Aug 2026 07:51:06 -0500 Subject: [PATCH 3/3] subarch --- clang/test/Preprocessor/init-amdgpu.c | 13 +++++++++++++ 1 file changed, 13 insertions(+) diff --git a/clang/test/Preprocessor/init-amdgpu.c b/clang/test/Preprocessor/init-amdgpu.c index 4ce767f0f5f46..e144e71d05a92 100644 --- a/clang/test/Preprocessor/init-amdgpu.c +++ b/clang/test/Preprocessor/init-amdgpu.c @@ -8,6 +8,11 @@ // RUN: %clang_cc1 -E -dM -triple=amdgcn-amd-amdhsa -target-cpu gfx90a < /dev/null \ // RUN: | FileCheck -match-full-lines -check-prefix=LP64 %s // +// RUN: %clang_cc1 -E -dM -triple=amdgpu9.0a-amd-amdhsa < /dev/null \ +// RUN: | FileCheck -match-full-lines -check-prefix=LP64 %s +// RUN: %clang_cc1 -x c++ -E -dM -triple=amdgpu12.50-amd-amdhsa < /dev/null \ +// RUN: | FileCheck -match-full-lines -check-prefix=LP64 %s +// // r600 has 32-bit pointers and keeps the generic defaults. // // RUN: %clang_cc1 -E -dM -triple=r600 < /dev/null \ @@ -17,6 +22,8 @@ // // RUN: %clang_cc1 -x cl -cl-std=CL2.0 -E -dM -triple=amdgcn-amd-amdhsa < /dev/null \ // RUN: | FileCheck -match-full-lines -check-prefix=OPENCL %s +// RUN: %clang_cc1 -x cl -cl-std=CL2.0 -E -dM -triple=amdgpu9.0a-amd-amdhsa < /dev/null \ +// RUN: | FileCheck -match-full-lines -check-prefix=OPENCL %s // // Offloading languages copy these from the auxiliary host target, so the // standalone definitions above do not leak into them. @@ -30,9 +37,15 @@ // RUN: %clang_cc1 -x hip -fcuda-is-device -E -dM -triple=amdgcn-amd-amdhsa \ // RUN: -aux-triple i386-unknown-linux-gnu < /dev/null \ // RUN: | FileCheck -match-full-lines -check-prefix=LLP64 %s +// RUN: %clang_cc1 -x hip -fcuda-is-device -E -dM -triple=amdgpu9.0a-amd-amdhsa \ +// RUN: -aux-triple x86_64-pc-windows-msvc < /dev/null \ +// RUN: | FileCheck -match-full-lines -check-prefix=LLP64 %s // RUN: %clang_cc1 -fopenmp -fopenmp-is-target-device -E -dM \ // RUN: -triple=amdgcn-amd-amdhsa -aux-triple x86_64-pc-windows-msvc < /dev/null \ // RUN: | FileCheck -match-full-lines -check-prefix=LLP64 %s +// RUN: %clang_cc1 -fopenmp -fopenmp-is-target-device -E -dM \ +// RUN: -triple=amdgpu9.0a-amd-amdhsa -aux-triple x86_64-pc-windows-msvc < /dev/null \ +// RUN: | FileCheck -match-full-lines -check-prefix=LLP64 %s // LP64-DAG: #define __INT64_TYPE__ long int // LP64-DAG: #define __UINT64_TYPE__ long unsigned int _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
