Author: Joseph Huber Date: 2026-08-10T09:49:24-05:00 New Revision: 419a1e6c5fe807397e09ad5ab8d1974f0fe0038e
URL: https://github.com/llvm/llvm-project/commit/419a1e6c5fe807397e09ad5ab8d1974f0fe0038e DIFF: https://github.com/llvm/llvm-project/commit/419a1e6c5fe807397e09ad5ab8d1974f0fe0038e.diff LOG: [AMDGPU] Mark int64_t and intmax_t as LP64 (#214865) 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. Added: clang/test/Preprocessor/init-amdgpu.c Modified: clang/lib/Basic/Targets/AMDGPU.cpp clang/test/CodeGen/AMDGPU/amdgcn-av-none-amdgpu-atomic.cpp clang/test/CodeGenCXX/builtin-amdgcn-atomic-inc-dec.cpp Removed: ################################################################################ diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp index c15ac0462918d..fce626031147e 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) diff --git a/clang/test/Preprocessor/init-amdgpu.c b/clang/test/Preprocessor/init-amdgpu.c new file mode 100644 index 0000000000000..9e938ab4502a5 --- /dev/null +++ b/clang/test/Preprocessor/init-amdgpu.c @@ -0,0 +1,67 @@ +// Standalone amdgcn is LP64, so 'long' is used for the 64-bit and maximal +// integer types. +// +// 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=amdgpu9.42-amd-amdhsa < /dev/null \ +// RUN: | FileCheck -match-full-lines -check-prefix=LP64 %s +// RUN: %clang_cc1 -E -dM -triple=amdgpu10.30-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 \ +// 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=amdgpu9.0a-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=amdgpu12.50-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=amdgpu9.0a-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=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 -x hip -fcuda-is-device -E -dM -triple=amdgpu12.50-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=amdgpu9.0a-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=amdgpu12.50-amd-amdhsa -aux-triple x86_64-unknown-linux-gnu < /dev/null \ +// RUN: | FileCheck -match-full-lines -check-prefix=LP64 %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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
