Cold-Air-Networks/mirror-llvm-project
2
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions 7

clang/AMDGPU: Rename amdgcn triples to amdgpu in cc1 tests (14) (#212053)

Browse source 
Replace the legacy amdgcn architecture name with amdgpu in various
test triples. This set of updates covers cases that did not have
a set -target-cpu.

Co-authored-by: Claude (Opus 4.8) <noreply@anthropic.com>
This commit is contained in:
Matt Arsenault 2026-07-27 10:45:30 +02:00 committed by GitHub
parent d211ecfc81
commit 3b6e754c0a
No known key found for this signature in database
GPG key ID: B5690EEEBB952194
267 changed files with 485 additions and 495 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

4
clang/test/AST/ByteCode/amdgpu-nullptr.cl
View file

@ -1,6 +1,6 @@
// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL2.0 -triple amdgcn -emit-llvm -o - | FileCheck %s
// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL2.0 -triple amdgpu -emit-llvm -o - | FileCheck %s
// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL2.0 -triple amdgcn -emit-llvm -fexperimental-new-constant-interpreter -o - | FileCheck %s
// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL2.0 -triple amdgpu -emit-llvm -fexperimental-new-constant-interpreter -o - | FileCheck %s
// CHECK: @fold_priv ={{.*}} local_unnamed_addr addrspace(1) global ptr addrspace(5) addrspacecast (ptr addrspace(1) null to ptr addrspace(5)), align 4

10
clang/test/AST/ast-dump-amdgpu-types.c
View file

@ -1,12 +1,12 @@
// REQUIRES: amdgpu-registered-target
// Test without serialization:
// RUN: %clang_cc1 -triple amdgcn -ast-dump -ast-dump-filter __amdgpu_buffer_rsrc_t %s | FileCheck %s -check-prefix=BUFFER-RSRC
// RUN: %clang_cc1 -triple amdgcn -ast-dump -ast-dump-filter __amdgpu_named_workgroup_barrier %s | FileCheck %s -check-prefix=WORKGROUP-BARRIER
// RUN: %clang_cc1 -triple amdgpu -ast-dump -ast-dump-filter __amdgpu_buffer_rsrc_t %s | FileCheck %s -check-prefix=BUFFER-RSRC
// RUN: %clang_cc1 -triple amdgpu -ast-dump -ast-dump-filter __amdgpu_named_workgroup_barrier %s | FileCheck %s -check-prefix=WORKGROUP-BARRIER
//
// Test with serialization:
// RUN: %clang_cc1 -triple amdgcn -emit-pch -o %t %s
// RUN: %clang_cc1 -x c -triple amdgcn -include-pch %t -ast-dump-all -ast-dump-filter __amdgpu_buffer_rsrc_t /dev/null | sed -e "s/ <undeserialized declarations>//" -e "s/ imported//" | FileCheck %s -check-prefix=BUFFER-RSRC
// RUN: %clang_cc1 -x c -triple amdgcn -include-pch %t -ast-dump-all -ast-dump-filter __amdgpu_named_workgroup_barrier /dev/null | sed -e "s/ <undeserialized declarations>//" -e "s/ imported//" | FileCheck %s -check-prefix=WORKGROUP-BARRIER
// RUN: %clang_cc1 -triple amdgpu -emit-pch -o %t %s
// RUN: %clang_cc1 -x c -triple amdgpu -include-pch %t -ast-dump-all -ast-dump-filter __amdgpu_buffer_rsrc_t /dev/null | sed -e "s/ <undeserialized declarations>//" -e "s/ imported//" | FileCheck %s -check-prefix=BUFFER-RSRC
// RUN: %clang_cc1 -x c -triple amdgpu -include-pch %t -ast-dump-all -ast-dump-filter __amdgpu_named_workgroup_barrier /dev/null | sed -e "s/ <undeserialized declarations>//" -e "s/ imported//" | FileCheck %s -check-prefix=WORKGROUP-BARRIER
// BUFFER-RSRC: TypedefDecl {{.*}} implicit __amdgpu_buffer_rsrc_t
// BUFFER-RSRC-NEXT: -BuiltinType {{.*}} '__amdgpu_buffer_rsrc_t'

4
clang/test/Analysis/addrspace-null.c
View file

@ -1,9 +1,9 @@
// RUN: %clang_analyze_cc1 -triple amdgcn-unknown-unknown \
// RUN: %clang_analyze_cc1 -triple amdgpu-unknown-unknown \
// RUN: -analyzer-checker=core -DAMDGCN_TRIPLE \
// RUN: -analyzer-checker=debug.ExprInspection \
// RUN: -Wno-implicit-int -Wno-int-conversion -verify %s
//
// RUN: %clang_analyze_cc1 -triple amdgcn-unknown-unknown \
// RUN: %clang_analyze_cc1 -triple amdgpu-unknown-unknown \
// RUN: -analyzer-checker=core -DDEFAULT_TRIPLE \
// RUN: -analyzer-checker=debug.ExprInspection \
// RUN: -Wno-implicit-int -Wno-int-conversion -verify %s

6
clang/test/Analysis/cast-value-notes.cpp
View file

@ -1,13 +1,13 @@
// RUN: %clang_analyze_cc1 -std=c++14 -triple amdgcn-unknown-unknown \
// RUN: %clang_analyze_cc1 -std=c++14 -triple amdgpu-unknown-unknown \
// RUN: -analyzer-checker=core,apiModeling.llvm.CastValue,debug.ExprInspection\
// RUN: -analyzer-output=text -verify -DX86 -DSUPPRESSED %s 2>&1 | FileCheck %s -check-prefix=X86-CHECK
//
// RUN: %clang_analyze_cc1 -std=c++14 -triple amdgcn-unknown-unknown \
// RUN: %clang_analyze_cc1 -std=c++14 -triple amdgpu-unknown-unknown \
// RUN: -analyzer-checker=core,apiModeling.llvm.CastValue,debug.ExprInspection\
// RUN: -analyzer-config suppress-dereferences-from-any-address-space=false\
// RUN: -analyzer-output=text -verify -DX86 -DNOT_SUPPRESSED %s 2>&1 | FileCheck %s -check-prefix=X86-CHECK
//
// RUN: %clang_analyze_cc1 -std=c++14 -triple amdgcn-unknown-unknown \
// RUN: %clang_analyze_cc1 -std=c++14 -triple amdgpu-unknown-unknown \
// RUN: -analyzer-checker=core,apiModeling.llvm.CastValue,debug.ExprInspection\
// RUN: -analyzer-config suppress-dereferences-from-any-address-space=true\
// RUN: -analyzer-output=text -verify -DX86 -DSUPPRESSED %s 2>&1 | FileCheck %s -check-prefix=X86-CHECK

2
clang/test/Analysis/element-region-address-space.c
View file

@ -1,4 +1,4 @@
// RUN: %clang_analyze_cc1 -triple amdgcn-unknown-unknown \
// RUN: %clang_analyze_cc1 -triple amdgpu-unknown-unknown \
// RUN: -Wno-incompatible-library-redeclaration \
// RUN: -analyzer-checker=core,unix -verify %s

2
clang/test/Analysis/ptr-cmp-const-trunc.cl
View file

@ -1,4 +1,4 @@
//RUN: %clang_analyze_cc1 -triple amdgcn-unknown-unknown -analyzer-checker=core -verify %s
//RUN: %clang_analyze_cc1 -triple amdgpu-unknown-unknown -analyzer-checker=core -verify %s
// expected-no-diagnostics
#include <stdint.h>

2
clang/test/Analysis/solver-sym-simplification-ptr-bool.cl
View file

@ -1,4 +1,4 @@
// RUN: %clang_analyze_cc1 -triple amdgcn-unknown-unknown -analyzer-checker=core %s
// RUN: %clang_analyze_cc1 -triple amdgpu-unknown-unknown -analyzer-checker=core %s
// expected-no-diagnostics

2
clang/test/Analysis/z3/cstring-addrspace.c
View file

@ -1,4 +1,4 @@
// RUN: %clang_analyze_cc1 -triple amdgcn-unknown-unknown \
// RUN: %clang_analyze_cc1 -triple amdgpu-unknown-unknown \
// RUN: -analyzer-checker=core,alpha.unix.cstring,debug.ExprInspection \
// RUN: -analyzer-config crosscheck-with-z3=true -verify %s \
// RUN: -Wno-incompatible-library-redeclaration

6
clang/test/CIR/CodeGen/amdgpu-address-spaces.cpp
View file

@ -1,13 +1,13 @@
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fclangir -emit-cir \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fclangir -emit-cir \
// RUN: -mmlir -mlir-print-ir-before=cir-target-lowering %s -o %t.cir 2> %t-pre.cir
// RUN: FileCheck --check-prefix=CIR-PRE --input-file=%t-pre.cir %s
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fclangir -emit-llvm %s -o %t.ll
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fclangir -emit-llvm %s -o %t.ll
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm %s -o %t.ll
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -emit-llvm %s -o %t.ll
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
// Test address space handling for AMDGPU target in C++ mode (non-OpenCL/HIP).

6
clang/test/CIR/CodeGen/amdgpu-array-addrspace.cpp
View file

@ -1,8 +1,8 @@
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fclangir -emit-cir %s -o %t.cir
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fclangir -emit-cir %s -o %t.cir
// RUN: FileCheck --input-file=%t.cir %s -check-prefix=CIR
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fclangir -emit-llvm %s -o %t.ll
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fclangir -emit-llvm %s -o %t.ll
// RUN: FileCheck --input-file=%t.ll %s -check-prefix=LLVM
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm %s -o %t.ll
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -emit-llvm %s -o %t.ll
// RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG
// Test that address spaces are preserved through array-to-pointer decay

6
clang/test/CIR/CodeGen/amdgpu-call-addrspace-cast.cpp
View file

@ -1,8 +1,8 @@
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fclangir -emit-cir %s -o %t.cir
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fclangir -emit-cir %s -o %t.cir
// RUN: FileCheck --input-file=%t.cir %s -check-prefix=CIR
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fclangir -emit-llvm %s -o %t.ll
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fclangir -emit-llvm %s -o %t.ll
// RUN: FileCheck --input-file=%t.ll %s -check-prefix=LLVM
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm %s -o %t.ll
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -emit-llvm %s -o %t.ll
// RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG
// Test that address space casts are inserted when passing pointers in

6
clang/test/CIR/CodeGen/amdgpu-stack-alloca-array-decay.cpp
View file

@ -1,12 +1,12 @@
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa \
// RUN: -fclangir -emit-cir %s -o %t.cir
// RUN: FileCheck --input-file=%t.cir %s --check-prefix=CIR
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa \
// RUN: -fclangir -emit-llvm %s -o %t.ll
// RUN: FileCheck --input-file=%t.ll %s --check-prefix=LLVM
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa \
// RUN: -emit-llvm %s -o %t.ll
// RUN: FileCheck --input-file=%t.ll %s --check-prefix=OGCG

4
clang/test/CIR/CodeGen/amdgpu-target-lowering-as.cpp
View file

@ -10,11 +10,11 @@
// offload_constant -> 4 (CONSTANT_ADDRESS)
// offload_private -> 5 (PRIVATE_ADDRESS)
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fclangir -emit-cir \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fclangir -emit-cir \
// RUN: -mmlir -mlir-print-ir-before=cir-target-lowering %s -o %t.cir 2> %t.pre.cir
// RUN: FileCheck --check-prefix=PRE --input-file=%t.pre.cir %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fclangir -emit-cir %s -o %t.cir
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fclangir -emit-cir %s -o %t.cir
// RUN: FileCheck --check-prefix=POST --input-file=%t.cir %s
// ---- GlobalOp: lang AS emitted before target lowering, target AS after ----

8
clang/test/CIR/CodeGenHIP/amdgpu-attrs.hip
View file

@ -1,22 +1,22 @@
#include "../CodeGenCUDA/Inputs/cuda.h"
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \
// RUN: %clang_cc1 -triple=amdgpu-amd-amdhsa -x hip -fclangir \
// RUN: -Wno-deprecated-declarations \
// RUN: -fcuda-is-device -emit-cir %s -o %t.cir
// RUN: FileCheck --check-prefix=CIR %s --input-file=%t.cir
// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \
// RUN: %clang_cc1 -triple=amdgpu-amd-amdhsa -x hip -fclangir \
// RUN: -Wno-deprecated-declarations \
// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll
// RUN: FileCheck --check-prefix=LLVM %s --input-file=%t.ll
// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip \
// RUN: %clang_cc1 -triple=amdgpu-amd-amdhsa -x hip \
// RUN: -Wno-deprecated-declarations \
// RUN: -fcuda-is-device -emit-llvm %s -o %t.ogcg.ll
// RUN: FileCheck --check-prefix=LLVM %s --input-file=%t.ogcg.ll
// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \
// RUN: %clang_cc1 -triple=amdgpu-amd-amdhsa -x hip -fclangir \
// RUN: -Wno-deprecated-declarations \
// RUN: -fcuda-is-device -mamdgpu-expand-waitcnt-profiling \
// RUN: -emit-cir %s -o %t.prof.cir

6
clang/test/CIR/CodeGenHIP/amdgpu-module-flags.hip
View file

@ -1,15 +1,15 @@
#include "../CodeGenCUDA/Inputs/cuda.h"
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \
// RUN: %clang_cc1 -triple=amdgpu-amd-amdhsa -x hip -fclangir \
// RUN: -fcuda-is-device -emit-cir %s -o %t.cir
// RUN: FileCheck --check-prefix=CIR %s --input-file=%t.cir
// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \
// RUN: %clang_cc1 -triple=amdgpu-amd-amdhsa -x hip -fclangir \
// RUN: -fcuda-is-device -emit-llvm %s -o %t.cir.ll
// RUN: FileCheck --check-prefix=LLVM %s --input-file=%t.cir.ll
// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip \
// RUN: %clang_cc1 -triple=amdgpu-amd-amdhsa -x hip \
// RUN: -fcuda-is-device -emit-llvm %s -o %t.ogcg.ll
// RUN: FileCheck --check-prefix=LLVM %s --input-file=%t.ogcg.ll

6
clang/test/CIR/CodeGenHIP/hip-cuid.hip
View file

@ -1,15 +1,15 @@
#include "../CodeGenCUDA/Inputs/cuda.h"
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \
// RUN: %clang_cc1 -triple=amdgpu-amd-amdhsa -x hip -fclangir \
// RUN: -fcuda-is-device -emit-cir %s -o %t.cir
// RUN: FileCheck --check-prefix=CIR %s --input-file=%t.cir
// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \
// RUN: %clang_cc1 -triple=amdgpu-amd-amdhsa -x hip -fclangir \
// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll
// RUN: FileCheck --check-prefix=LLVM %s --input-file=%t.ll
// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip \
// RUN: %clang_cc1 -triple=amdgpu-amd-amdhsa -x hip \
// RUN: -fcuda-is-device -emit-llvm %s -o %t.ogcg.ll
// RUN: FileCheck --check-prefix=OGCG %s --input-file=%t.ogcg.ll

4
clang/test/CIR/CodeGenHIP/simple.cpp
View file

@ -5,7 +5,7 @@
// RUN: -I%S/../CodeGenCUDA/Inputs/ -emit-cir %s -o %t.cir
// RUN: FileCheck --check-prefix=CIR-HOST --input-file=%t.cir %s
// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \
// RUN: %clang_cc1 -triple=amdgpu-amd-amdhsa -x hip -fclangir \
// RUN: -fcuda-is-device -fhip-new-launch-api \
// RUN: -I%S/../CodeGenCUDA/Inputs/ -emit-cir %s -o %t.cir
// RUN: FileCheck --check-prefix=CIR-DEVICE --input-file=%t.cir %s
@ -15,7 +15,7 @@
// RUN: -I%S/../CodeGenCUDA/Inputs/ %s -o %t.ll
// RUN: FileCheck --check-prefix=OGCG-HOST --input-file=%t.ll %s
// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip \
// RUN: %clang_cc1 -triple=amdgpu-amd-amdhsa -x hip \
// RUN: -fcuda-is-device -fhip-new-launch-api \
// RUN: -fvisibility=hidden \
// RUN: -I%S/../CodeGenCUDA/Inputs/ -emit-llvm %s -o %t.ll

6
clang/test/CIR/CodeGenOpenMP/emit-device-functions.cpp
View file

@ -1,10 +1,10 @@
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fopenmp -fopenmp-is-target-device \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fopenmp -fopenmp-is-target-device \
// RUN: -fclangir -emit-cir %s -o %t.cir
// RUN: FileCheck --input-file=%t.cir %s -check-prefix=CIR
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fopenmp -fopenmp-is-target-device \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fopenmp -fopenmp-is-target-device \
// RUN: -fclangir -emit-llvm %s -o %t.ll
// RUN: FileCheck --input-file=%t.ll %s -check-prefix=LLVM
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fopenmp -fopenmp-is-target-device \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fopenmp -fopenmp-is-target-device \
// RUN: -emit-llvm %s -o %t.ll
// RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG

8
clang/test/CIR/CodeGenOpenMP/omp-module-attrs.c
View file

@ -11,16 +11,16 @@
// Host with target triples
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fopenmp -fclangir -emit-cir \
// RUN: -fopenmp-targets=amdgcn-amd-amdhsa,nvptx64-nvidia-cuda %s -o - \
// RUN: -fopenmp-targets=amdgpu-amd-amdhsa,nvptx64-nvidia-cuda %s -o - \
// RUN: | FileCheck %s --check-prefix=HOST-TRIPLES
// HOST-TRIPLES: module {{.*}} attributes {
// HOST-TRIPLES-SAME: omp.is_gpu = false
// HOST-TRIPLES-SAME: omp.is_target_device = false
// HOST-TRIPLES-SAME: omp.target_triples = ["amdgcn-amd-amdhsa", "nvptx64-nvidia-cuda"]
// HOST-TRIPLES-SAME: omp.target_triples = ["amdgpu-amd-amdhsa", "nvptx64-nvidia-cuda"]
// Device, AMDGPU
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fopenmp -fclangir -emit-cir \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fopenmp -fclangir -emit-cir \
// RUN: -fopenmp-is-target-device %s -o - \
// RUN: | FileCheck %s --check-prefix=AMDGPU-DEVICE
@ -47,7 +47,7 @@
// CPU-DEVICE-SAME: omp.is_target_device = true
// Device with omp.flags
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fopenmp -fclangir -emit-cir \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fopenmp -fclangir -emit-cir \
// RUN: -fopenmp-is-target-device \
// RUN: -fopenmp-assume-no-thread-state \
// RUN: -fopenmp-assume-no-nested-parallelism %s -o - \

2
clang/test/CodeGen/AMDGPU/sanitizer.c
View file

@ -1,5 +1,5 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
// RUN: %clang_cc1 -cc1 -triple amdgcn-amd-amdhsa -emit-llvm -fsanitize=function %s -o - | FileCheck %s
// RUN: %clang_cc1 -cc1 -triple amdgpu-amd-amdhsa -emit-llvm -fsanitize=function %s -o - | FileCheck %s
// CHECK-LABEL: define dso_local void @function(
// CHECK-SAME: ) #[[ATTR0:[0-9]+]] !func_sanitize [[META2:![0-9]+]] {

2
clang/test/CodeGen/address-space.c
View file

@ -1,5 +1,5 @@
// RUN: %clang_cc1 -triple x86_64-apple-darwin -emit-llvm < %s | FileCheck -enable-var-scope -check-prefixes=CHECK,X86 %s
// RUN: %clang_cc1 -triple amdgcn -emit-llvm < %s | FileCheck -enable-var-scope -check-prefixes=CHECK,AMDGCN %s
// RUN: %clang_cc1 -triple amdgpu -emit-llvm < %s | FileCheck -enable-var-scope -check-prefixes=CHECK,AMDGCN %s
// CHECK: @foo ={{.*}} addrspace(1) global
int foo __attribute__((address_space(1)));

4
clang/test/CodeGen/amdgpu-abi-version.c
View file

@ -1,6 +1,6 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals all --version 5
// RUN: %clang_cc1 -cc1 -triple amdgcn-amd-amdhsa -emit-llvm -mcode-object-version=none %s -o - | FileCheck -check-prefixes=CHECK,LLVM %s
// RUN: %clang_cc1 -cc1 -triple amdgcn-amd-amdhsa-llvm -emit-llvm -mcode-object-version=none %s -o - | FileCheck -check-prefixes=CHECK,LLVMENV %s
// RUN: %clang_cc1 -cc1 -triple amdgpu-amd-amdhsa -emit-llvm -mcode-object-version=none %s -o - | FileCheck -check-prefixes=CHECK,LLVM %s
// RUN: %clang_cc1 -cc1 -triple amdgpu-amd-amdhsa-llvm -emit-llvm -mcode-object-version=none %s -o - | FileCheck -check-prefixes=CHECK,LLVMENV %s
//.
// LLVM: @__oclc_ABI_version = external addrspace(4) global i32

2
clang/test/CodeGen/amdgpu-address-spaces.cpp
View file

@ -1,5 +1,5 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals --version 3
// RUN: %clang_cc1 -cc1 -triple amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s
// RUN: %clang_cc1 -cc1 -triple amdgpu-amd-amdhsa -emit-llvm %s -o - | FileCheck %s
int [[clang::opencl_global]] a = 100;
int [[clang::opencl_generic]] b = 42;

2
clang/test/CodeGen/amdgpu-image-rsrc-type-debug-info.c
View file

@ -1,6 +1,6 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn -emit-llvm -o - %s -debug-info-kind=limited | FileCheck %s
// RUN: %clang_cc1 -triple amdgpu -emit-llvm -o - %s -debug-info-kind=limited | FileCheck %s
// CHECK-LABEL: define dso_local void @test_locals(
// CHECK-SAME: ) #[[ATTR0:[0-9]+]] !dbg [[DBG5:![0-9]+]] {

2
clang/test/CodeGen/amdgpu-variadic-call.c
View file

@ -1,6 +1,6 @@
// REQUIRES: amdgpu-registered-target
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature
// RUN: %clang_cc1 -cc1 -std=c23 -triple amdgcn-amd-amdhsa -emit-llvm -O1 %s -o - | FileCheck %s
// RUN: %clang_cc1 -cc1 -std=c23 -triple amdgpu-amd-amdhsa -emit-llvm -O1 %s -o - | FileCheck %s
void sink_0(...);
void sink_1(int, ...);

2
clang/test/CodeGen/annotations-field.c
View file

@ -1,5 +1,5 @@
// RUN: %clang_cc1 -triple x86_64 -emit-llvm -o - %s | FileCheck %s '-D$MANGLE_AS=p0' '-D$CONST_AS=' --check-prefixes=CHECK,X86
// RUN: %clang_cc1 -triple amdgcn -emit-llvm -o - %s | FileCheck %s '-D$MANGLE_AS=p4' '-D$CONST_AS= addrspace(4)' --check-prefixes=CHECK,AMDGPU
// RUN: %clang_cc1 -triple amdgpu -emit-llvm -o - %s | FileCheck %s '-D$MANGLE_AS=p4' '-D$CONST_AS= addrspace(4)' --check-prefixes=CHECK,AMDGPU
// END.
// CHECK: private unnamed_addr[[$CONST_AS]] constant [8 x i8] c"v_ann_{{.}}\00", section "llvm.metadata"

2
clang/test/CodeGen/annotations-loc.c
View file

@ -1,5 +1,5 @@
// RUN: %clang_cc1 -triple x86_64 -emit-llvm -o - %s | FileCheck %s '-D$CONST_AS='
// RUN: %clang_cc1 -triple amdgcn -emit-llvm -o - %s | FileCheck %s '-D$CONST_AS= addrspace(4)'
// RUN: %clang_cc1 -triple amdgpu -emit-llvm -o - %s | FileCheck %s '-D$CONST_AS= addrspace(4)'
// END.
# 1 "t.c"
# 1 "<built-in>"

2
clang/test/CodeGen/bitfield-access-unit.c
View file

@ -27,7 +27,7 @@
// Configs that have expensive unaligned access
// 64-bit Little Endian
// RUN: %clang_cc1 -triple=aarch64-linux-gnu %s -target-feature +strict-align -emit-llvm -o /dev/null -fdump-record-layouts-simple | FileCheck --check-prefixes CHECK,LAYOUT,LAYOUT-STRICT,CHECK-64,LAYOUT-64,LAYOUT-64-STRICT %s
// RUN: %clang_cc1 -triple=amdgcn-elf %s -emit-llvm -o /dev/null -fdump-record-layouts-simple | FileCheck --check-prefixes CHECK,LAYOUT,LAYOUT-STRICT,CHECK-64,LAYOUT-64,LAYOUT-64-STRICT %s
// RUN: %clang_cc1 -triple=amdgpu-elf %s -emit-llvm -o /dev/null -fdump-record-layouts-simple | FileCheck --check-prefixes CHECK,LAYOUT,LAYOUT-STRICT,CHECK-64,LAYOUT-64,LAYOUT-64-STRICT %s
// RUN: %clang_cc1 -triple=loongarch64-elf -target-feature -ual %s -emit-llvm -o /dev/null -fdump-record-layouts-simple | FileCheck --check-prefixes CHECK,LAYOUT,LAYOUT-STRICT,CHECK-64,LAYOUT-64,LAYOUT-64-STRICT %s
// RUN: %clang_cc1 -triple=riscv64 %s -emit-llvm -o /dev/null -fdump-record-layouts-simple | FileCheck --check-prefixes CHECK,LAYOUT,LAYOUT-STRICT,CHECK-64,LAYOUT-64,LAYOUT-64-STRICT %s

2
clang/test/CodeGen/default-address-space.c
View file

@ -1,4 +1,4 @@
// RUN: %clang_cc1 -triple amdgcn---amdgiz -emit-llvm < %s | FileCheck -check-prefixes=CHECK %s
// RUN: %clang_cc1 -triple amdgpu---amdgiz -emit-llvm < %s | FileCheck -check-prefixes=CHECK %s
// CHECK-DAG: @foo ={{.*}} addrspace(1) global i32 0
int foo;

6
clang/test/CodeGen/dynamic-alloca-with-address-space.c
View file

@ -1,8 +1,8 @@
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm %s -o - \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -emit-llvm %s -o - \
// RUN: | FileCheck %s --check-prefix=CHECK
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -DOCL12 -x cl -std=cl1.2 \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -DOCL12 -x cl -std=cl1.2 \
// RUN: -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK-CL12
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x cl -std=cl2.0 \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -x cl -std=cl2.0 \
// RUN: -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK-CL20
#if defined(OCL12)

2
clang/test/CodeGen/fp-atomic-ops.c
View file

@ -1,4 +1,4 @@
// RUN: %clang_cc1 %s -emit-llvm -DDOUBLE -O0 -o - -triple=amdgcn-amd-amdhsa \
// RUN: %clang_cc1 %s -emit-llvm -DDOUBLE -O0 -o - -triple=amdgpu-amd-amdhsa \
// RUN: | FileCheck -check-prefixes=FLOAT,DOUBLE %s
// RUN: %clang_cc1 %s -emit-llvm -DDOUBLE -O0 -o - -triple=aarch64-linux-gnu \

20
clang/test/CodeGen/link-builtin-bitcode.c
View file

@ -1,13 +1,13 @@
// Build two version of the bitcode library, one with a target-cpu set and one without
// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx803 -DBITCODE -emit-llvm-bc -o %t-lib.bc %s
// RUN: %clang_cc1 -triple amdgcn-- -DBITCODE -emit-llvm-bc -o %t-lib.no-cpu.bc %s
// RUN: %clang_cc1 -triple amdgpu8.03-- -DBITCODE -emit-llvm-bc -o %t-lib.bc %s
// RUN: %clang_cc1 -triple amdgpu-- -DBITCODE -emit-llvm-bc -o %t-lib.no-cpu.bc %s
// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx90a -emit-llvm-bc -o %t.bc %s
// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx90a -emit-llvm \
// RUN: %clang_cc1 -triple amdgpu9.0a-- -emit-llvm-bc -o %t.bc %s
// RUN: %clang_cc1 -triple amdgpu9.0a-- -emit-llvm \
// RUN: -mlink-builtin-bitcode %t-lib.bc -o - %t.bc | FileCheck %s
// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx90a -emit-llvm-bc -o %t.bc %s
// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx90a -emit-llvm \
// RUN: %clang_cc1 -triple amdgpu9.0a-- -emit-llvm-bc -o %t.bc %s
// RUN: %clang_cc1 -triple amdgpu9.0a-- -emit-llvm \
// RUN: -mlink-builtin-bitcode %t-lib.no-cpu.bc -o - %t.bc | FileCheck %s
#ifdef BITCODE
@ -43,7 +43,7 @@ int bar() { return no_attr() + attr_in_target() + attr_not_in_target() + attr_in
// CHECK-LABEL: @attr_incompatible
// CHECK-SAME: () #[[ATTR_INCOMPATIBLE:[0-9]+]] {
// CHECK: attributes #[[ATTR_BAR]] = { {{.*}} "target-cpu"="gfx90a" }
// CHECK: attributes #[[ATTR_COMPATIBLE]] = { {{.*}} "target-cpu"="gfx90a" "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-global-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64" }
// CHECK: attributes #[[ATTR_EXTEND]] = { {{.*}} "target-cpu"="gfx90a" "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+extended-image-insts,+flat-global-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64" }
// CHECK: attributes #[[ATTR_INCOMPATIBLE]] = { {{.*}} "target-cpu"="gfx90a" "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-global-insts,+gfx8-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64,-gfx9-insts" }
// CHECK: attributes #[[ATTR_BAR]] = { {{.*}} "no-trapping-math"="true" {{.*}} }
// CHECK: attributes #[[ATTR_COMPATIBLE]] = { {{.*}} "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-global-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64" }
// CHECK: attributes #[[ATTR_EXTEND]] = { {{.*}} "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+extended-image-insts,+flat-global-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64" }
// CHECK: attributes #[[ATTR_INCOMPATIBLE]] = { {{.*}} "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-global-insts,+gfx8-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64,-gfx9-insts" }

8
clang/test/CodeGen/linking-bitcode-postopt.cpp
View file

@ -2,27 +2,27 @@
// Test that -mlink-bitcode-postopt correctly enables LinkInModulesPass
// RUN: %clang_cc1 -triple amdgcn-- -emit-llvm-bc -o /dev/null \
// RUN: %clang_cc1 -triple amdgpu-- -emit-llvm-bc -o /dev/null \
// RUN: -mllvm -print-pipeline-passes \
// RUN: %s 2>&1 | FileCheck --check-prefixes=DEFAULT %s
// DEFAULT-NOT: LinkInModulesPass
// RUN: %clang_cc1 -triple amdgcn-- -emit-llvm-bc -o /dev/null \
// RUN: %clang_cc1 -triple amdgpu-- -emit-llvm-bc -o /dev/null \
// RUN: -mllvm -print-pipeline-passes \
// RUN: -mlink-builtin-bitcode-postopt \
// RUN: %s 2>&1 | FileCheck --check-prefixes=OPTION-POSITIVE %s
// OPTION-POSITIVE: LinkInModulesPass,assign-guid
// RUN: %clang_cc1 -triple amdgcn-- -emit-llvm-bc -o /dev/null \
// RUN: %clang_cc1 -triple amdgpu-- -emit-llvm-bc -o /dev/null \
// RUN: -mllvm -print-pipeline-passes \
// RUN: -mno-link-builtin-bitcode-postopt \
// RUN: %s 2>&1 | FileCheck --check-prefixes=OPTION-NEGATIVE %s
// OPTION-NEGATIVE-NOT: LinkInModulesPass
// RUN: %clang_cc1 -triple amdgcn-- -emit-llvm-bc -o /dev/null \
// RUN: %clang_cc1 -triple amdgpu-- -emit-llvm-bc -o /dev/null \
// RUN: -mllvm -print-pipeline-passes \
// RUN: -mlink-builtin-bitcode-postopt \
// RUN: -mno-link-builtin-bitcode-postopt \

10
clang/test/CodeGen/logb_scalbn.c
View file

@ -1,9 +1,9 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
// RUN: %clang -cc1 -triple amdgcn-amd-amdhsa -o - -emit-llvm %s | FileCheck %s -check-prefixes=DEFAULT
// RUN: %clang -cc1 -triple amdgcn-amd-amdhsa -o - -ffp-exception-behavior=ignore -emit-llvm %s | FileCheck %s -check-prefixes=IGNORE
// RUN: %clang -cc1 -triple amdgcn-amd-amdhsa -o - -ffp-exception-behavior=strict -emit-llvm %s | FileCheck %s -check-prefixes=STRICT
// RUN: %clang -cc1 -triple amdgcn-amd-amdhsa -o - -ffp-exception-behavior=maytrap -emit-llvm %s | FileCheck %s -check-prefixes=MAYTRAP
// RUN: %clang -cc1 -triple amdgcn-amd-amdhsa -o - -fmath-errno -emit-llvm %s | FileCheck %s -check-prefixes=ERRNO
// RUN: %clang -cc1 -triple amdgpu-amd-amdhsa -o - -emit-llvm %s | FileCheck %s -check-prefixes=DEFAULT
// RUN: %clang -cc1 -triple amdgpu-amd-amdhsa -o - -ffp-exception-behavior=ignore -emit-llvm %s | FileCheck %s -check-prefixes=IGNORE
// RUN: %clang -cc1 -triple amdgpu-amd-amdhsa -o - -ffp-exception-behavior=strict -emit-llvm %s | FileCheck %s -check-prefixes=STRICT
// RUN: %clang -cc1 -triple amdgpu-amd-amdhsa -o - -ffp-exception-behavior=maytrap -emit-llvm %s | FileCheck %s -check-prefixes=MAYTRAP
// RUN: %clang -cc1 -triple amdgpu-amd-amdhsa -o - -fmath-errno -emit-llvm %s | FileCheck %s -check-prefixes=ERRNO
// RUN: %clang -cc1 -triple spirv64-amd-amdhsa -o - -emit-llvm %s | FileCheck %s -check-prefixes=AMDGCNSPIRV-DEFAULT
// RUN: %clang -cc1 -triple spirv64-amd-amdhsa -o - -ffp-exception-behavior=ignore -emit-llvm %s | FileCheck %s -check-prefixes=AMDGCNSPIRV-IGNORE
// RUN: %clang -cc1 -triple spirv64-amd-amdhsa -o - -ffp-exception-behavior=strict -emit-llvm %s | FileCheck %s -check-prefixes=AMDGCNSPIRV-STRICT

2
clang/test/CodeGen/pragma-detect_mismatch.c
View file

@ -2,7 +2,7 @@
// RUN: | FileCheck %s
// RUN: %clang_cc1 %s -triple thumbv7-windows -fms-extensions -emit-llvm -o - \
// RUN: | FileCheck %s
// RUN: %clang_cc1 %s -triple amdgcn-amd-amdhsa -fms-extensions -emit-llvm -o \
// RUN: %clang_cc1 %s -triple amdgpu-amd-amdhsa -fms-extensions -emit-llvm -o \
// RUN: - | FileCheck -check-prefix=AMD %s
#pragma detect_mismatch("test", "1")

4
clang/test/CodeGen/scoped-atomic-ops.c
View file

@ -1,7 +1,7 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
// RUN: %clang_cc1 %s -emit-llvm -o - -triple=amdgcn-amd-amdhsa -ffreestanding \
// RUN: %clang_cc1 %s -emit-llvm -o - -triple=amdgpu-amd-amdhsa -ffreestanding \
// RUN: -fvisibility=hidden | FileCheck --check-prefixes=AMDGCN,AMDGCN_CL_DEF %s
// RUN: %clang_cc1 %s -emit-llvm -o - -triple=amdgcn-amd-amdhsa -ffreestanding \
// RUN: %clang_cc1 %s -emit-llvm -o - -triple=amdgpu-amd-amdhsa -ffreestanding \
// RUN: -cl-std=CL2.0 -fvisibility=hidden | FileCheck --check-prefixes=AMDGCN,AMDGCN_CL_20 %s
// RUN: %clang_cc1 %s -emit-llvm -o - -triple=nvptx64-nvidia-cuda -ffreestanding \
// RUN: -cl-std=CL2.0 -fvisibility=hidden | FileCheck --check-prefixes=NVPTX %s

4
clang/test/CodeGen/scoped-fence-ops.c
View file

@ -1,7 +1,7 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
// RUN: %clang_cc1 %s -emit-llvm -o - -triple=amdgcn-amd-amdhsa -ffreestanding \
// RUN: %clang_cc1 %s -emit-llvm -o - -triple=amdgpu-amd-amdhsa -ffreestanding \
// RUN: -fvisibility=hidden | FileCheck --check-prefixes=AMDGCN,AMDGCN_CL_DEF %s
// RUN: %clang_cc1 %s -emit-llvm -o - -triple=amdgcn-amd-amdhsa -ffreestanding \
// RUN: %clang_cc1 %s -emit-llvm -o - -triple=amdgpu-amd-amdhsa -ffreestanding \
// RUN: -cl-std=CL2.0 -fvisibility=hidden | FileCheck --check-prefixes=AMDGCN,AMDGCN_CL_20 %s
// RUN: %clang_cc1 %s -emit-llvm -o - -triple=spirv64-unknown-unknown -ffreestanding \
// RUN: -fvisibility=hidden | FileCheck --check-prefix=SPIRV %s

2
clang/test/CodeGen/sret.c
View file

@ -1,5 +1,5 @@
// RUN: %clang_cc1 %s -Wno-strict-prototypes -emit-llvm -o - | FileCheck %s
// RUN: %clang_cc1 %s -Wno-strict-prototypes -triple amdgcn-amd-amdhsa -emit-llvm -o - | FileCheck --check-prefix=NONZEROALLOCAAS %s
// RUN: %clang_cc1 %s -Wno-strict-prototypes -triple amdgpu-amd-amdhsa -emit-llvm -o - | FileCheck --check-prefix=NONZEROALLOCAAS %s
struct abc {
long a;

2
clang/test/CodeGen/target-addrspace.cpp
View file

@ -1,7 +1,7 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -emit-llvm \
// RUN: -fvisibility=hidden -o - %s | FileCheck %s --check-prefix=NVPTX
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -emit-llvm \
// RUN: -fvisibility=hidden -o - %s | FileCheck %s --check-prefix=AMDGPU
// NVPTX-LABEL: define hidden void @_Z1fPv(

2
clang/test/CodeGen/target-data.c
View file

@ -163,7 +163,7 @@
// R600SI: target datalayout = "e-m:e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-p7:160:256:256:32-p8:128:128:128:48-p9:192:256:256:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5-G1-ni:7:8:9"
// Test default -target-cpu
// RUN: %clang_cc1 -triple amdgcn-unknown -o - -emit-llvm %s \
// RUN: %clang_cc1 -triple amdgpu-unknown -o - -emit-llvm %s \
// RUN: | FileCheck %s -check-prefix=R600SIDefault
// R600SIDefault: target datalayout = "e-m:e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-p7:160:256:256:32-p8:128:128:128:48-p9:192:256:256:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5-G1-ni:7:8:9"

2
clang/test/CodeGenCUDA/address-spaces.cu
View file

@ -1,5 +1,5 @@
// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx-unknown-unknown | FileCheck %s
// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple amdgcn | FileCheck %s
// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple amdgpu | FileCheck %s
// Verifies Clang emits correct address spaces and addrspacecast instructions
// for CUDA code.

2
clang/test/CodeGenCUDA/alias.cu
View file

@ -6,7 +6,7 @@
// RUN: -o - %s | FileCheck %s
// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -emit-llvm -target-sdk-version=10.1 \
// RUN: -o - %s | FileCheck %s
// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
// RUN: %clang_cc1 -fcuda-is-device -triple amdgpu-amd-amdhsa -emit-llvm \
// RUN: -o - %s | FileCheck %s
#include "Inputs/cuda.h"

12
clang/test/CodeGenCUDA/amdgpu-code-object-version.cu
View file

@ -1,21 +1,21 @@
// Create module flag for code object version.
// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
// RUN: %clang_cc1 -fcuda-is-device -triple amdgpu-amd-amdhsa -emit-llvm \
// RUN: -o - %s | FileCheck %s -check-prefix=V6
// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
// RUN: %clang_cc1 -fcuda-is-device -triple amdgpu-amd-amdhsa -emit-llvm \
// RUN: -mcode-object-version=4 -o - %s | FileCheck -check-prefix=V4 %s
// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
// RUN: %clang_cc1 -fcuda-is-device -triple amdgpu-amd-amdhsa -emit-llvm \
// RUN: -mcode-object-version=5 -o - %s | FileCheck -check-prefix=V5 %s
// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
// RUN: %clang_cc1 -fcuda-is-device -triple amdgpu-amd-amdhsa -emit-llvm \
// RUN: -mcode-object-version=6 -o - %s | FileCheck -check-prefix=V6 %s
// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
// RUN: %clang_cc1 -fcuda-is-device -triple amdgpu-amd-amdhsa -emit-llvm \
// RUN: -mcode-object-version=none -o - %s | FileCheck %s -check-prefix=NONE
// RUN: not %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
// RUN: not %clang_cc1 -fcuda-is-device -triple amdgpu-amd-amdhsa -emit-llvm \
// RUN: -mcode-object-version=4.1 -o - %s 2>&1| FileCheck %s -check-prefix=INV
// V4: !{{.*}} = !{i32 1, !"amdhsa_code_object_version", i32 400}

4
clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu
View file

@ -2,9 +2,9 @@
// REQUIRES: x86-registered-target
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm -x hip %s -o - | FileCheck --check-prefixes=CHECK %s
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -emit-llvm -x hip %s -o - | FileCheck --check-prefixes=CHECK %s
// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -fcuda-is-device -emit-llvm -x hip %s -o - | FileCheck --check-prefixes=CHECK-SPIRV %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm -x hip %s -disable-O0-optnone -o - | opt -S -O2 | FileCheck %s --check-prefixes=OPT
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -emit-llvm -x hip %s -disable-O0-optnone -o - | opt -S -O2 | FileCheck %s --check-prefixes=OPT
// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -fcuda-is-device -emit-llvm -x hip %s -disable-O0-optnone -o - | opt -S -O2 | FileCheck %s --check-prefixes=OPT-SPIRV
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -x hip %s -o - | FileCheck -check-prefix=HOST %s

8
clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu
View file

@ -1,7 +1,7 @@
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa \
// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \
// RUN: | FileCheck -check-prefixes=CHECK,DEFAULT %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa --gpu-max-threads-per-block=1024 \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa --gpu-max-threads-per-block=1024 \
// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \
// RUN: | FileCheck -check-prefixes=CHECK,MAX1024 %s
// RUN: %clang_cc1 -triple spirv64-amd-amdhsa --gpu-max-threads-per-block=1024 \
@ -13,10 +13,10 @@
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm \
// RUN: -verify -Wno-deprecated-declarations -o - -x hip %s | FileCheck -check-prefix=NAMD %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -foffload-uniform-block \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -foffload-uniform-block \
// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \
// RUN: | FileCheck -check-prefixes=CHECK,DEFAULT %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fno-offload-uniform-block \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fno-offload-uniform-block \
// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \
// RUN: | FileCheck -check-prefixes=NOUB %s

6
clang/test/CodeGenCUDA/amdgpu-visibility.cu
View file

@ -1,6 +1,6 @@
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -fapply-global-visibility-to-externs -fvisibility=default -emit-llvm -o - %s | FileCheck --check-prefix=CHECK-DEFAULT %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -fapply-global-visibility-to-externs -fvisibility=protected -emit-llvm -o - %s | FileCheck --check-prefix=CHECK-PROTECTED %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -fapply-global-visibility-to-externs -fvisibility=hidden -emit-llvm -o - %s | FileCheck --check-prefix=CHECK-HIDDEN %s
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -fapply-global-visibility-to-externs -fvisibility=default -emit-llvm -o - %s | FileCheck --check-prefix=CHECK-DEFAULT %s
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -fapply-global-visibility-to-externs -fvisibility=protected -emit-llvm -o - %s | FileCheck --check-prefix=CHECK-PROTECTED %s
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -fapply-global-visibility-to-externs -fvisibility=hidden -emit-llvm -o - %s | FileCheck --check-prefix=CHECK-HIDDEN %s
#include "Inputs/cuda.h"

6
clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu
View file

@ -1,13 +1,13 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa \
// RUN: -fcuda-is-device -mcode-object-version=4 -emit-llvm -o - -x hip %s \
// RUN: | FileCheck -check-prefix=PRECOV5 %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa \
// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \
// RUN: | FileCheck -check-prefix=COV5 %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa \
// RUN: -fcuda-is-device -mcode-object-version=6 -emit-llvm -o - -x hip %s \
// RUN: | FileCheck -check-prefix=COV5 %s

2
clang/test/CodeGenCUDA/anon-ns.cu
View file

@ -1,4 +1,4 @@
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -cuid=abc \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -cuid=abc \
// RUN: -aux-triple x86_64-unknown-linux-gnu -std=c++17 -fgpu-rdc \
// RUN: -emit-llvm -o - -x hip %s > %t.dev

2
clang/test/CodeGenCUDA/atomic-ops.cu
View file

@ -1,4 +1,4 @@
// RUN: %clang_cc1 -x hip -std=c++11 -triple amdgcn -fcuda-is-device -emit-llvm %s -o - | FileCheck -enable-var-scope %s
// RUN: %clang_cc1 -x hip -std=c++11 -triple amdgpu -fcuda-is-device -emit-llvm %s -o - | FileCheck -enable-var-scope %s
#include "Inputs/cuda.h"
// CHECK-LABEL: @_Z24atomic32_op_singlethreadPiii

4
clang/test/CodeGenCUDA/atomic-options.hip
View file

@ -1,9 +1,9 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
// RUN: -emit-llvm -o - %s | FileCheck --check-prefix=HOST %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa \
// RUN: -emit-llvm -o - -fcuda-is-device %s | FileCheck --check-prefix=DEV %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa \
// RUN: -fatomic-fine-grained-memory -fatomic-ignore-denormal-mode \
// RUN: -emit-llvm -o - -fcuda-is-device %s | FileCheck --check-prefix=OPT %s
// RUN: %clang_cc1 -triple spirv64-amd-amdhsa \

4
clang/test/CodeGenCUDA/const-var.cu
View file

@ -1,12 +1,12 @@
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -x hip %s \
// RUN: -emit-llvm -o - | FileCheck -check-prefix=DEV %s
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip %s \
// RUN: -emit-llvm -o - | FileCheck -check-prefix=HOST %s
// Negative tests.
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -x hip %s \
// RUN: -emit-llvm -o - | FileCheck -check-prefix=DEV-NEG %s
#include "Inputs/cuda.h"

2
clang/test/CodeGenCUDA/dependent-libs.cu
View file

@ -1,6 +1,6 @@
// REQUIRES: x86-registered-target
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm -o - -fcuda-is-device -x hip %s | FileCheck --check-prefix=DEV %s
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -emit-llvm -o - -fcuda-is-device -x hip %s | FileCheck --check-prefix=DEV %s
// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm -o - -x hip %s | FileCheck --check-prefix=HOST %s
// DEV-NOT: llvm.dependent-libraries

8
clang/test/CodeGenCUDA/device-const-var-linkage.cu
View file

@ -2,21 +2,21 @@
// by host code. Variables only used in device code retain internal linkage.
// Non-RDC mode
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -x hip %s \
// RUN: -std=c++17 -emit-llvm -o - | FileCheck -check-prefix=DEV %s
// RDC mode
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -x hip %s \
// RUN: -std=c++17 -fgpu-rdc -emit-llvm -o - | FileCheck -check-prefix=RDC %s
// With -fvisibility=hidden
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -x hip %s \
// RUN: -std=c++17 -fvisibility=hidden -fapply-global-visibility-to-externs \
// RUN: -emit-llvm -o - | FileCheck -check-prefix=HIDDEN %s
// Negative test: const device vars NOT referenced by host should not be
// externalized.
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -x hip %s \
// RUN: -std=c++17 -emit-llvm -o - | FileCheck -check-prefix=NEG %s
#include "Inputs/cuda.h"

2
clang/test/CodeGenCUDA/device-init-fun.cu
View file

@ -1,6 +1,6 @@
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -std=c++11 \
// RUN: %clang_cc1 -triple amdgpu -fcuda-is-device -std=c++11 \
// RUN: -fgpu-allow-device-init -x hip \
// RUN: -fno-threadsafe-statics -emit-llvm -o - %s \
// RUN: | FileCheck %s

4
clang/test/CodeGenCUDA/device-use-host-var.cu
View file

@ -1,6 +1,6 @@
// RUN: %clang_cc1 -std=c++14 -triple amdgcn-amd-amdhsa \
// RUN: %clang_cc1 -std=c++14 -triple amdgpu-amd-amdhsa \
// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s | FileCheck %s
// RUN: %clang_cc1 -std=c++14 -triple amdgcn-amd-amdhsa \
// RUN: %clang_cc1 -std=c++14 -triple amdgpu-amd-amdhsa \
// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s | FileCheck -check-prefix=NEG %s
#include "Inputs/cuda.h"

2
clang/test/CodeGenCUDA/device-var-init.cu
View file

@ -9,7 +9,7 @@
// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -std=c++11 \
// RUN: -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck -check-prefixes=HOST %s
// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -std=c++11 \
// RUN: %clang_cc1 -triple amdgpu -fcuda-is-device -std=c++11 \
// RUN: -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck -check-prefixes=DEVICE,AMDGCN %s
#ifdef __clang__

4
clang/test/CodeGenCUDA/device-var-linkage.cu
View file

@ -1,7 +1,7 @@
// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device \
// RUN: %clang_cc1 -triple amdgpu -fcuda-is-device \
// RUN: -emit-llvm -o - -x hip %s \
// RUN: | FileCheck -check-prefixes=DEV,NORDC %s
// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device \
// RUN: %clang_cc1 -triple amdgpu -fcuda-is-device \
// RUN: -fgpu-rdc -cuid=abc -emit-llvm -o - -x hip %s \
// RUN: | FileCheck -check-prefixes=DEV,RDC %s
// RUN: %clang_cc1 -triple x86_64-unknown-gnu-linux \

4
clang/test/CodeGenCUDA/dft-func-attr-skip-intrinsic.hip
View file

@ -1,6 +1,6 @@
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -x ir -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm-bc -disable-llvm-passes -o %t.bc %S/Inputs/device-lib-code.ll
// RUN: %clang_cc1 -x hip -fcuda-is-device -triple amdgcn-amd-amdhsa -mlink-builtin-bitcode %t.bc -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s
// RUN: %clang_cc1 -x ir -fcuda-is-device -triple amdgpu-amd-amdhsa -emit-llvm-bc -disable-llvm-passes -o %t.bc %S/Inputs/device-lib-code.ll
// RUN: %clang_cc1 -x hip -fcuda-is-device -triple amdgpu-amd-amdhsa -mlink-builtin-bitcode %t.bc -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s
#include "Inputs/cuda.h"

4
clang/test/CodeGenCUDA/host-used-device-var.cu
View file

@ -1,5 +1,5 @@
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -x hip %s \
// RUN: -std=c++17 -O3 -mllvm -amdgpu-internalize-symbols -emit-llvm -o - \
// RUN: -cuid=123 | FileCheck -check-prefix=DEV %s
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip %s \
@ -7,7 +7,7 @@
// Negative tests.
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -x hip %s \
// RUN: -std=c++17 -O3 -mllvm -amdgpu-internalize-symbols -emit-llvm -o - \
// RUN: | FileCheck -check-prefix=DEV-NEG %s
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip %s \

4
clang/test/CodeGenCUDA/implicit-hd-dtor-trap-stub-hip.hip
View file

@ -5,9 +5,9 @@
// slot uses the real mangled name, dtor body is real. With or without an
// organic device caller, the emitted IR is the same.
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -x hip \
// RUN: -emit-llvm -o - %s | FileCheck %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -x hip \
// RUN: -DWITH_DEVICE_CALLER -emit-llvm -o - %s | FileCheck %s
// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -x hip \
// RUN: -emit-llvm -o - %s | FileCheck %s

2
clang/test/CodeGenCUDA/implicit-hd-explicit-inst-hip.hip
View file

@ -5,7 +5,7 @@
// get a real body if their device codegen is clean, or a trap body if it
// would otherwise raise a target-mismatch diagnostic.
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -x hip \
// RUN: -emit-llvm -o - %s | FileCheck %s
// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -x hip \
// RUN: -emit-llvm -o - %s | FileCheck %s

2
clang/test/CodeGenCUDA/implicit-hd-overload-ambig-hip.hip
View file

@ -4,7 +4,7 @@
// Overload ambiguity inside an implicit-H+D explicit-inst member without an
// organic device caller: error is silenced, function body is a trap stub.
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -x hip \
// RUN: -emit-llvm -o - %s | FileCheck %s
// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -x hip \
// RUN: -emit-llvm -o - %s | FileCheck %s

2
clang/test/CodeGenCUDA/kernel-amdgcn.cu
View file

@ -1,4 +1,4 @@
// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -emit-llvm -x hip %s -o - | FileCheck %s
// RUN: %clang_cc1 -triple amdgpu -fcuda-is-device -emit-llvm -x hip %s -o - | FileCheck %s
// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -fcuda-is-device -emit-llvm -x hip %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
#include "Inputs/cuda.h"

4
clang/test/CodeGenCUDA/kernel-arg-name-metadata.cu
View file

@ -1,8 +1,8 @@
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fhip-kernel-arg-name \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fhip-kernel-arg-name \
// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \
// RUN: | FileCheck %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa \
// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \
// RUN: | FileCheck -check-prefix=NEG %s

2
clang/test/CodeGenCUDA/kernel-args.cu
View file

@ -1,4 +1,4 @@
// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -fcuda-is-device \
// RUN: %clang_cc1 -x hip -triple amdgpu-amd-amdhsa -fcuda-is-device \
// RUN: -emit-llvm %s -o - | FileCheck -check-prefix=AMDGCN %s
// RUN: %clang_cc1 -x hip -triple spirv64-amd-amdhsa -fcuda-is-device \
// RUN: -emit-llvm %s -o - | FileCheck -check-prefix=AMDGCNSPIRV %s

6
clang/test/CodeGenCUDA/kernel-dbg-info.cu
View file

@ -3,7 +3,7 @@
// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -O0 \
// RUN: -fcuda-include-gpubinary %t -debug-info-kind=limited \
// RUN: -o - -x hip | FileCheck -check-prefixes=CHECK,O0 %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm %s -O0 \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -emit-llvm %s -O0 \
// RUN: -fcuda-include-gpubinary %t -debug-info-kind=limited \
// RUN: -o - -x hip -fcuda-is-device | FileCheck -check-prefix=DEV %s
@ -11,7 +11,7 @@
// RUN: -fcuda-include-gpubinary %t -debug-info-kind=limited \
// RUN: -o - -x hip -debugger-tuning=gdb -dwarf-version=4 \
// RUN: | FileCheck -check-prefixes=CHECK,O0 %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm %s -O0 \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -emit-llvm %s -O0 \
// RUN: -fcuda-include-gpubinary %t -debug-info-kind=limited \
// RUN: -o - -x hip -debugger-tuning=gdb -dwarf-version=4 \
// RUN: -fcuda-is-device | FileCheck -check-prefix=DEV %s
@ -19,7 +19,7 @@
// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -O3 \
// RUN: -fcuda-include-gpubinary %t -debug-info-kind=limited \
// RUN: -o - -x hip -debugger-tuning=gdb -dwarf-version=4 | FileCheck %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm %s -O3 \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -emit-llvm %s -O3 \
// RUN: -fcuda-include-gpubinary %t -debug-info-kind=limited \
// RUN: -o - -x hip -debugger-tuning=gdb -dwarf-version=4 \
// RUN: -fcuda-is-device | FileCheck -check-prefix=DEV %s

2
clang/test/CodeGenCUDA/lambda-constexpr-capture.cu
View file

@ -1,6 +1,6 @@
// RUN: %clang_cc1 -emit-llvm -x hip %s -o - -triple x86_64-linux-gnu \
// RUN: | FileCheck -check-prefixes=CHECK,HOST %s
// RUN: %clang_cc1 -emit-llvm -x hip %s -o - -triple amdgcn-amd-amdhsa -fcuda-is-device \
// RUN: %clang_cc1 -emit-llvm -x hip %s -o - -triple amdgpu-amd-amdhsa -fcuda-is-device \
// RUN: | FileCheck -check-prefixes=CHECK,DEV %s
#include "Inputs/cuda.h"

2
clang/test/CodeGenCUDA/lambda-noinline.cu
View file

@ -2,7 +2,7 @@
// RUN: -triple x86_64-linux-gnu \
// RUN: | FileCheck -check-prefix=HOST %s
// RUN: %clang_cc1 -x hip -emit-llvm -std=c++11 %s -o - \
// RUN: -triple amdgcn-amd-amdhsa -fcuda-is-device \
// RUN: -triple amdgpu-amd-amdhsa -fcuda-is-device \
// RUN: | FileCheck -check-prefix=DEV %s
#include "Inputs/cuda.h"

2
clang/test/CodeGenCUDA/lambda-reference-var.cu
View file

@ -2,7 +2,7 @@
// RUN: -triple x86_64-linux-gnu \
// RUN: | FileCheck -check-prefix=HOST %s
// RUN: %clang_cc1 -x hip -emit-llvm -std=c++11 %s -o - \
// RUN: -triple amdgcn-amd-amdhsa -fcuda-is-device \
// RUN: -triple amdgpu-amd-amdhsa -fcuda-is-device \
// RUN: | FileCheck -check-prefix=DEV %s
#include "Inputs/cuda.h"

2
clang/test/CodeGenCUDA/lambda.cu
View file

@ -2,7 +2,7 @@
// RUN: -triple x86_64-linux-gnu \
// RUN: | FileCheck -check-prefix=HOST %s
// RUN: %clang_cc1 -x hip -emit-llvm -std=c++11 %s -o - \
// RUN: -triple amdgcn-amd-amdhsa -fcuda-is-device \
// RUN: -triple amdgpu-amd-amdhsa -fcuda-is-device \
// RUN: | FileCheck -check-prefix=DEV %s
#include "Inputs/cuda.h"

33
clang/test/CodeGenCUDA/link-builtin-bitcode-denormal-fp-mode.cu
View file

@ -5,31 +5,31 @@
// Build the fake device library in the way rocm-device-libs should be built.
//
// RUN: %clang_cc1 -x cl -triple amdgcn-amd-amdhsa -fdenormal-fp-math-f32=dynamic \
// RUN: %clang_cc1 -x cl -triple amdgpu-amd-amdhsa -fdenormal-fp-math-f32=dynamic \
// RUN: -mcode-object-version=none -emit-llvm-bc \
// RUN: %S/Inputs/ocml-sample.cl -o %t.dynamic.f32.bc
//
// RUN: %clang_cc1 -x cl -triple amdgcn-amd-amdhsa -fdenormal-fp-math=dynamic \
// RUN: %clang_cc1 -x cl -triple amdgpu-amd-amdhsa -fdenormal-fp-math=dynamic \
// RUN: -mcode-object-version=none -emit-llvm-bc \
// RUN: %S/Inputs/ocml-sample.cl -o %t.dynamic.full.bc
// Check the default behavior with no denormal-fp-math arguments.
// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 -fcuda-is-device \
// RUN: %clang_cc1 -x hip -triple amdgpu8.03-amd-amdhsa -fcuda-is-device \
// RUN: -mlink-builtin-bitcode %t.dynamic.f32.bc \
// RUN: -emit-llvm %s -o - | FileCheck -implicit-check-not=denormal-fp-math %s --check-prefixes=CHECK,INTERNALIZE
// Check an explicit full ieee request
// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 -fcuda-is-device \
// RUN: %clang_cc1 -x hip -triple amdgpu8.03-amd-amdhsa -fcuda-is-device \
// RUN: -fdenormal-fp-math=ieee \
// RUN: -mlink-builtin-bitcode %t.dynamic.f32.bc \
// RUN: -emit-llvm %s -o - | FileCheck -implicit-check-not=denormal-fp-math %s --check-prefixes=CHECK,INTERNALIZE
// Check explicit f32-only flushing request
// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 \
// RUN: %clang_cc1 -x hip -triple amdgpu8.03-amd-amdhsa \
// RUN: -fcuda-is-device -fdenormal-fp-math-f32=preserve-sign \
// RUN: -mlink-builtin-bitcode %t.dynamic.f32.bc -emit-llvm %s -o - \
// RUN: | FileCheck -implicit-check-not=denormal-fp-math --enable-var-scope %s --check-prefixes=CHECK,INTERNALIZE,IEEEF64-PSZF32
@ -38,21 +38,21 @@
// Check explicit flush all request. Only the f32 component of the library is
// dynamic, so the linked functions should use IEEE as the base mode and the new
// functions preserve-sign.
// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 \
// RUN: %clang_cc1 -x hip -triple amdgpu8.03-amd-amdhsa \
// RUN: -fcuda-is-device -fdenormal-fp-math=preserve-sign \
// RUN: -mlink-builtin-bitcode %t.dynamic.f32.bc -emit-llvm %s -o - \
// RUN: | FileCheck -implicit-check-not=denormal-fp-math --enable-var-scope %s --check-prefixes=CHECK,INTERNALIZE,PSZ
// Check explicit f32-only, ieee-other flushing request
// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 \
// RUN: %clang_cc1 -x hip -triple amdgpu8.03-amd-amdhsa \
// RUN: -fcuda-is-device -fdenormal-fp-math=ieee -fdenormal-fp-math-f32=preserve-sign \
// RUN: -mlink-builtin-bitcode %t.dynamic.f32.bc -emit-llvm %s -o - \
// RUN: | FileCheck -implicit-check-not=denormal-fp-math --enable-var-scope %s --check-prefixes=CHECK,INTERNALIZE,IEEEF64-PSZF32
// Check inverse of normal usage. Requesting IEEE f32, with flushed f16/f64
// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 \
// RUN: %clang_cc1 -x hip -triple amdgpu8.03-amd-amdhsa \
// RUN: -fcuda-is-device -fdenormal-fp-math=preserve-sign -fdenormal-fp-math-f32=ieee \
// RUN: -mlink-builtin-bitcode %t.dynamic.f32.bc -emit-llvm %s -o - \
// RUN: | FileCheck -implicit-check-not=denormal-fp-math --enable-var-scope %s --check-prefixes=CHECK,INTERNALIZE,IEEEF32-PSZF64-DYNF32
@ -60,7 +60,7 @@
// Check backwards from the normal usage where both library components can be
// overridden.
// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 \
// RUN: %clang_cc1 -x hip -triple amdgpu8.03-amd-amdhsa \
// RUN: -fcuda-is-device -fdenormal-fp-math=preserve-sign -fdenormal-fp-math-f32=ieee \
// RUN: -mlink-builtin-bitcode %t.dynamic.full.bc -emit-llvm %s -o - \
// RUN: | FileCheck -implicit-check-not=denormal-fp-math --enable-var-scope %s --check-prefixes=CHECK,INTERNALIZE,IEEEF32-PSZF64-DYNFULL
@ -68,7 +68,7 @@
// Check the case where no internalization is performed
// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 \
// RUN: %clang_cc1 -x hip -triple amdgpu8.03-amd-amdhsa \
// RUN: -fcuda-is-device -fdenormal-fp-math=preserve-sign -fdenormal-fp-math-f32=ieee \
// RUN: -mlink-bitcode-file %t.dynamic.full.bc -emit-llvm %s -o - \
// RUN: | FileCheck -implicit-check-not=denormal-fp-math --enable-var-scope %s --check-prefixes=CHECK,NOINTERNALIZE,NOINTERNALIZE-IEEEF32-PSZF64-DYNFULL
@ -133,33 +133,24 @@ __global__ void kernel_f64(double* out, double* a, double* b, double* c) {
// Default mode relies on the implicit check-not for the denormal-fp-math.
// PSZ: #[[$KERNELATTR]] = { {{.*}} denormal_fpenv(preservesign)
// PSZ-SAME: "target-cpu"="gfx803"
// PSZ: #[[$FUNCATTR]] = { {{.*}} denormal_fpenv(float: preservesign)
// PSZ-SAME: "target-cpu"="gfx803"
// PSZ: #[[$WEAK_FUNCATTR]] = { {{.*}} denormal_fpenv(float: preservesign)
// PSZ-SAME: "target-cpu"="gfx803"
// FIXME: Should check-not denormal_fpenv within the line
// IEEEF64-PSZF32: #[[$KERNELATTR]] = { {{.*}} denormal_fpenv(float: preservesign)
// IEEEF64-PSZF32-SAME: "target-cpu"="gfx803"
// IEEEF64-PSZF32: #[[$FUNCATTR]] = { {{.*}} denormal_fpenv(float: preservesign)
// IEEEF64-PSZF32-SAME: "target-cpu"="gfx803"
// IEEEF64-PSZF32: #[[$WEAK_FUNCATTR]] = { {{.*}} denormal_fpenv(float: preservesign)
// IEEEF64-PSZF32-SAME: "target-cpu"="gfx803"
// IEEEF32-PSZF64-DYNF32: #[[$KERNELATTR]] = { {{.*}} denormal_fpenv(preservesign, float: ieee) {{.*}} "target-cpu"="gfx803" {{.*}} }
// IEEEF32-PSZF64-DYNF32: #[[$KERNELATTR]] = { {{.*}} denormal_fpenv(preservesign, float: ieee) {{.*}} }
// implicit-not
// implicit check-not
// IEEEF32-PSZF64-DYNFULL: #[[$KERNELATTR]] = { {{.*}} denormal_fpenv(preservesign, float: ieee)
// IEEEF32-PSZF64-DYNFULL-SAME: "target-cpu"="gfx803"
// IEEEF32-PSZF64-DYNFULL: #[[$FUNCATTR]] = { {{.*}} denormal_fpenv(preservesign, float: ieee)
// IEEEF32-PSZF64-DYNFULL-SAME: "target-cpu"="gfx803"
// IEEEF32-PSZF64-DYNFULL: #[[$WEAK_FUNCATTR]] = { {{.*}} denormal_fpenv(preservesign, float: ieee)
// IEEEF32-PSZF64-DYNFULL-SAME: "target-cpu"="gfx803"
// -mlink-bitcode-file doesn't internalize or propagate attributes.
// NOINTERNALIZE-IEEEF32-PSZF64-DYNFULL: #[[$KERNELATTR]] = { {{.*}} denormal_fpenv(preservesign, float: ieee) {{.*}} "target-cpu"="gfx803" {{.*}} }
// NOINTERNALIZE-IEEEF32-PSZF64-DYNFULL: #[[$KERNELATTR]] = { {{.*}} denormal_fpenv(preservesign, float: ieee) {{.*}} }
// NOINTERNALIZE-IEEEF32-PSZF64-DYNFULL: #[[$FUNCATTR]] = { {{.*}} denormal_fpenv(dynamic) {{.*}} }
// NOINTERNALIZE-IEEEF32-PSZF64-DYNFULL: #[[$WEAK_FUNCATTR]] = { {{.*}} denormal_fpenv(dynamic) {{.*}} }

12
clang/test/CodeGenCUDA/link-builtin-bitcode-gpu-attrs-preserved.cu
View file

@ -9,29 +9,29 @@
// Build the fake device library in the way rocm-device-libs should be built.
//
// RUN: %clang_cc1 -x cl -triple amdgcn-amd-amdhsa\
// RUN: %clang_cc1 -x cl -triple amdgpu-amd-amdhsa\
// RUN: -mcode-object-version=none -emit-llvm-bc \
// RUN: %S/Inputs/ocml-sample-target-attrs.cl -o %t.bc
// Check the default behavior
// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 -fcuda-is-device \
// RUN: %clang_cc1 -x hip -triple amdgpu8.03-amd-amdhsa -fcuda-is-device \
// RUN: -mlink-builtin-bitcode %t.bc \
// RUN: -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,INTERNALIZE
// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx1101 -fcuda-is-device \
// RUN: %clang_cc1 -x hip -triple amdgpu11.01-amd-amdhsa -fcuda-is-device \
// RUN: -mlink-builtin-bitcode %t.bc -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,INTERNALIZE
// Check the case where no internalization is performed
// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 \
// RUN: %clang_cc1 -x hip -triple amdgpu8.03-amd-amdhsa \
// RUN: -fcuda-is-device -mlink-bitcode-file %t.bc -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,NOINTERNALIZE
// Check the case where no internalization is performed
// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx1101 \
// RUN: %clang_cc1 -x hip -triple amdgpu11.01-amd-amdhsa \
// RUN: -fcuda-is-device -mlink-bitcode-file %t.bc -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,NOINTERNALIZE
// CHECK: define {{.*}} i64 @do_intrin_stuff() #[[ATTR:[0-9]+]]
// INTERNALIZE: attributes #[[ATTR]] = {{.*}} "target-cpu"="gfx{{.*}}" "target-features"="{{.*}}+gfx11-insts{{.*}}"
// INTERNALIZE: attributes #[[ATTR]] = {{.*}} "target-features"="{{.*}}+gfx11-insts{{.*}}"
// NOINTERNALIZE: attributes #[[ATTR]] = {{.*}} "target-features"="+gfx11-insts"
#define __device__ __attribute__((device))

4
clang/test/CodeGenCUDA/managed-var.cu
View file

@ -1,8 +1,8 @@
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -std=c++11 \
// RUN: -emit-llvm -o - -x hip %s | FileCheck \
// RUN: -check-prefixes=COMMON,DEV,NORDC-D %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -std=c++11 \
// RUN: -emit-llvm -fgpu-rdc -cuid=abc -o - -x hip %s > %t.dev
// RUN: cat %t.dev | FileCheck -check-prefixes=COMMON,DEV,RDC-D %s

4
clang/test/CodeGenCUDA/ms-linker-options.cu
View file

@ -1,10 +1,10 @@
// RUN: %clang_cc1 -emit-llvm -o - -fcuda-is-device -fms-extensions -x hip %s \
// RUN: -fno-autolink -triple amdgcn-amd-amdhsa \
// RUN: -fno-autolink -triple amdgpu-amd-amdhsa \
// RUN: | FileCheck -check-prefix=DEV %s
// RUN: %clang_cc1 -emit-llvm -o - -fms-extensions -x hip %s -triple \
// RUN: x86_64-pc-windows-msvc | FileCheck -check-prefix=HOST %s
// RUN: %clang_cc1 -emit-llvm -o - -fcuda-is-device -fms-extensions %s \
// RUN: -fno-autolink -triple amdgcn-amd-amdhsa \
// RUN: -fno-autolink -triple amdgpu-amd-amdhsa \
// RUN: | FileCheck -check-prefix=DEV %s
// RUN: %clang_cc1 -emit-llvm -o - -fms-extensions %s -triple \
// RUN: x86_64-pc-windows-msvc | FileCheck -check-prefix=HOST %s

2
clang/test/CodeGenCUDA/noinline.cu
View file

@ -3,7 +3,7 @@
// RUN: %clang_cc1 -triple nvptx-nvidia-cuda -fcuda-is-device \
// RUN: -O2 -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device \
// RUN: -O2 -disable-llvm-passes -emit-llvm -o - -x hip %s | FileCheck %s
// RUN: %clang_cc1 -triple x86_64-unknown-gnu-linux \

2
clang/test/CodeGenCUDA/norecurse.cu
View file

@ -4,7 +4,7 @@
// RUN: %clang_cc1 -triple nvptx-nvidia-cuda -fcuda-is-device \
// RUN: -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device \
// RUN: -emit-llvm -disable-llvm-passes -o - -x hip %s | FileCheck %s
#include "Inputs/cuda.h"

2
clang/test/CodeGenCUDA/nosanitize-cuid.hip
View file

@ -1,4 +1,4 @@
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fsanitize=address -fcuda-is-device \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fsanitize=address -fcuda-is-device \
// RUN: -emit-llvm -cuid=abcd -o - %s | FileCheck %s
#include "Inputs/cuda.h"

4
clang/test/CodeGenCUDA/static-device-var-no-rdc.cu
View file

@ -1,7 +1,7 @@
// REQUIRES: x86-registered-target
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -std=c++11 \
// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefix=DEV %s
// RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \
@ -9,7 +9,7 @@
// Negative tests.
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -std=c++11 \
// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefix=DEV-NEG %s
// RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \

8
clang/test/CodeGenCUDA/static-device-var-rdc.cu
View file

@ -1,7 +1,7 @@
// REQUIRES: x86-registered-target
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device \
// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o %t.nocuid.dev -x hip %s
// RUN: cat %t.nocuid.dev | FileCheck -check-prefixes=DEV,INT-DEV %s
@ -9,7 +9,7 @@
// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o %t.nocuid.host -x hip %s
// RUN: cat %t.nocuid.host | FileCheck -check-prefixes=HOST,INT-HOST %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -cuid=abc \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -cuid=abc \
// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s > %t.dev
// RUN: cat %t.dev | FileCheck -check-prefixes=DEV,EXT-DEV %s
@ -25,7 +25,7 @@
// Negative tests.
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device \
// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s | FileCheck \
// RUN: -check-prefix=DEV-NEG %s
@ -33,7 +33,7 @@
// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s | FileCheck \
// RUN: -check-prefix=HOST-NEG %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -cuid=abc \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -cuid=abc \
// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s > %t.dev
// RUN: cat %t.dev | FileCheck -check-prefix=DEV-NEG %s

4
clang/test/CodeGenCUDA/template-class-static-member.cu
View file

@ -1,4 +1,4 @@
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device \
// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefix=DEV %s
// RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \
@ -6,7 +6,7 @@
// Negative tests.
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device \
// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefix=DEV-NEG %s
#include "Inputs/cuda.h"

2
clang/test/CodeGenCUDA/types.cu
View file

@ -1,4 +1,4 @@
// RUN: %clang_cc1 -triple amdgcn -aux-triple x86_64 -fcuda-is-device -emit-llvm %s -o - | FileCheck -check-prefix=DEV %s
// RUN: %clang_cc1 -triple amdgpu -aux-triple x86_64 -fcuda-is-device -emit-llvm %s -o - | FileCheck -check-prefix=DEV %s
// RUN: %clang_cc1 -triple x86_64 -aux-triple amdgpu -emit-llvm %s -o - | FileCheck -check-prefix=HOST %s
#include "Inputs/cuda.h"

2
clang/test/CodeGenCUDA/unnamed-types.cu
View file

@ -1,6 +1,6 @@
// RUN: %clang_cc1 -std=c++11 -x hip -triple x86_64-linux-gnu -aux-triple amdgpu-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=HOST
// RUN: %clang_cc1 -std=c++11 -x hip -triple x86_64-pc-windows-msvc -aux-triple amdgpu-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=MSVC
// RUN: %clang_cc1 -std=c++11 -x hip -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm %s -o - | FileCheck %s --check-prefix=DEVICE
// RUN: %clang_cc1 -std=c++11 -x hip -triple amdgpu-amd-amdhsa -fcuda-is-device -emit-llvm %s -o - | FileCheck %s --check-prefix=DEVICE
#include "Inputs/cuda.h"

2
clang/test/CodeGenCXX/address-space-cast-coerce.cpp
View file

@ -1,4 +1,4 @@
// RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -emit-llvm -o - | FileCheck %s
// RUN: %clang_cc1 %s -triple=amdgpu-amd-amdhsa -emit-llvm -o - | FileCheck %s
template<typename T, unsigned int n> struct my_vector_base;

2
clang/test/CodeGenCXX/address-space-cast.cpp
View file

@ -1,4 +1,4 @@
// RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -emit-llvm -o - | FileCheck %s
// RUN: %clang_cc1 %s -triple=amdgpu-amd-amdhsa -emit-llvm -o - | FileCheck %s
#define __private__ __attribute__((address_space(5)))

2
clang/test/CodeGenCXX/amdgcn-automatic-variable.cpp
View file

@ -1,5 +1,5 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
// RUN: %clang_cc1 -O0 -triple amdgcn---amdgiz -emit-llvm %s -o - | FileCheck %s
// RUN: %clang_cc1 -O0 -triple amdgpu---amdgiz -emit-llvm %s -o - | FileCheck %s
// CHECK-LABEL: @_Z5func1Pi(
// CHECK-NEXT: entry:

2
clang/test/CodeGenCXX/amdgcn-func-arg.cpp
View file

@ -1,5 +1,5 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
// RUN: %clang_cc1 -O0 -triple amdgcn -emit-llvm %s -o - | FileCheck %s
// RUN: %clang_cc1 -O0 -triple amdgpu -emit-llvm %s -o - | FileCheck %s
class A {
public:

2
clang/test/CodeGenCXX/amdgcn-string-literal.cpp
View file

@ -1,4 +1,4 @@
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -emit-llvm %s -o - | FileCheck %s
// CHECK: @.str = private unnamed_addr addrspace(4) constant [6 x i8] c"g_str\00", align 1
// CHECK: @g_str ={{.*}} addrspace(1) global ptr addrspacecast (ptr addrspace(4) @.str to ptr), align 8

2
clang/test/CodeGenCXX/amdgcn_declspec_get.cpp
View file

@ -1,4 +1,4 @@
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm -O3 -fdeclspec \
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -emit-llvm -O3 -fdeclspec \
// RUN: -disable-llvm-passes -o - %s | FileCheck %s
int get_x();

2
clang/test/CodeGenCXX/amdgpu-barrier-typeinfo.cpp
View file

@ -1,5 +1,5 @@
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn %s -emit-llvm -o - | FileCheck %s
// RUN: %clang_cc1 -triple amdgpu %s -emit-llvm -o - | FileCheck %s
namespace std { class type_info; };

2
clang/test/CodeGenCXX/amdgpu-buffer-rsrc-typeinfo.cpp
View file

@ -1,5 +1,5 @@
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn %s -emit-llvm -o - | FileCheck %s
// RUN: %clang_cc1 -triple amdgpu %s -emit-llvm -o - | FileCheck %s
namespace std { class type_info; };

2
clang/test/CodeGenCXX/amdgpu-image-rsrc-typeinfo.cpp
View file

@ -1,6 +1,6 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn %s -emit-llvm -o - | FileCheck %s
// RUN: %clang_cc1 -triple amdgpu %s -emit-llvm -o - | FileCheck %s
namespace std { class type_info; }
auto &a = typeid(__amdgpu_texture_t);
//// NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:

2
clang/test/CodeGenCXX/amdgpu-kernel-arg-pointer-type.cpp
View file

@ -1,6 +1,6 @@
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck --check-prefixes=COMMON,CHECK %s
// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -emit-llvm %s -o - | FileCheck --check-prefixes=COMMON,CHECK %s
// Derived from CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu by deleting references to HOST
// The original test passes the result through opt O2, but that seems to introduce invalid

2
clang/test/CodeGenCXX/bitfield-access-empty.cpp
View file

@ -21,7 +21,7 @@
// Configs that have expensive unaligned access
// Little Endian
// RUN: %clang_cc1 -triple=amdgcn-elf %s -emit-llvm -o /dev/null -fdump-record-layouts-simple | FileCheck --check-prefixes CHECK,LAYOUT %s
// RUN: %clang_cc1 -triple=amdgpu-elf %s -emit-llvm -o /dev/null -fdump-record-layouts-simple | FileCheck --check-prefixes CHECK,LAYOUT %s
// RUN: %clang_cc1 -triple=arc-elf %s -emit-llvm -o /dev/null -fdump-record-layouts-simple | FileCheck --check-prefixes CHECK,LAYOUT %s
// RUN: %clang_cc1 -triple=bpf %s -emit-llvm -o /dev/null -fdump-record-layouts-simple | FileCheck --check-prefixes CHECK,LAYOUT %s
// RUN: %clang_cc1 -triple=csky %s -emit-llvm -o /dev/null -fdump-record-layouts-simple | FileCheck --check-prefixes CHECK,LAYOUT %s

2
clang/test/CodeGenCXX/bitfield-access-tail.cpp
View file

@ -21,7 +21,7 @@
// Configs that have expensive unaligned access
// Little Endian
// RUN: %clang_cc1 -triple=amdgcn-elf %s -emit-llvm -o /dev/null -fdump-record-layouts-simple | FileCheck --check-prefixes CHECK,LAYOUT,LAYOUT64 %s
// RUN: %clang_cc1 -triple=amdgpu-elf %s -emit-llvm -o /dev/null -fdump-record-layouts-simple | FileCheck --check-prefixes CHECK,LAYOUT,LAYOUT64 %s
// RUN: %clang_cc1 -triple=arc-elf %s -emit-llvm -o /dev/null -fdump-record-layouts-simple | FileCheck --check-prefixes CHECK,LAYOUT,LAYOUT32 %s
// RUN: %clang_cc1 -triple=bpf %s -emit-llvm -o /dev/null -fdump-record-layouts-simple | FileCheck --check-prefixes CHECK,LAYOUT,LAYOUT64 %s
// RUN: %clang_cc1 -triple=csky %s -emit-llvm -o /dev/null -fdump-record-layouts-simple | FileCheck --check-prefixes CHECK,LAYOUT,LAYOUT32 %s

2
clang/test/CodeGenCXX/builtin-amdgcn-atomic-inc-dec.cpp
View file

@ -2,7 +2,7 @@
// REQUIRES: amdgpu-registered-target
// REQUIRES: spirv-registered-target
// RUN: %clang_cc1 %s -x hip -fcuda-is-device -emit-llvm -O0 -o - \
// RUN: -triple=amdgcn-amd-amdhsa | FileCheck --check-prefix=GCN %s
// RUN: -triple=amdgpu-amd-amdhsa | FileCheck --check-prefix=GCN %s
// RUN: %clang_cc1 %s -x hip -fcuda-is-device -emit-llvm -O0 -o - \
// RUN: -triple=spirv64-amd-amdhsa | FileCheck --check-prefix=AMDGCNSPIRV %s

2
clang/test/CodeGenCXX/builtin-amdgcn-fence.cpp
View file

@ -2,7 +2,7 @@
// REQUIRES: amdgpu-registered-target
// REQUIRES: spirv-registered-target
// RUN: %clang_cc1 %s -emit-llvm -O0 -o - \
// RUN: -triple=amdgcn-amd-amdhsa | FileCheck --check-prefix=GCN %s
// RUN: -triple=amdgpu-amd-amdhsa | FileCheck --check-prefix=GCN %s
// RUN: %clang_cc1 %s -emit-llvm -O0 -o - \
// RUN: -triple=spirv64-amd-amdhsa | FileCheck --check-prefix=AMDGCNSPIRV %s

2
clang/test/CodeGenCXX/conditional-temporaries.cpp
View file

@ -1,7 +1,7 @@
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -emit-llvm %s -o - -triple=x86_64-apple-darwin9 -O2 -disable-llvm-passes | FileCheck %s --check-prefixes=CHECK,CHECK-NOOPT
// RUN: %clang_cc1 -emit-llvm %s -o - -triple=x86_64-apple-darwin9 -O2 | FileCheck %s --check-prefixes=CHECK,CHECK-OPT
// RUN: %clang_cc1 -emit-llvm %s -o - -triple=amdgcn-amd-amdhsa -O2 | FileCheck %s --check-prefixes=CHECK,CHECK-OPT
// RUN: %clang_cc1 -emit-llvm %s -o - -triple=amdgpu-amd-amdhsa -O2 | FileCheck %s --check-prefixes=CHECK,CHECK-OPT
namespace {

2
clang/test/CodeGenCXX/cxx0x-initializer-stdinitializerlist.cpp
View file

@ -1,5 +1,5 @@
// RUN: %clang_cc1 -std=c++11 -triple x86_64-none-linux-gnu -fmerge-all-constants -emit-llvm -o - %s | FileCheck -check-prefixes=X86,CHECK %s
// RUN: %clang_cc1 -std=c++11 -triple amdgcn-amd-amdhsa -DNO_TLS -fmerge-all-constants -emit-llvm -o - %s | FileCheck -check-prefixes=AMDGCN,CHECK %s
// RUN: %clang_cc1 -std=c++11 -triple amdgpu-amd-amdhsa -DNO_TLS -fmerge-all-constants -emit-llvm -o - %s | FileCheck -check-prefixes=AMDGCN,CHECK %s
namespace std {
typedef decltype(sizeof(int)) size_t;

4
clang/test/CodeGenCXX/cxx11-extern-constexpr.cpp
View file

@ -1,7 +1,7 @@
// RUN: %clang_cc1 -std=c++11 %s -emit-llvm -o - -triple x86_64-linux-gnu | FileCheck %s --check-prefixes=X86,CXX11X86
// RUN: %clang_cc1 -std=c++1z %s -emit-llvm -o - -triple x86_64-linux-gnu | FileCheck %s --check-prefixes=X86,CXX17X86
// RUN: %clang_cc1 -std=c++11 %s -emit-llvm -o - -triple amdgcn-amd-amdhsa | FileCheck %s --check-prefixes=AMD,CXX11AMD
// RUN: %clang_cc1 -std=c++1z %s -emit-llvm -o - -triple amdgcn-amd-amdhsa | FileCheck %s --check-prefixes=AMD,CXX17AMD
// RUN: %clang_cc1 -std=c++11 %s -emit-llvm -o - -triple amdgpu-amd-amdhsa | FileCheck %s --check-prefixes=AMD,CXX11AMD
// RUN: %clang_cc1 -std=c++1z %s -emit-llvm -o - -triple amdgpu-amd-amdhsa | FileCheck %s --check-prefixes=AMD,CXX17AMD
struct A {
static const int Foo = 123;

2
clang/test/CodeGenCXX/dynamic-cast-address-space.cpp
View file

@ -1,5 +1,5 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals all --no-generate-body-for-unused-prefixes --version 4
// RUN: %clang_cc1 -I%S %s -triple amdgcn-amd-amdhsa -emit-llvm -fcxx-exceptions -fexceptions -o - | FileCheck %s
// RUN: %clang_cc1 -I%S %s -triple amdgpu-amd-amdhsa -emit-llvm -fcxx-exceptions -fexceptions -o - | FileCheck %s
// RUN: %clang_cc1 -I%S %s -triple spirv64-amd-amdhsa -emit-llvm -fcxx-exceptions -fexceptions -o - | FileCheck %s --check-prefix=WITH-NONZERO-DEFAULT-AS
struct A { virtual void f(); };

Some files were not shown because too many files have changed in this diff Show more