Author: Matt Arsenault Date: 2026-07-25T21:50:11+02:00 New Revision: 56f4b9c4e875477c0fbc73160f96524d4915040c
URL: https://github.com/llvm/llvm-project/commit/56f4b9c4e875477c0fbc73160f96524d4915040c DIFF: https://github.com/llvm/llvm-project/commit/56f4b9c4e875477c0fbc73160f96524d4915040c.diff LOG: clang/AMDGPU: Migrate cc1 tests to subarch triples (11) (#211956) Rewrite the second half of CodeGenOpenCL cc1 test RUN lines to the new subarch triple form, dropping the redundant -target-cpu. Co-authored-by: Claude (Opus 4.8) <[email protected]> Added: Modified: clang/test/CodeGenOpenCL/amdgcn-buffer-rsrc-type.cl clang/test/CodeGenOpenCL/amdgpu-cluster-dims.cl clang/test/CodeGenOpenCL/amdgpu-features-default-delta.cl clang/test/CodeGenOpenCL/amdgpu-module-flag-xnack-sramecc.cl clang/test/CodeGenOpenCL/amdgpu-readonly-features-written-with-no-target.cl clang/test/CodeGenOpenCL/amdgpu-xnack-any-only.cl clang/test/CodeGenOpenCL/atomics-cas-remarks-gfx90a.cl clang/test/CodeGenOpenCL/atomics-unsafe-hw-remarks-gfx90a.cl clang/test/CodeGenOpenCL/builtins-alloca.cl clang/test/CodeGenOpenCL/builtins-amdgcn-asyncmark.cl clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-gfx12.cl clang/test/CodeGenOpenCL/builtins-amdgcn-fiji.cl clang/test/CodeGenOpenCL/builtins-amdgcn-flat-address-space.cl clang/test/CodeGenOpenCL/builtins-amdgcn-fp-atomics-gfx11-err.cl clang/test/CodeGenOpenCL/builtins-amdgcn-fp-atomics-gfx90a-err.cl clang/test/CodeGenOpenCL/builtins-amdgcn-fp8.cl clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11-param-err.cl clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-5-generic-err.cl clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-param-err.cl clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-wmma-w32.cl clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-wmma-w64.cl clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-load-monitor.cl clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1251-wmma-w32.cl clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-store.cl clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-gfx12-w32-err.cl clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-gfx12-w64-err.cl clang/test/CodeGenOpenCL/builtins-amdgcn-interp.cl clang/test/CodeGenOpenCL/builtins-amdgcn-make-buffer-rsrc.cl clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-max.cl clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-load.cl clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-store-format.cl clang/test/CodeGenOpenCL/builtins-amdgcn-raytracing.cl clang/test/CodeGenOpenCL/builtins-amdgcn-s-buffer-load.cl clang/test/CodeGenOpenCL/builtins-amdgcn-smfmac-err.cl clang/test/CodeGenOpenCL/builtins-amdgcn-swmmac-w32-gfx10-err.cl clang/test/CodeGenOpenCL/builtins-amdgcn-swmmac-w64-gfx10-err.cl clang/test/CodeGenOpenCL/builtins-amdgcn-swmmac-w64.cl clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl clang/test/CodeGenOpenCL/builtins-amdgcn-wmma-w32-gfx10-err.cl clang/test/CodeGenOpenCL/builtins-amdgcn-wmma-w64.cl clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx1030.cl clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx12.cl clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx8.cl clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx90a.cl clang/test/CodeGenOpenCL/check-atomic-alignment.cl Removed: ################################################################################ diff --git a/clang/test/CodeGenOpenCL/amdgcn-buffer-rsrc-type.cl b/clang/test/CodeGenOpenCL/amdgcn-buffer-rsrc-type.cl index 99d4a65c41801..6d767f4bcb5db 100644 --- a/clang/test/CodeGenOpenCL/amdgcn-buffer-rsrc-type.cl +++ b/clang/test/CodeGenOpenCL/amdgcn-buffer-rsrc-type.cl @@ -1,6 +1,6 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 // REQUIRES: amdgpu-registered-target - // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu verde -emit-llvm -o - %s | FileCheck %s + // RUN: %clang_cc1 -triple amdgpu6.01-unknown-unknown -emit-llvm -o - %s | FileCheck %s typedef struct AA_ty { int x; diff --git a/clang/test/CodeGenOpenCL/amdgpu-cluster-dims.cl b/clang/test/CodeGenOpenCL/amdgpu-cluster-dims.cl index fa64969555a89..e6536c2c494a9 100644 --- a/clang/test/CodeGenOpenCL/amdgpu-cluster-dims.cl +++ b/clang/test/CodeGenOpenCL/amdgpu-cluster-dims.cl @@ -1,5 +1,5 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-attributes --check-globals all --include-generated-funcs --prefix-filecheck-ir-name VAR --version 6 -// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx1250 -disable-llvm-passes -fno-ident -emit-llvm %s -o - | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgpu12.50-amd-amdhsa -disable-llvm-passes -fno-ident -emit-llvm %s -o - | FileCheck %s kernel void foo(global int *p) { *p = 1; } // CHECK: Function Attrs: convergent norecurse nounwind @@ -24,8 +24,8 @@ kernel void foo(global int *p) { *p = 1; } // CHECK-NEXT: ret void // //. -// CHECK: attributes #[[ATTR0]] = { convergent norecurse nounwind "amdgpu-cluster-dims"="0,0,0" "amdgpu-flat-work-group-size"="1,256" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx1250" } -// CHECK: attributes #[[ATTR1]] = { alwaysinline convergent norecurse nounwind "amdgpu-cluster-dims"="0,0,0" "amdgpu-flat-work-group-size"="1,256" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx1250" } +// CHECK: attributes #[[ATTR0]] = { convergent norecurse nounwind "amdgpu-cluster-dims"="0,0,0" "amdgpu-flat-work-group-size"="1,256" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +// CHECK: attributes #[[ATTR1]] = { alwaysinline convergent norecurse nounwind "amdgpu-cluster-dims"="0,0,0" "amdgpu-flat-work-group-size"="1,256" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } // CHECK: attributes #[[ATTR2]] = { convergent nounwind } //. // CHECK: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600} diff --git a/clang/test/CodeGenOpenCL/amdgpu-features-default-delta.cl b/clang/test/CodeGenOpenCL/amdgpu-features-default-delta.cl index 56f8ed2c0e979..381294169dd72 100644 --- a/clang/test/CodeGenOpenCL/amdgpu-features-default-delta.cl +++ b/clang/test/CodeGenOpenCL/amdgpu-features-default-delta.cl @@ -5,26 +5,27 @@ // Default behavior for gfx90a: test_default has no target-features, // test_explicit_attr has only the delta (+gfx11-insts). -// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx90a -emit-llvm -o - %s | FileCheck --check-prefix=GFX90A %s +// RUN: %clang_cc1 -triple amdgpu9.0a -emit-llvm -o - %s | FileCheck --check-prefix=GFX90A %s // With -target-feature, both functions get the delta feature. // gfx1030 defaults to wavefrontsize32, so +wavefrontsize64 is a delta. -// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx1030 -target-feature +wavefrontsize64 -emit-llvm -o - %s | FileCheck --check-prefix=CMDLINE %s +// RUN: %clang_cc1 -triple amdgpu10.30 -target-feature +wavefrontsize64 -emit-llvm -o - %s | FileCheck --check-prefix=CMDLINE %s // GFX90A-LABEL: define {{.*}} @test_default() // GFX90A-SAME: #[[ATTR_DEFAULT:[0-9]+]] // GFX90A-LABEL: define {{.*}} @test_explicit_attr() // GFX90A-SAME: #[[ATTR_EXPLICIT:[0-9]+]] // -// test_default should have target-cpu but NO target-features +// The processor is encoded in the triple's subarch, so there is no +// target-cpu attribute. test_default has NO target-features. // GFX90A: attributes #[[ATTR_DEFAULT]] = { -// GFX90A-SAME: "target-cpu"="gfx90a" +// GFX90A-NOT: "target-cpu" // GFX90A-NOT: "target-features" // GFX90A-SAME: } // -// test_explicit_attr should have target-cpu and ONLY the delta target-features +// test_explicit_attr should have ONLY the delta target-features // GFX90A: attributes #[[ATTR_EXPLICIT]] = { -// GFX90A-SAME: "target-cpu"="gfx90a" +// GFX90A-NOT: "target-cpu" // GFX90A-SAME: "target-features"="+gfx11-insts" // GFX90A-SAME: } @@ -36,12 +37,12 @@ // CMDLINE-SAME: #[[ATTR_EXPLICIT:[0-9]+]] // // CMDLINE: attributes #[[ATTR_DEFAULT]] = { -// CMDLINE-SAME: "target-cpu"="gfx1030" +// CMDLINE-NOT: "target-cpu" // CMDLINE-SAME: "target-features"="+wavefrontsize64" // CMDLINE-SAME: } // // CMDLINE: attributes #[[ATTR_EXPLICIT]] = { -// CMDLINE-SAME: "target-cpu"="gfx1030" +// CMDLINE-NOT: "target-cpu" // CMDLINE-SAME: "target-features"="{{[^"]*}}+gfx11-insts{{[^"]*}}+wavefrontsize64{{[^"]*}}" // CMDLINE-SAME: } diff --git a/clang/test/CodeGenOpenCL/amdgpu-module-flag-xnack-sramecc.cl b/clang/test/CodeGenOpenCL/amdgpu-module-flag-xnack-sramecc.cl index 439a815f5c42d..4d085907ddf89 100644 --- a/clang/test/CodeGenOpenCL/amdgpu-module-flag-xnack-sramecc.cl +++ b/clang/test/CodeGenOpenCL/amdgpu-module-flag-xnack-sramecc.cl @@ -1,13 +1,13 @@ // Test that xnack and sramecc module flags are emitted based on -m flags -// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \ +// RUN: %clang_cc1 -triple amdgpu9.0a-amd-amdhsa \ // RUN: -mxnack -mno-sramecc \ // RUN: -emit-llvm -o - %s | FileCheck %s --check-prefixes=XNACK-ON,SRAMECC-OFF -// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \ +// RUN: %clang_cc1 -triple amdgpu9.0a-amd-amdhsa \ // RUN: -mno-xnack -msramecc \ // RUN: -emit-llvm -o - %s | FileCheck %s --check-prefixes=XNACK-OFF,SRAMECC-ON -// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \ +// RUN: %clang_cc1 -triple amdgpu9.0a-amd-amdhsa \ // RUN: -emit-llvm -o - %s | FileCheck %s --check-prefix=NO-FLAGS // XNACK-ON-DAG: !{i32 1, !"amdgpu.xnack", i32 1} diff --git a/clang/test/CodeGenOpenCL/amdgpu-readonly-features-written-with-no-target.cl b/clang/test/CodeGenOpenCL/amdgpu-readonly-features-written-with-no-target.cl index 1239df2a96d2e..5347a624e5c78 100644 --- a/clang/test/CodeGenOpenCL/amdgpu-readonly-features-written-with-no-target.cl +++ b/clang/test/CodeGenOpenCL/amdgpu-readonly-features-written-with-no-target.cl @@ -4,9 +4,9 @@ // 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 %s -// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx1100 -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx1200 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu9.42 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu11.00 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu12.00 -emit-llvm -o - %s | FileCheck %s __attribute__((target("gws,image-insts,vmem-to-lds-load-insts"))) void test() {} diff --git a/clang/test/CodeGenOpenCL/amdgpu-xnack-any-only.cl b/clang/test/CodeGenOpenCL/amdgpu-xnack-any-only.cl index 709ca1acf4cb3..f34938f59460d 100644 --- a/clang/test/CodeGenOpenCL/amdgpu-xnack-any-only.cl +++ b/clang/test/CodeGenOpenCL/amdgpu-xnack-any-only.cl @@ -4,19 +4,19 @@ // TODO: In the future, clang should not emit the flag for targets that don't support // xnack control. -// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx12-5-generic \ +// RUN: %clang_cc1 -triple amdgpu12.5-amd-amdhsa \ // RUN: -mxnack -emit-llvm -o - %s | FileCheck %s --check-prefix=XNACK-ON -// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx12-5-generic \ +// RUN: %clang_cc1 -triple amdgpu12.5-amd-amdhsa \ // RUN: -mno-xnack -emit-llvm -o - %s | FileCheck %s --check-prefix=XNACK-OFF -// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx12-5-generic \ +// RUN: %clang_cc1 -triple amdgpu12.5-amd-amdhsa \ // RUN: -emit-llvm -o - %s | FileCheck %s --check-prefix=NO-FLAGS -// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx1250 \ +// RUN: %clang_cc1 -triple amdgpu12.50-amd-amdhsa \ // RUN: -mxnack -emit-llvm -o - %s | FileCheck %s --check-prefix=XNACK-ON -// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx1250 \ +// RUN: %clang_cc1 -triple amdgpu12.50-amd-amdhsa \ // RUN: -mno-xnack -emit-llvm -o - %s | FileCheck %s --check-prefix=XNACK-OFF // Module flags are emitted regardless of target support diff --git a/clang/test/CodeGenOpenCL/atomics-cas-remarks-gfx90a.cl b/clang/test/CodeGenOpenCL/atomics-cas-remarks-gfx90a.cl index 59bf87b554af3..bca5072a09257 100644 --- a/clang/test/CodeGenOpenCL/atomics-cas-remarks-gfx90a.cl +++ b/clang/test/CodeGenOpenCL/atomics-cas-remarks-gfx90a.cl @@ -1,8 +1,8 @@ -// RUN: %clang_cc1 %s -cl-std=CL2.0 -O0 -triple=amdgcn-amd-amdhsa -target-cpu gfx90a \ +// RUN: %clang_cc1 %s -cl-std=CL2.0 -O0 -triple=amdgpu9.0a-amd-amdhsa \ // RUN: -Rpass=atomic-expand -S -o - 2>&1 | \ // RUN: FileCheck %s --check-prefix=REMARK -// RUN: %clang_cc1 %s -cl-std=CL2.0 -O0 -triple=amdgcn-amd-amdhsa -target-cpu gfx90a \ +// RUN: %clang_cc1 %s -cl-std=CL2.0 -O0 -triple=amdgpu9.0a-amd-amdhsa \ // RUN: -Rpass=atomic-expand -emit-llvm -o - 2>&1 | \ // RUN: FileCheck %s --check-prefix=GFX90A-CAS diff --git a/clang/test/CodeGenOpenCL/atomics-unsafe-hw-remarks-gfx90a.cl b/clang/test/CodeGenOpenCL/atomics-unsafe-hw-remarks-gfx90a.cl index 0dafb44f12a3c..7fd20035ab430 100644 --- a/clang/test/CodeGenOpenCL/atomics-unsafe-hw-remarks-gfx90a.cl +++ b/clang/test/CodeGenOpenCL/atomics-unsafe-hw-remarks-gfx90a.cl @@ -1,8 +1,8 @@ -// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple=amdgcn-amd-amdhsa -target-cpu gfx90a \ +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple=amdgpu9.0a-amd-amdhsa \ // RUN: -Rpass=si-lower -munsafe-fp-atomics %s -emit-llvm -o - 2>&1 | \ // RUN: FileCheck %s --check-prefix=GFX90A-HW -// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple=amdgcn-amd-amdhsa -target-cpu gfx90a \ +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple=amdgpu9.0a-amd-amdhsa \ // RUN: -Rpass=si-lower -munsafe-fp-atomics %s -S -o - 2>&1 | \ // RUN: FileCheck %s --check-prefix=GFX90A-HW-REMARK diff --git a/clang/test/CodeGenOpenCL/builtins-alloca.cl b/clang/test/CodeGenOpenCL/builtins-alloca.cl index 01f12788342d9..ff591d6d537f2 100644 --- a/clang/test/CodeGenOpenCL/builtins-alloca.cl +++ b/clang/test/CodeGenOpenCL/builtins-alloca.cl @@ -3,7 +3,7 @@ // RUN: -emit-llvm -o - | FileCheck --check-prefixes=OPENCL12 %s // RUN: %clang_cc1 %s -O0 -triple amdgcn-amd-amdhsa -cl-std=CL2.0 \ // RUN: -emit-llvm -o - | FileCheck --check-prefixes=OPENCL20 %s -// RUN: %clang_cc1 %s -O0 -triple amdgcn-amd-mesa3d -target-cpu gfx600 -cl-std=CL3.0 \ +// RUN: %clang_cc1 %s -O0 -triple amdgpu6.00-amd-mesa3d -cl-std=CL3.0 \ // RUN: -emit-llvm -o - | FileCheck --check-prefixes=OPENCL30 %s // RUN: %clang_cc1 %s -O0 -triple amdgcn-amd-amdhsa -cl-std=CL3.0 \ // RUN: -emit-llvm -o - | FileCheck --check-prefixes=OPENCL30GAS %s diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-asyncmark.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-asyncmark.cl index e2e8911b5a7d4..838bfdbdcb8f1 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-asyncmark.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-asyncmark.cl @@ -1,8 +1,8 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py -// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx900 -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx942 -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx1250 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgpu9.00-unknown-unknown -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgpu9.42-unknown-unknown -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgpu10.10-unknown-unknown -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgpu12.50-unknown-unknown -emit-llvm -o - %s | FileCheck %s // REQUIRES: amdgpu-registered-target // CHECK-LABEL: @test_invocation( diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl index e8c17fd39a07a..ae04f2b9bc69d 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl @@ -1,8 +1,8 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu hawaii -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu fiji -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx906 -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgpu7.01-unknown-unknown -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgpu8.03-unknown-unknown -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgpu9.06-unknown-unknown -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgpu10.10-unknown-unknown -emit-llvm -o - %s | FileCheck %s typedef unsigned int uint; typedef unsigned long ulong; diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl index b8c46039dac53..a7d320e97775d 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl @@ -1,7 +1,7 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -verify -emit-llvm -o - %s -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -verify -emit-llvm -o - %s +// RUN: %clang_cc1 -triple amdgpu9.00-unknown-unknown -verify -emit-llvm -o - %s +// RUN: %clang_cc1 -triple amdgpu10.10-unknown-unknown -verify -emit-llvm -o - %s typedef unsigned int uint; typedef half __attribute__((ext_vector_type(2))) half2; diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-gfx12.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-gfx12.cl index 576fb475f883d..f8564c7853c68 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-gfx12.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-gfx12.cl @@ -1,7 +1,7 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1170 -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1200 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu11.70-unknown-unknown -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu12.00-unknown-unknown -emit-llvm -o - %s | FileCheck %s typedef unsigned int uint; diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-fiji.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-fiji.cl index 2178718f90d5a..a364f2d85f3bd 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-fiji.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-fiji.cl @@ -1,5 +1,5 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu fiji -emit-llvm -o - %s | FileCheck -enable-var-scope --check-prefixes=CHECK %s +// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgpu8.03-unknown-unknown -emit-llvm -o - %s | FileCheck -enable-var-scope --check-prefixes=CHECK %s #pragma OPENCL EXTENSION cl_khr_fp64 : enable diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-flat-address-space.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-flat-address-space.cl index 28f418ffb9e6a..f7ae305c3a139 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-flat-address-space.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-flat-address-space.cl @@ -1,5 +1,5 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu tahiti -emit-llvm -disable-llvm-passes -o - %s | FileCheck -enable-var-scope %s +// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgpu6.00-unknown-unknown -emit-llvm -disable-llvm-passes -o - %s | FileCheck -enable-var-scope %s // SI did not actually support flat addressing, but we can codegen the address // space test builtins. The target specfic part is a load from the implicit diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-fp-atomics-gfx11-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-fp-atomics-gfx11-err.cl index 39191322ca6e4..aca30a231f8b2 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-fp-atomics-gfx11-err.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-fp-atomics-gfx11-err.cl @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx1100 \ +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgpu11.00-amd-amdhsa \ // RUN: -verify -S -o - %s // REQUIRES: amdgpu-registered-target diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-fp-atomics-gfx90a-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-fp-atomics-gfx90a-err.cl index f651ce349e206..8d9c4107189d4 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-fp-atomics-gfx90a-err.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-fp-atomics-gfx90a-err.cl @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \ +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgpu9.0a-amd-amdhsa \ // RUN: -verify -S -o - %s // REQUIRES: amdgpu-registered-target diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-fp8.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-fp8.cl index 83a44be930ae0..c18f8e5fa34be 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-fp8.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-fp8.cl @@ -1,8 +1,8 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx942 -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1170 -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1200 -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1250 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu9.42-unknown-unknown -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu11.70-unknown-unknown -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu12.00-unknown-unknown -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu12.50-unknown-unknown -emit-llvm -o - %s | FileCheck %s typedef float v2f __attribute__((ext_vector_type(2))); diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11-param-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11-param-err.cl index fc835e1a0933c..e7ddafa22528d 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11-param-err.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11-param-err.cl @@ -1,6 +1,6 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -verify -emit-llvm -o - %s +// RUN: %clang_cc1 -triple amdgpu11.00-unknown-unknown -verify -emit-llvm -o - %s typedef unsigned int uint; typedef uint uint2 __attribute__((ext_vector_type(2))); diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-5-generic-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-5-generic-err.cl index 3bb9ab8b74910..35ebff87c8a88 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-5-generic-err.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-5-generic-err.cl @@ -1,5 +1,5 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx12-5-generic -verify -S -o - %s +// RUN: %clang_cc1 -triple amdgpu12.5-amd-amdhsa -verify -S -o - %s typedef short v2s __attribute__((ext_vector_type(2))); typedef unsigned int v4u __attribute__((ext_vector_type(4))); diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl index 9f97265759ba0..d8fb243865f76 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl @@ -1,6 +1,6 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1200 -verify -emit-llvm -o - %s +// RUN: %clang_cc1 -triple amdgpu12.00-unknown-unknown -verify -emit-llvm -o - %s typedef unsigned int uint; diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-param-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-param-err.cl index 1a5043328895a..480801b4eecbb 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-param-err.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-param-err.cl @@ -1,6 +1,6 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1200 -verify -emit-llvm -o - %s +// RUN: %clang_cc1 -triple amdgpu12.00-unknown-unknown -verify -emit-llvm -o - %s kernel void builtins_amdgcn_s_barrier_signal_err(global int* in, global int* out, int barrier) { diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-wmma-w32.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-wmma-w32.cl index 3f36c502f4a01..9f80bc374249a 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-wmma-w32.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-wmma-w32.cl @@ -1,7 +1,7 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1170 -target-feature +wavefrontsize32 -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1200 -target-feature +wavefrontsize32 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu11.70-unknown-unknown -target-feature +wavefrontsize32 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu12.00-unknown-unknown -target-feature +wavefrontsize32 -emit-llvm -o - %s | FileCheck %s typedef int v2i __attribute__((ext_vector_type(2))); typedef float v8f __attribute__((ext_vector_type(8))); diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-wmma-w64.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-wmma-w64.cl index 2e1309ebf0001..9e9fe74b3e4c5 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-wmma-w64.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-wmma-w64.cl @@ -1,7 +1,7 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1170 -target-feature +wavefrontsize64 -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1200 -target-feature +wavefrontsize64 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu11.70-unknown-unknown -target-feature +wavefrontsize64 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu12.00-unknown-unknown -target-feature +wavefrontsize64 -emit-llvm -o - %s | FileCheck %s typedef float v4f __attribute__((ext_vector_type(4))); typedef half v4h __attribute__((ext_vector_type(4))); diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl index 332a2fa94ee92..bcaa638d848e4 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl @@ -1,5 +1,5 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py -// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx1200 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgpu12.00-unknown-unknown -emit-llvm -o - %s | FileCheck %s // REQUIRES: amdgpu-registered-target diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-load-monitor.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-load-monitor.cl index 4e285fcb217e2..784632d52ef3f 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-load-monitor.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-load-monitor.cl @@ -1,13 +1,13 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu gfx1250 -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-GFX1250 +// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgpu12.50-unknown-unknown -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-GFX1250 typedef int v2i __attribute__((ext_vector_type(2))); typedef int v4i __attribute__((ext_vector_type(4))); // CHECK-GFX1250-LABEL: @test_amdgcn_global_load_monitor_b32( // CHECK-GFX1250-NEXT: entry: -// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.global.load.monitor.b32.i32(ptr addrspace(1) [[INPTR:%.*]], i32 0, metadata [[META8:![0-9]+]]) +// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.global.load.monitor.b32.i32(ptr addrspace(1) [[INPTR:%.*]], i32 0, metadata [[META7:![0-9]+]]) // CHECK-GFX1250-NEXT: ret i32 [[TMP0]] // int test_amdgcn_global_load_monitor_b32(global int* inptr) @@ -17,7 +17,7 @@ int test_amdgcn_global_load_monitor_b32(global int* inptr) // CHECK-GFX1250-LABEL: @test_amdgcn_global_load_monitor_b64( // CHECK-GFX1250-NEXT: entry: -// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.global.load.monitor.b64.v2i32(ptr addrspace(1) [[INPTR:%.*]], i32 2, metadata [[META9:![0-9]+]]) +// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.global.load.monitor.b64.v2i32(ptr addrspace(1) [[INPTR:%.*]], i32 2, metadata [[META8:![0-9]+]]) // CHECK-GFX1250-NEXT: ret <2 x i32> [[TMP0]] // v2i test_amdgcn_global_load_monitor_b64(global v2i* inptr) @@ -27,7 +27,7 @@ v2i test_amdgcn_global_load_monitor_b64(global v2i* inptr) // CHECK-GFX1250-LABEL: @test_amdgcn_global_load_monitor_b128( // CHECK-GFX1250-NEXT: entry: -// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.global.load.monitor.b128.v4i32(ptr addrspace(1) [[INPTR:%.*]], i32 2, metadata [[META10:![0-9]+]]) +// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.global.load.monitor.b128.v4i32(ptr addrspace(1) [[INPTR:%.*]], i32 2, metadata [[META9:![0-9]+]]) // CHECK-GFX1250-NEXT: ret <4 x i32> [[TMP0]] // v4i test_amdgcn_global_load_monitor_b128(global v4i* inptr) @@ -37,7 +37,7 @@ v4i test_amdgcn_global_load_monitor_b128(global v4i* inptr) // CHECK-GFX1250-LABEL: @test_amdgcn_flat_load_monitor_b32( // CHECK-GFX1250-NEXT: entry: -// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.flat.load.monitor.b32.i32(ptr [[INPTR:%.*]], i32 0, metadata [[META8]]) +// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.flat.load.monitor.b32.i32(ptr [[INPTR:%.*]], i32 0, metadata [[META7]]) // CHECK-GFX1250-NEXT: ret i32 [[TMP0]] // int test_amdgcn_flat_load_monitor_b32(int* inptr) @@ -47,7 +47,7 @@ int test_amdgcn_flat_load_monitor_b32(int* inptr) // CHECK-GFX1250-LABEL: @test_amdgcn_flat_load_monitor_b64( // CHECK-GFX1250-NEXT: entry: -// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.flat.load.monitor.b64.v2i32(ptr [[INPTR:%.*]], i32 5, metadata [[META11:![0-9]+]]) +// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.flat.load.monitor.b64.v2i32(ptr [[INPTR:%.*]], i32 5, metadata [[META10:![0-9]+]]) // CHECK-GFX1250-NEXT: ret <2 x i32> [[TMP0]] // v2i test_amdgcn_flat_load_monitor_b64(v2i* inptr) @@ -57,7 +57,7 @@ v2i test_amdgcn_flat_load_monitor_b64(v2i* inptr) // CHECK-GFX1250-LABEL: @test_amdgcn_flat_load_monitor_b128( // CHECK-GFX1250-NEXT: entry: -// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.flat.load.monitor.b128.v4i32(ptr [[INPTR:%.*]], i32 0, metadata [[META8]]) +// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.flat.load.monitor.b128.v4i32(ptr [[INPTR:%.*]], i32 0, metadata [[META7]]) // CHECK-GFX1250-NEXT: ret <4 x i32> [[TMP0]] // v4i test_amdgcn_flat_load_monitor_b128(v4i* inptr) @@ -67,7 +67,7 @@ v4i test_amdgcn_flat_load_monitor_b128(v4i* inptr) // CHECK-GFX1250-LABEL: @test_amdgcn_global_load_monitor_b32_wavefront( // CHECK-GFX1250-NEXT: entry: -// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.global.load.monitor.b32.i32(ptr addrspace(1) [[INPTR:%.*]], i32 0, metadata [[META12:![0-9]+]]) +// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.global.load.monitor.b32.i32(ptr addrspace(1) [[INPTR:%.*]], i32 0, metadata [[META11:![0-9]+]]) // CHECK-GFX1250-NEXT: ret i32 [[TMP0]] // int test_amdgcn_global_load_monitor_b32_wavefront(global int* inptr) @@ -77,7 +77,7 @@ int test_amdgcn_global_load_monitor_b32_wavefront(global int* inptr) // CHECK-GFX1250-LABEL: @test_amdgcn_global_load_monitor_b32_single( // CHECK-GFX1250-NEXT: entry: -// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.global.load.monitor.b32.i32(ptr addrspace(1) [[INPTR:%.*]], i32 0, metadata [[META13:![0-9]+]]) +// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.global.load.monitor.b32.i32(ptr addrspace(1) [[INPTR:%.*]], i32 0, metadata [[META12:![0-9]+]]) // CHECK-GFX1250-NEXT: ret i32 [[TMP0]] // int test_amdgcn_global_load_monitor_b32_single(global int* inptr) @@ -85,9 +85,3 @@ int test_amdgcn_global_load_monitor_b32_single(global int* inptr) return __builtin_amdgcn_global_load_monitor_b32(inptr, __ATOMIC_RELAXED, __MEMORY_SCOPE_SINGLE); } -// CHECK-GFX1250: [[META8]] = !{!""} -// CHECK-GFX1250: [[META9]] = !{!"agent"} -// CHECK-GFX1250: [[META10]] = !{!"workgroup"} -// CHECK-GFX1250: [[META11]] = !{!"cluster"} -// CHECK-GFX1250: [[META12]] = !{!"wavefront"} -// CHECK-GFX1250: [[META13]] = !{!"singlethread"} diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1251-wmma-w32.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1251-wmma-w32.cl index 720a923cae542..743b72c945488 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1251-wmma-w32.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1251-wmma-w32.cl @@ -1,6 +1,6 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1251 -target-feature +wavefrontsize32 -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-GFX1251 +// RUN: %clang_cc1 -triple amdgpu12.51-unknown-unknown -target-feature +wavefrontsize32 -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-GFX1251 typedef double v8d __attribute__((ext_vector_type(8))); typedef double v2d __attribute__((ext_vector_type(2))); diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl index a86e16a933728..37e41f0cbe2e4 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl @@ -1,6 +1,6 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu9.00-unknown-unknown -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu10.10-unknown-unknown -emit-llvm -o - %s | FileCheck %s #pragma OPENCL EXTENSION cl_khr_fp16 : enable typedef unsigned int uint; diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-store.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-store.cl index 63d7fcac16874..cd34500a94a24 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-store.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-store.cl @@ -1,10 +1,10 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals smart // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu gfx900 -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu gfx1200 -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu gfx1250 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgpu9.00-unknown-unknown -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgpu10.10-unknown-unknown -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgpu11.00-unknown-unknown -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgpu12.00-unknown-unknown -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgpu12.50-unknown-unknown -emit-llvm -o - %s | FileCheck %s typedef __attribute__((__vector_size__(4 * sizeof(unsigned int)))) unsigned int v4u32; typedef v4u32 __global *global_ptr_to_v4u32; diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-gfx12-w32-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-gfx12-w32-err.cl index 6f7a93ef897ac..88f13c3a8faa1 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-gfx12-w32-err.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-gfx12-w32-err.cl @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx1200 -target-feature +wavefrontsize32 \ +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgpu12.00-amd-amdhsa -target-feature +wavefrontsize32 \ // RUN: -verify -S -o - %s // REQUIRES: amdgpu-registered-target diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-gfx12-w64-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-gfx12-w64-err.cl index b7323f1b41c2a..dec0898568da4 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-gfx12-w64-err.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-tr-gfx12-w64-err.cl @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx1200 -target-feature +wavefrontsize64 \ +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgpu12.00-amd-amdhsa -target-feature +wavefrontsize64 \ // RUN: -verify -S -o - %s // REQUIRES: amdgpu-registered-target diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-interp.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-interp.cl index ca4e890af691d..a4f3f27742aa2 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-interp.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-interp.cl @@ -1,5 +1,5 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu9.00-unknown-unknown -emit-llvm -o - %s | FileCheck %s #pragma OPENCL EXTENSION cl_khr_fp16 : enable diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-make-buffer-rsrc.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-make-buffer-rsrc.cl index 4b5232c0010aa..7dcc1ca823c7b 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-make-buffer-rsrc.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-make-buffer-rsrc.cl @@ -1,6 +1,6 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -target-cpu verde -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu6.01-unknown-unknown -cl-std=CL2.0 -emit-llvm -o - %s | FileCheck %s // CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p0( // CHECK-NEXT: entry: diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-max.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-max.cl index 90c423511477d..6928032d8c45a 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-max.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-max.cl @@ -1,7 +1,7 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5 -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx10-1-generic -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx10-3-generic -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1250 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu10.1-unknown-unknown -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu10.3-unknown-unknown -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu12.50-unknown-unknown -emit-llvm -o - %s | FileCheck %s // REQUIRES: amdgpu-registered-target // CHECK-LABEL: define dso_local float @test_atomic_fmax_f32( diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-load.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-load.cl index 3403b69e07e4b..cfb356c684f64 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-load.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-load.cl @@ -1,6 +1,6 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu verde -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu6.01-unknown-unknown -emit-llvm -o - %s | FileCheck %s typedef unsigned char u8; typedef unsigned short u16; diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-store-format.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-store-format.cl index b10c6d59635f4..01e235c1d7024 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-store-format.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-store-format.cl @@ -1,6 +1,6 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu verde -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu6.01-unknown-unknown -emit-llvm -o - %s | FileCheck %s #pragma OPENCL EXTENSION cl_khr_fp16 : enable diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-raytracing.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-raytracing.cl index 2cf7f3dc6f80e..69b8238fd3a0f 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-raytracing.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-raytracing.cl @@ -1,11 +1,11 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1030 \ +// RUN: %clang_cc1 -triple amdgpu10.30-unknown-unknown \ // RUN: -emit-llvm -cl-std=CL2.0 -o - %s | FileCheck %s -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1030 -S \ +// RUN: %clang_cc1 -triple amdgpu10.30-unknown-unknown -S \ // RUN: -cl-std=CL2.0 -o - %s | FileCheck -check-prefix=ISA %s -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1200 -emit-llvm \ +// RUN: %clang_cc1 -triple amdgpu12.00-unknown-unknown -emit-llvm \ // RUN: -cl-std=CL2.0 -o - %s | FileCheck -check-prefix=GFX12 %s -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1200 -S \ +// RUN: %clang_cc1 -triple amdgpu12.00-unknown-unknown -S \ // RUN: -cl-std=CL2.0 -o - %s | FileCheck -check-prefix=GFX12ISA %s // Test llvm.amdgcn.image.bvh.intersect.ray intrinsic. diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-s-buffer-load.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-s-buffer-load.cl index 0c640e5020267..62e38ad756845 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-s-buffer-load.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-s-buffer-load.cl @@ -1,6 +1,6 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu verde -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu6.01-unknown-unknown -emit-llvm -o - %s | FileCheck %s #pragma OPENCL EXTENSION cl_khr_fp16 : enable @@ -23,10 +23,10 @@ typedef half v4f16 __attribute__((ext_vector_type(4))); // CHECK-LABEL: @test_amdgcn_s_buffer_load_i32( // CHECK-NEXT: entry: -// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 -// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32> -// CHECK-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.s.buffer.load.i32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0) -// CHECK-NEXT: ret i32 [[TMP0]] +// CHECK-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i128 [[TMP0]] to <4 x i32> +// CHECK-NEXT: [[TMP2:%.*]] = tail call i32 @llvm.amdgcn.s.buffer.load.i32(<4 x i32> [[TMP1]], i32 0, i32 0) +// CHECK-NEXT: ret i32 [[TMP2]] // int test_amdgcn_s_buffer_load_i32(__amdgpu_buffer_rsrc_t rsrc) { return __builtin_amdgcn_s_buffer_load_i32(rsrc, /*offset=*/0, /*aux=*/0); @@ -34,10 +34,10 @@ int test_amdgcn_s_buffer_load_i32(__amdgpu_buffer_rsrc_t rsrc) { // CHECK-LABEL: @test_amdgcn_s_buffer_load_i32_non_const_offset( // CHECK-NEXT: entry: -// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 -// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32> -// CHECK-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.s.buffer.load.i32(<4 x i32> [[RSRC_VEC:%.*]], i32 [[OFFSET:%.*]], i32 0) -// CHECK-NEXT: ret i32 [[TMP0]] +// CHECK-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i128 [[TMP0]] to <4 x i32> +// CHECK-NEXT: [[TMP2:%.*]] = tail call i32 @llvm.amdgcn.s.buffer.load.i32(<4 x i32> [[TMP1]], i32 [[OFFSET:%.*]], i32 0) +// CHECK-NEXT: ret i32 [[TMP2]] // int test_amdgcn_s_buffer_load_i32_non_const_offset(__amdgpu_buffer_rsrc_t rsrc, int offset) { return __builtin_amdgcn_s_buffer_load_i32(rsrc, offset, /*aux=*/0); @@ -45,10 +45,10 @@ int test_amdgcn_s_buffer_load_i32_non_const_offset(__amdgpu_buffer_rsrc_t rsrc, // CHECK-LABEL: @test_amdgcn_s_buffer_load_v2i32( // CHECK-NEXT: entry: -// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 -// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32> -// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.s.buffer.load.v2i32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0) -// CHECK-NEXT: ret <2 x i32> [[TMP0]] +// CHECK-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i128 [[TMP0]] to <4 x i32> +// CHECK-NEXT: [[TMP2:%.*]] = tail call <2 x i32> @llvm.amdgcn.s.buffer.load.v2i32(<4 x i32> [[TMP1]], i32 0, i32 0) +// CHECK-NEXT: ret <2 x i32> [[TMP2]] // v2i32 test_amdgcn_s_buffer_load_v2i32(__amdgpu_buffer_rsrc_t rsrc) { return __builtin_amdgcn_s_buffer_load_v2i32(rsrc, /*offset=*/0, /*aux=*/0); @@ -56,10 +56,10 @@ v2i32 test_amdgcn_s_buffer_load_v2i32(__amdgpu_buffer_rsrc_t rsrc) { // CHECK-LABEL: @test_amdgcn_s_buffer_load_v3i32( // CHECK-NEXT: entry: -// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 -// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32> -// CHECK-NEXT: [[TMP0:%.*]] = tail call <3 x i32> @llvm.amdgcn.s.buffer.load.v3i32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0) -// CHECK-NEXT: ret <3 x i32> [[TMP0]] +// CHECK-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i128 [[TMP0]] to <4 x i32> +// CHECK-NEXT: [[TMP2:%.*]] = tail call <3 x i32> @llvm.amdgcn.s.buffer.load.v3i32(<4 x i32> [[TMP1]], i32 0, i32 0) +// CHECK-NEXT: ret <3 x i32> [[TMP2]] // v3i32 test_amdgcn_s_buffer_load_v3i32(__amdgpu_buffer_rsrc_t rsrc) { return __builtin_amdgcn_s_buffer_load_v3i32(rsrc, /*offset=*/0, /*aux=*/0); @@ -67,10 +67,10 @@ v3i32 test_amdgcn_s_buffer_load_v3i32(__amdgpu_buffer_rsrc_t rsrc) { // CHECK-LABEL: @test_amdgcn_s_buffer_load_v4i32( // CHECK-NEXT: entry: -// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 -// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32> -// CHECK-NEXT: [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.s.buffer.load.v4i32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0) -// CHECK-NEXT: ret <4 x i32> [[TMP0]] +// CHECK-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i128 [[TMP0]] to <4 x i32> +// CHECK-NEXT: [[TMP2:%.*]] = tail call <4 x i32> @llvm.amdgcn.s.buffer.load.v4i32(<4 x i32> [[TMP1]], i32 0, i32 0) +// CHECK-NEXT: ret <4 x i32> [[TMP2]] // v4i32 test_amdgcn_s_buffer_load_v4i32(__amdgpu_buffer_rsrc_t rsrc) { return __builtin_amdgcn_s_buffer_load_v4i32(rsrc, /*offset=*/0, /*aux=*/0); @@ -78,10 +78,10 @@ v4i32 test_amdgcn_s_buffer_load_v4i32(__amdgpu_buffer_rsrc_t rsrc) { // CHECK-LABEL: @test_amdgcn_s_buffer_load_v8i32( // CHECK-NEXT: entry: -// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 -// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32> -// CHECK-NEXT: [[TMP0:%.*]] = tail call <8 x i32> @llvm.amdgcn.s.buffer.load.v8i32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0) -// CHECK-NEXT: ret <8 x i32> [[TMP0]] +// CHECK-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i128 [[TMP0]] to <4 x i32> +// CHECK-NEXT: [[TMP2:%.*]] = tail call <8 x i32> @llvm.amdgcn.s.buffer.load.v8i32(<4 x i32> [[TMP1]], i32 0, i32 0) +// CHECK-NEXT: ret <8 x i32> [[TMP2]] // v8i32 test_amdgcn_s_buffer_load_v8i32(__amdgpu_buffer_rsrc_t rsrc) { return __builtin_amdgcn_s_buffer_load_v8i32(rsrc, /*offset=*/0, /*aux=*/0); @@ -89,10 +89,10 @@ v8i32 test_amdgcn_s_buffer_load_v8i32(__amdgpu_buffer_rsrc_t rsrc) { // CHECK-LABEL: @test_amdgcn_s_buffer_load_v16i32( // CHECK-NEXT: entry: -// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 -// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32> -// CHECK-NEXT: [[TMP0:%.*]] = tail call <16 x i32> @llvm.amdgcn.s.buffer.load.v16i32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0) -// CHECK-NEXT: ret <16 x i32> [[TMP0]] +// CHECK-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i128 [[TMP0]] to <4 x i32> +// CHECK-NEXT: [[TMP2:%.*]] = tail call <16 x i32> @llvm.amdgcn.s.buffer.load.v16i32(<4 x i32> [[TMP1]], i32 0, i32 0) +// CHECK-NEXT: ret <16 x i32> [[TMP2]] // v16i32 test_amdgcn_s_buffer_load_v16i32(__amdgpu_buffer_rsrc_t rsrc) { return __builtin_amdgcn_s_buffer_load_v16i32(rsrc, /*offset=*/0, /*aux=*/0); @@ -100,10 +100,10 @@ v16i32 test_amdgcn_s_buffer_load_v16i32(__amdgpu_buffer_rsrc_t rsrc) { // CHECK-LABEL: @test_amdgcn_s_buffer_load_f32( // CHECK-NEXT: entry: -// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 -// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32> -// CHECK-NEXT: [[TMP0:%.*]] = tail call float @llvm.amdgcn.s.buffer.load.f32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0) -// CHECK-NEXT: ret float [[TMP0]] +// CHECK-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i128 [[TMP0]] to <4 x i32> +// CHECK-NEXT: [[TMP2:%.*]] = tail call float @llvm.amdgcn.s.buffer.load.f32(<4 x i32> [[TMP1]], i32 0, i32 0) +// CHECK-NEXT: ret float [[TMP2]] // float test_amdgcn_s_buffer_load_f32(__amdgpu_buffer_rsrc_t rsrc) { return __builtin_amdgcn_s_buffer_load_f32(rsrc, /*offset=*/0, /*aux=*/0); @@ -111,10 +111,10 @@ float test_amdgcn_s_buffer_load_f32(__amdgpu_buffer_rsrc_t rsrc) { // CHECK-LABEL: @test_amdgcn_s_buffer_load_v2f32( // CHECK-NEXT: entry: -// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 -// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32> -// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x float> @llvm.amdgcn.s.buffer.load.v2f32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0) -// CHECK-NEXT: ret <2 x float> [[TMP0]] +// CHECK-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i128 [[TMP0]] to <4 x i32> +// CHECK-NEXT: [[TMP2:%.*]] = tail call <2 x float> @llvm.amdgcn.s.buffer.load.v2f32(<4 x i32> [[TMP1]], i32 0, i32 0) +// CHECK-NEXT: ret <2 x float> [[TMP2]] // v2f32 test_amdgcn_s_buffer_load_v2f32(__amdgpu_buffer_rsrc_t rsrc) { return __builtin_amdgcn_s_buffer_load_v2f32(rsrc, /*offset=*/0, /*aux=*/0); @@ -122,10 +122,10 @@ v2f32 test_amdgcn_s_buffer_load_v2f32(__amdgpu_buffer_rsrc_t rsrc) { // CHECK-LABEL: @test_amdgcn_s_buffer_load_v3f32( // CHECK-NEXT: entry: -// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 -// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32> -// CHECK-NEXT: [[TMP0:%.*]] = tail call <3 x float> @llvm.amdgcn.s.buffer.load.v3f32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0) -// CHECK-NEXT: ret <3 x float> [[TMP0]] +// CHECK-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i128 [[TMP0]] to <4 x i32> +// CHECK-NEXT: [[TMP2:%.*]] = tail call <3 x float> @llvm.amdgcn.s.buffer.load.v3f32(<4 x i32> [[TMP1]], i32 0, i32 0) +// CHECK-NEXT: ret <3 x float> [[TMP2]] // v3f32 test_amdgcn_s_buffer_load_v3f32(__amdgpu_buffer_rsrc_t rsrc) { return __builtin_amdgcn_s_buffer_load_v3f32(rsrc, /*offset=*/0, /*aux=*/0); @@ -133,10 +133,10 @@ v3f32 test_amdgcn_s_buffer_load_v3f32(__amdgpu_buffer_rsrc_t rsrc) { // CHECK-LABEL: @test_amdgcn_s_buffer_load_v4f32( // CHECK-NEXT: entry: -// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 -// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32> -// CHECK-NEXT: [[TMP0:%.*]] = tail call <4 x float> @llvm.amdgcn.s.buffer.load.v4f32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0) -// CHECK-NEXT: ret <4 x float> [[TMP0]] +// CHECK-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i128 [[TMP0]] to <4 x i32> +// CHECK-NEXT: [[TMP2:%.*]] = tail call <4 x float> @llvm.amdgcn.s.buffer.load.v4f32(<4 x i32> [[TMP1]], i32 0, i32 0) +// CHECK-NEXT: ret <4 x float> [[TMP2]] // v4f32 test_amdgcn_s_buffer_load_v4f32(__amdgpu_buffer_rsrc_t rsrc) { return __builtin_amdgcn_s_buffer_load_v4f32(rsrc, /*offset=*/0, /*aux=*/0); @@ -144,10 +144,10 @@ v4f32 test_amdgcn_s_buffer_load_v4f32(__amdgpu_buffer_rsrc_t rsrc) { // CHECK-LABEL: @test_amdgcn_s_buffer_load_v8f32( // CHECK-NEXT: entry: -// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 -// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32> -// CHECK-NEXT: [[TMP0:%.*]] = tail call <8 x float> @llvm.amdgcn.s.buffer.load.v8f32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0) -// CHECK-NEXT: ret <8 x float> [[TMP0]] +// CHECK-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i128 [[TMP0]] to <4 x i32> +// CHECK-NEXT: [[TMP2:%.*]] = tail call <8 x float> @llvm.amdgcn.s.buffer.load.v8f32(<4 x i32> [[TMP1]], i32 0, i32 0) +// CHECK-NEXT: ret <8 x float> [[TMP2]] // v8f32 test_amdgcn_s_buffer_load_v8f32(__amdgpu_buffer_rsrc_t rsrc) { return __builtin_amdgcn_s_buffer_load_v8f32(rsrc, /*offset=*/0, /*aux=*/0); @@ -155,10 +155,10 @@ v8f32 test_amdgcn_s_buffer_load_v8f32(__amdgpu_buffer_rsrc_t rsrc) { // CHECK-LABEL: @test_amdgcn_s_buffer_load_v16f32( // CHECK-NEXT: entry: -// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 -// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32> -// CHECK-NEXT: [[TMP0:%.*]] = tail call <16 x float> @llvm.amdgcn.s.buffer.load.v16f32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0) -// CHECK-NEXT: ret <16 x float> [[TMP0]] +// CHECK-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i128 [[TMP0]] to <4 x i32> +// CHECK-NEXT: [[TMP2:%.*]] = tail call <16 x float> @llvm.amdgcn.s.buffer.load.v16f32(<4 x i32> [[TMP1]], i32 0, i32 0) +// CHECK-NEXT: ret <16 x float> [[TMP2]] // v16f32 test_amdgcn_s_buffer_load_v16f32(__amdgpu_buffer_rsrc_t rsrc) { return __builtin_amdgcn_s_buffer_load_v16f32(rsrc, /*offset=*/0, /*aux=*/0); @@ -166,10 +166,10 @@ v16f32 test_amdgcn_s_buffer_load_v16f32(__amdgpu_buffer_rsrc_t rsrc) { // CHECK-LABEL: @test_amdgcn_s_buffer_load_i8( // CHECK-NEXT: entry: -// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 -// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32> -// CHECK-NEXT: [[TMP0:%.*]] = tail call i8 @llvm.amdgcn.s.buffer.load.i8(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0) -// CHECK-NEXT: ret i8 [[TMP0]] +// CHECK-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i128 [[TMP0]] to <4 x i32> +// CHECK-NEXT: [[TMP2:%.*]] = tail call i8 @llvm.amdgcn.s.buffer.load.i8(<4 x i32> [[TMP1]], i32 0, i32 0) +// CHECK-NEXT: ret i8 [[TMP2]] // char test_amdgcn_s_buffer_load_i8(__amdgpu_buffer_rsrc_t rsrc) { return __builtin_amdgcn_s_buffer_load_i8(rsrc, /*offset=*/0, /*aux=*/0); @@ -177,10 +177,10 @@ char test_amdgcn_s_buffer_load_i8(__amdgpu_buffer_rsrc_t rsrc) { // CHECK-LABEL: @test_amdgcn_s_buffer_load_u8( // CHECK-NEXT: entry: -// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 -// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32> -// CHECK-NEXT: [[TMP0:%.*]] = tail call i8 @llvm.amdgcn.s.buffer.load.i8(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0) -// CHECK-NEXT: ret i8 [[TMP0]] +// CHECK-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i128 [[TMP0]] to <4 x i32> +// CHECK-NEXT: [[TMP2:%.*]] = tail call i8 @llvm.amdgcn.s.buffer.load.i8(<4 x i32> [[TMP1]], i32 0, i32 0) +// CHECK-NEXT: ret i8 [[TMP2]] // unsigned char test_amdgcn_s_buffer_load_u8(__amdgpu_buffer_rsrc_t rsrc) { return __builtin_amdgcn_s_buffer_load_u8(rsrc, /*offset=*/0, /*aux=*/0); @@ -188,10 +188,10 @@ unsigned char test_amdgcn_s_buffer_load_u8(__amdgpu_buffer_rsrc_t rsrc) { // CHECK-LABEL: @test_amdgcn_s_buffer_load_i16( // CHECK-NEXT: entry: -// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 -// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32> -// CHECK-NEXT: [[TMP0:%.*]] = tail call i16 @llvm.amdgcn.s.buffer.load.i16(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0) -// CHECK-NEXT: ret i16 [[TMP0]] +// CHECK-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i128 [[TMP0]] to <4 x i32> +// CHECK-NEXT: [[TMP2:%.*]] = tail call i16 @llvm.amdgcn.s.buffer.load.i16(<4 x i32> [[TMP1]], i32 0, i32 0) +// CHECK-NEXT: ret i16 [[TMP2]] // short test_amdgcn_s_buffer_load_i16(__amdgpu_buffer_rsrc_t rsrc) { return __builtin_amdgcn_s_buffer_load_i16(rsrc, /*offset=*/0, /*aux=*/0); @@ -199,10 +199,10 @@ short test_amdgcn_s_buffer_load_i16(__amdgpu_buffer_rsrc_t rsrc) { // CHECK-LABEL: @test_amdgcn_s_buffer_load_u16( // CHECK-NEXT: entry: -// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 -// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32> -// CHECK-NEXT: [[TMP0:%.*]] = tail call i16 @llvm.amdgcn.s.buffer.load.i16(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0) -// CHECK-NEXT: ret i16 [[TMP0]] +// CHECK-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i128 [[TMP0]] to <4 x i32> +// CHECK-NEXT: [[TMP2:%.*]] = tail call i16 @llvm.amdgcn.s.buffer.load.i16(<4 x i32> [[TMP1]], i32 0, i32 0) +// CHECK-NEXT: ret i16 [[TMP2]] // unsigned short test_amdgcn_s_buffer_load_u16(__amdgpu_buffer_rsrc_t rsrc) { return __builtin_amdgcn_s_buffer_load_u16(rsrc, /*offset=*/0, /*aux=*/0); @@ -210,10 +210,10 @@ unsigned short test_amdgcn_s_buffer_load_u16(__amdgpu_buffer_rsrc_t rsrc) { // CHECK-LABEL: @test_amdgcn_s_buffer_load_f16( // CHECK-NEXT: entry: -// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 -// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32> -// CHECK-NEXT: [[TMP0:%.*]] = tail call half @llvm.amdgcn.s.buffer.load.f16(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0) -// CHECK-NEXT: ret half [[TMP0]] +// CHECK-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i128 [[TMP0]] to <4 x i32> +// CHECK-NEXT: [[TMP2:%.*]] = tail call half @llvm.amdgcn.s.buffer.load.f16(<4 x i32> [[TMP1]], i32 0, i32 0) +// CHECK-NEXT: ret half [[TMP2]] // half test_amdgcn_s_buffer_load_f16(__amdgpu_buffer_rsrc_t rsrc) { return __builtin_amdgcn_s_buffer_load_f16(rsrc, /*offset=*/0, /*aux=*/0); @@ -221,10 +221,10 @@ half test_amdgcn_s_buffer_load_f16(__amdgpu_buffer_rsrc_t rsrc) { // CHECK-LABEL: @test_amdgcn_s_buffer_load_v2f16( // CHECK-NEXT: entry: -// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 -// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32> -// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x half> @llvm.amdgcn.s.buffer.load.v2f16(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0) -// CHECK-NEXT: ret <2 x half> [[TMP0]] +// CHECK-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i128 [[TMP0]] to <4 x i32> +// CHECK-NEXT: [[TMP2:%.*]] = tail call <2 x half> @llvm.amdgcn.s.buffer.load.v2f16(<4 x i32> [[TMP1]], i32 0, i32 0) +// CHECK-NEXT: ret <2 x half> [[TMP2]] // v2f16 test_amdgcn_s_buffer_load_v2f16(__amdgpu_buffer_rsrc_t rsrc) { return __builtin_amdgcn_s_buffer_load_v2f16(rsrc, /*offset=*/0, /*aux=*/0); @@ -232,10 +232,10 @@ v2f16 test_amdgcn_s_buffer_load_v2f16(__amdgpu_buffer_rsrc_t rsrc) { // CHECK-LABEL: @test_amdgcn_s_buffer_load_v3f16( // CHECK-NEXT: entry: -// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 -// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32> -// CHECK-NEXT: [[TMP0:%.*]] = tail call <3 x half> @llvm.amdgcn.s.buffer.load.v3f16(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0) -// CHECK-NEXT: ret <3 x half> [[TMP0]] +// CHECK-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i128 [[TMP0]] to <4 x i32> +// CHECK-NEXT: [[TMP2:%.*]] = tail call <3 x half> @llvm.amdgcn.s.buffer.load.v3f16(<4 x i32> [[TMP1]], i32 0, i32 0) +// CHECK-NEXT: ret <3 x half> [[TMP2]] // v3f16 test_amdgcn_s_buffer_load_v3f16(__amdgpu_buffer_rsrc_t rsrc) { return __builtin_amdgcn_s_buffer_load_v3f16(rsrc, /*offset=*/0, /*aux=*/0); @@ -243,10 +243,10 @@ v3f16 test_amdgcn_s_buffer_load_v3f16(__amdgpu_buffer_rsrc_t rsrc) { // CHECK-LABEL: @test_amdgcn_s_buffer_load_v4f16( // CHECK-NEXT: entry: -// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 -// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32> -// CHECK-NEXT: [[TMP0:%.*]] = tail call <4 x half> @llvm.amdgcn.s.buffer.load.v4f16(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0) -// CHECK-NEXT: ret <4 x half> [[TMP0]] +// CHECK-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i128 [[TMP0]] to <4 x i32> +// CHECK-NEXT: [[TMP2:%.*]] = tail call <4 x half> @llvm.amdgcn.s.buffer.load.v4f16(<4 x i32> [[TMP1]], i32 0, i32 0) +// CHECK-NEXT: ret <4 x half> [[TMP2]] // v4f16 test_amdgcn_s_buffer_load_v4f16(__amdgpu_buffer_rsrc_t rsrc) { return __builtin_amdgcn_s_buffer_load_v4f16(rsrc, /*offset=*/0, /*aux=*/0); @@ -254,10 +254,10 @@ v4f16 test_amdgcn_s_buffer_load_v4f16(__amdgpu_buffer_rsrc_t rsrc) { // CHECK-LABEL: @test_amdgcn_s_buffer_load_v2i8( // CHECK-NEXT: entry: -// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 -// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32> -// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x i8> @llvm.amdgcn.s.buffer.load.v2i8(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0) -// CHECK-NEXT: ret <2 x i8> [[TMP0]] +// CHECK-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i128 [[TMP0]] to <4 x i32> +// CHECK-NEXT: [[TMP2:%.*]] = tail call <2 x i8> @llvm.amdgcn.s.buffer.load.v2i8(<4 x i32> [[TMP1]], i32 0, i32 0) +// CHECK-NEXT: ret <2 x i8> [[TMP2]] // v2i8 test_amdgcn_s_buffer_load_v2i8(__amdgpu_buffer_rsrc_t rsrc) { return __builtin_amdgcn_s_buffer_load_v2i8(rsrc, /*offset=*/0, /*aux=*/0); @@ -265,10 +265,10 @@ v2i8 test_amdgcn_s_buffer_load_v2i8(__amdgpu_buffer_rsrc_t rsrc) { // CHECK-LABEL: @test_amdgcn_s_buffer_load_v3i8( // CHECK-NEXT: entry: -// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 -// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32> -// CHECK-NEXT: [[TMP0:%.*]] = tail call <3 x i8> @llvm.amdgcn.s.buffer.load.v3i8(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0) -// CHECK-NEXT: ret <3 x i8> [[TMP0]] +// CHECK-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i128 [[TMP0]] to <4 x i32> +// CHECK-NEXT: [[TMP2:%.*]] = tail call <3 x i8> @llvm.amdgcn.s.buffer.load.v3i8(<4 x i32> [[TMP1]], i32 0, i32 0) +// CHECK-NEXT: ret <3 x i8> [[TMP2]] // v3i8 test_amdgcn_s_buffer_load_v3i8(__amdgpu_buffer_rsrc_t rsrc) { return __builtin_amdgcn_s_buffer_load_v3i8(rsrc, /*offset=*/0, /*aux=*/0); @@ -276,10 +276,10 @@ v3i8 test_amdgcn_s_buffer_load_v3i8(__amdgpu_buffer_rsrc_t rsrc) { // CHECK-LABEL: @test_amdgcn_s_buffer_load_v4i8( // CHECK-NEXT: entry: -// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 -// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32> -// CHECK-NEXT: [[TMP0:%.*]] = tail call <4 x i8> @llvm.amdgcn.s.buffer.load.v4i8(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0) -// CHECK-NEXT: ret <4 x i8> [[TMP0]] +// CHECK-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128 +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i128 [[TMP0]] to <4 x i32> +// CHECK-NEXT: [[TMP2:%.*]] = tail call <4 x i8> @llvm.amdgcn.s.buffer.load.v4i8(<4 x i32> [[TMP1]], i32 0, i32 0) +// CHECK-NEXT: ret <4 x i8> [[TMP2]] // v4i8 test_amdgcn_s_buffer_load_v4i8(__amdgpu_buffer_rsrc_t rsrc) { return __builtin_amdgcn_s_buffer_load_v4i8(rsrc, /*offset=*/0, /*aux=*/0); diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-smfmac-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-smfmac-err.cl index 487d31539cf4d..c69a277fe0c9b 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-smfmac-err.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-smfmac-err.cl @@ -1,6 +1,6 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx908 -verify -S -o - %s -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx90a -verify -S -o - %s +// RUN: %clang_cc1 -triple amdgpu9.08-unknown-unknown -verify -S -o - %s +// RUN: %clang_cc1 -triple amdgpu9.0a-unknown-unknown -verify -S -o - %s typedef int v2i __attribute__((ext_vector_type(2))); typedef float v4f __attribute__((ext_vector_type(4))); diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-swmmac-w32-gfx10-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-swmmac-w32-gfx10-err.cl index fa4f3b5e40233..61921f80fdbfc 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-swmmac-w32-gfx10-err.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-swmmac-w32-gfx10-err.cl @@ -1,5 +1,5 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1030 \ +// RUN: %clang_cc1 -triple amdgpu10.30-unknown-unknown \ // RUN: -verify -S -o - %s typedef int v2i __attribute__((ext_vector_type(2))); diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-swmmac-w64-gfx10-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-swmmac-w64-gfx10-err.cl index 4cde9d78abf2f..21bb0c7ab9a03 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-swmmac-w64-gfx10-err.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-swmmac-w64-gfx10-err.cl @@ -1,5 +1,5 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1030 \ +// RUN: %clang_cc1 -triple amdgpu10.30-unknown-unknown \ // RUN: -verify -S -o - %s typedef int v2i __attribute__((ext_vector_type(2))); diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-swmmac-w64.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-swmmac-w64.cl index 3928b77ad0689..01cf6d0ce8580 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-swmmac-w64.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-swmmac-w64.cl @@ -1,6 +1,6 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1200 -target-feature +wavefrontsize64 -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-GFX1200 +// RUN: %clang_cc1 -triple amdgpu12.00-unknown-unknown -target-feature +wavefrontsize64 -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-GFX1200 typedef int v2i __attribute__((ext_vector_type(2))); typedef int v4i __attribute__((ext_vector_type(4))); diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl index 758b5aa532d73..7d52aaafa7d1c 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl @@ -1,8 +1,8 @@ // RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-feature +wavefrontsize64 -emit-llvm -o - %s | FileCheck -enable-var-scope %s -// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu gfx900 -emit-llvm -o - %s | FileCheck -enable-var-scope %s -// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu gfx900 -target-feature +wavefrontsize64 -emit-llvm -o - %s | FileCheck -enable-var-scope %s -// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -target-feature +wavefrontsize64 -emit-llvm -o - %s | FileCheck -enable-var-scope %s -// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -target-feature +wavefrontsize64 -emit-llvm -o - %s | FileCheck -enable-var-scope %s +// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgpu9.00-unknown-unknown -emit-llvm -o - %s | FileCheck -enable-var-scope %s +// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgpu9.00-unknown-unknown -target-feature +wavefrontsize64 -emit-llvm -o - %s | FileCheck -enable-var-scope %s +// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgpu10.10-unknown-unknown -target-feature +wavefrontsize64 -emit-llvm -o - %s | FileCheck -enable-var-scope %s +// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgpu11.00-unknown-unknown -target-feature +wavefrontsize64 -emit-llvm -o - %s | FileCheck -enable-var-scope %s typedef unsigned long ulong; diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-wmma-w32-gfx10-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-wmma-w32-gfx10-err.cl index ed72a8ee7dbd2..94801c47b314a 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-wmma-w32-gfx10-err.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-wmma-w32-gfx10-err.cl @@ -1,6 +1,6 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1030 \ +// RUN: %clang_cc1 -triple amdgpu10.30-unknown-unknown \ // RUN: -verify -DWMMA_GFX1100_TESTS -S -o - %s diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-wmma-w64.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-wmma-w64.cl index 89909e1ef9e33..3337634e68ac2 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-wmma-w64.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-wmma-w64.cl @@ -1,6 +1,6 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -target-feature +wavefrontsize64 -DWMMA_GFX1100_TESTS -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-GFX1100 +// RUN: %clang_cc1 -triple amdgpu11.00-unknown-unknown -target-feature +wavefrontsize64 -DWMMA_GFX1100_TESTS -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-GFX1100 typedef float v4f __attribute__((ext_vector_type(4))); typedef half v8h __attribute__((ext_vector_type(8))); diff --git a/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx1030.cl b/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx1030.cl index 14efaa74d90fd..6dfb2c7f8e4e8 100644 --- a/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx1030.cl +++ b/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx1030.cl @@ -1,7 +1,7 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx1030 \ +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgpu10.30-amd-amdhsa \ // RUN: -S -o - %s -// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx1030 \ +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgpu10.30-amd-amdhsa \ // RUN: -S -o - %s | FileCheck -check-prefix=GFX1030 %s // CHECK-LABEL: test_ds_addf_local diff --git a/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx12.cl b/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx12.cl index 461abc3708128..57a91b07f3535 100644 --- a/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx12.cl +++ b/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx12.cl @@ -1,7 +1,7 @@ -// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx1200 \ +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgpu12.00-amd-amdhsa \ // RUN: %s -emit-llvm -o - | FileCheck %s -// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx1200 \ +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgpu12.00-amd-amdhsa \ // RUN: -S -o - %s | FileCheck -check-prefix=GFX12 %s // REQUIRES: amdgpu-registered-target diff --git a/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx8.cl b/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx8.cl index 2728490d5c02e..d3cf134098300 100644 --- a/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx8.cl +++ b/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx8.cl @@ -1,6 +1,6 @@ -// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx810 \ +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgpu8.10-amd-amdhsa \ // RUN: %s -emit-llvm -o - | FileCheck %s -// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx810 \ +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgpu8.10-amd-amdhsa \ // RUN: -S -o - %s | FileCheck -check-prefix=GFX8 %s // REQUIRES: amdgpu-registered-target diff --git a/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx90a.cl b/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx90a.cl index ca8ddb3951e85..c40c67089d384 100644 --- a/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx90a.cl +++ b/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx90a.cl @@ -1,7 +1,7 @@ -// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \ +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgpu9.0a-amd-amdhsa \ // RUN: %s -emit-llvm -o - | FileCheck %s -check-prefix=CHECK -// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \ +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgpu9.0a-amd-amdhsa \ // RUN: -S -o - %s | FileCheck -check-prefix=GFX90A %s // REQUIRES: amdgpu-registered-target diff --git a/clang/test/CodeGenOpenCL/check-atomic-alignment.cl b/clang/test/CodeGenOpenCL/check-atomic-alignment.cl index 66753052cba26..51889d6547700 100644 --- a/clang/test/CodeGenOpenCL/check-atomic-alignment.cl +++ b/clang/test/CodeGenOpenCL/check-atomic-alignment.cl @@ -1,5 +1,5 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5 -// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx942 \ +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgpu9.42-amd-amdhsa \ // RUN: %s -emit-llvm -o - -disable-llvm-passes | FileCheck %s // REQUIRES: amdgpu-registered-target _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
