aboutsummaryrefslogtreecommitdiff
path: root/test/CodeGenCUDA
diff options
context:
space:
mode:
Diffstat (limited to 'test/CodeGenCUDA')
-rw-r--r--test/CodeGenCUDA/Inputs/cuda.h5
-rw-r--r--test/CodeGenCUDA/address-spaces.cu19
-rw-r--r--test/CodeGenCUDA/alias.cu3
-rw-r--r--test/CodeGenCUDA/amdgpu-kernel-attrs.cu37
-rw-r--r--test/CodeGenCUDA/convergent.cu4
-rw-r--r--test/CodeGenCUDA/device-stub.cu173
-rw-r--r--test/CodeGenCUDA/device-var-init.cu110
-rw-r--r--test/CodeGenCUDA/device-vtable.cu11
-rw-r--r--test/CodeGenCUDA/filter-decl.cu8
-rw-r--r--test/CodeGenCUDA/flush-denormals.cu15
-rw-r--r--test/CodeGenCUDA/kernel-amdgcn.cu41
-rw-r--r--test/CodeGenCUDA/kernel-args.cu39
-rw-r--r--test/CodeGenCUDA/kernel-call.cu13
-rw-r--r--test/CodeGenCUDA/library-builtin.cu22
-rw-r--r--test/CodeGenCUDA/link-device-bitcode.cu2
-rw-r--r--test/CodeGenCUDA/propagate-metadata.cu2
16 files changed, 398 insertions, 106 deletions
diff --git a/test/CodeGenCUDA/Inputs/cuda.h b/test/CodeGenCUDA/Inputs/cuda.h
index 9b9f43a1aaa9..3adbdc5b6d18 100644
--- a/test/CodeGenCUDA/Inputs/cuda.h
+++ b/test/CodeGenCUDA/Inputs/cuda.h
@@ -16,7 +16,12 @@ struct dim3 {
typedef struct cudaStream *cudaStream_t;
+#ifdef __HIP__
+int hipConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0,
+ cudaStream_t stream = 0);
+#else
int cudaConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0,
cudaStream_t stream = 0);
+#endif
extern "C" __device__ int printf(const char*, ...);
diff --git a/test/CodeGenCUDA/address-spaces.cu b/test/CodeGenCUDA/address-spaces.cu
index 449529bb24b4..58b0a4370762 100644
--- a/test/CodeGenCUDA/address-spaces.cu
+++ b/test/CodeGenCUDA/address-spaces.cu
@@ -1,4 +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
// Verifies Clang emits correct address spaces and addrspacecast instructions
// for CUDA code.
@@ -19,11 +20,11 @@ struct MyStruct {
int data2;
};
-// CHECK: @_ZZ5func0vE1a = internal addrspace(3) global %struct.MyStruct zeroinitializer
-// CHECK: @_ZZ5func1vE1a = internal addrspace(3) global float 0.000000e+00
-// CHECK: @_ZZ5func2vE1a = internal addrspace(3) global [256 x float] zeroinitializer
-// CHECK: @_ZZ5func3vE1a = internal addrspace(3) global float 0.000000e+00
-// CHECK: @_ZZ5func4vE1a = internal addrspace(3) global float 0.000000e+00
+// CHECK: @_ZZ5func0vE1a = internal addrspace(3) global %struct.MyStruct undef
+// CHECK: @_ZZ5func1vE1a = internal addrspace(3) global float undef
+// CHECK: @_ZZ5func2vE1a = internal addrspace(3) global [256 x float] undef
+// CHECK: @_ZZ5func3vE1a = internal addrspace(3) global float undef
+// CHECK: @_ZZ5func4vE1a = internal addrspace(3) global float undef
// CHECK: @b = addrspace(3) global float undef
__device__ void foo() {
@@ -48,7 +49,7 @@ __device__ void func0() {
ap->data2 = 2;
}
// CHECK: define void @_Z5func0v()
-// CHECK: store %struct.MyStruct* addrspacecast (%struct.MyStruct addrspace(3)* @_ZZ5func0vE1a to %struct.MyStruct*), %struct.MyStruct** %ap
+// CHECK: store %struct.MyStruct* addrspacecast (%struct.MyStruct addrspace(3)* @_ZZ5func0vE1a to %struct.MyStruct*), %struct.MyStruct** %{{.*}}
__device__ void callee(float *ap) {
*ap = 1.0f;
@@ -67,7 +68,7 @@ __device__ void func2() {
*ap = 1.0f;
}
// CHECK: define void @_Z5func2v()
-// CHECK: store float* getelementptr inbounds ([256 x float], [256 x float]* addrspacecast ([256 x float] addrspace(3)* @_ZZ5func2vE1a to [256 x float]*), i32 0, i32 128), float** %ap
+// CHECK: store float* getelementptr inbounds ([256 x float], [256 x float]* addrspacecast ([256 x float] addrspace(3)* @_ZZ5func2vE1a to [256 x float]*), i{{32|64}} 0, i{{32|64}} 128), float** %{{.*}}
__device__ void func3() {
__shared__ float a;
@@ -75,7 +76,7 @@ __device__ void func3() {
*ap = 1.0f;
}
// CHECK: define void @_Z5func3v()
-// CHECK: store float* addrspacecast (float addrspace(3)* @_ZZ5func3vE1a to float*), float** %ap
+// CHECK: store float* addrspacecast (float addrspace(3)* @_ZZ5func3vE1a to float*), float** %{{.*}}
__device__ void func4() {
__shared__ float a;
@@ -83,7 +84,7 @@ __device__ void func4() {
*ap = 1.0f;
}
// CHECK: define void @_Z5func4v()
-// CHECK: store float* addrspacecast (float addrspace(3)* @_ZZ5func4vE1a to float*), float** %ap
+// CHECK: store float* addrspacecast (float addrspace(3)* @_ZZ5func4vE1a to float*), float** %{{.*}}
__shared__ float b;
diff --git a/test/CodeGenCUDA/alias.cu b/test/CodeGenCUDA/alias.cu
index 6efff6b92aa8..56e595ed3639 100644
--- a/test/CodeGenCUDA/alias.cu
+++ b/test/CodeGenCUDA/alias.cu
@@ -1,8 +1,11 @@
// REQUIRES: x86-registered-target
// REQUIRES: nvptx-registered-target
+// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -emit-llvm \
// RUN: -o - %s | FileCheck %s
+// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn -emit-llvm \
+// RUN: -o - %s | FileCheck %s
#include "Inputs/cuda.h"
diff --git a/test/CodeGenCUDA/amdgpu-kernel-attrs.cu b/test/CodeGenCUDA/amdgpu-kernel-attrs.cu
new file mode 100644
index 000000000000..70eb9091d8d4
--- /dev/null
+++ b/test/CodeGenCUDA/amdgpu-kernel-attrs.cu
@@ -0,0 +1,37 @@
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
+// RUN: -fcuda-is-device -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple nvptx \
+// RUN: -fcuda-is-device -emit-llvm -o - %s | FileCheck %s \
+// RUN: -check-prefix=NAMD
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm \
+// RUN: -verify -o - %s | FileCheck -check-prefix=NAMD %s
+
+#include "Inputs/cuda.h"
+
+__attribute__((amdgpu_flat_work_group_size(32, 64))) // expected-no-diagnostics
+__global__ void flat_work_group_size_32_64() {
+// CHECK: define amdgpu_kernel void @_Z26flat_work_group_size_32_64v() [[FLAT_WORK_GROUP_SIZE_32_64:#[0-9]+]]
+}
+__attribute__((amdgpu_waves_per_eu(2))) // expected-no-diagnostics
+__global__ void waves_per_eu_2() {
+// CHECK: define amdgpu_kernel void @_Z14waves_per_eu_2v() [[WAVES_PER_EU_2:#[0-9]+]]
+}
+__attribute__((amdgpu_num_sgpr(32))) // expected-no-diagnostics
+__global__ void num_sgpr_32() {
+// CHECK: define amdgpu_kernel void @_Z11num_sgpr_32v() [[NUM_SGPR_32:#[0-9]+]]
+}
+__attribute__((amdgpu_num_vgpr(64))) // expected-no-diagnostics
+__global__ void num_vgpr_64() {
+// CHECK: define amdgpu_kernel void @_Z11num_vgpr_64v() [[NUM_VGPR_64:#[0-9]+]]
+}
+
+// Make sure this is silently accepted on other targets.
+// NAMD-NOT: "amdgpu-flat-work-group-size"
+// NAMD-NOT: "amdgpu-waves-per-eu"
+// NAMD-NOT: "amdgpu-num-vgpr"
+// NAMD-NOT: "amdgpu-num-sgpr"
+
+// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="32,64"
+// CHECK-DAG: attributes [[WAVES_PER_EU_2]] = { convergent noinline nounwind optnone "amdgpu-waves-per-eu"="2"
+// CHECK-DAG: attributes [[NUM_SGPR_32]] = { convergent noinline nounwind optnone "amdgpu-num-sgpr"="32"
+// CHECK-DAG: attributes [[NUM_VGPR_64]] = { convergent noinline nounwind optnone "amdgpu-num-vgpr"="64"
diff --git a/test/CodeGenCUDA/convergent.cu b/test/CodeGenCUDA/convergent.cu
index 62818f9e5af4..dd410cc644ac 100644
--- a/test/CodeGenCUDA/convergent.cu
+++ b/test/CodeGenCUDA/convergent.cu
@@ -2,11 +2,11 @@
// REQUIRES: nvptx-registered-target
// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -emit-llvm \
-// RUN: -disable-llvm-passes -o - %s | FileCheck -check-prefix DEVICE %s
+// RUN: -disable-llvm-passes -o - %s | FileCheck -allow-deprecated-dag-overlap -check-prefix DEVICE %s
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm \
// RUN: -disable-llvm-passes -o - %s | \
-// RUN: FileCheck -check-prefix HOST %s
+// RUN: FileCheck -allow-deprecated-dag-overlap -check-prefix HOST %s
#include "Inputs/cuda.h"
diff --git a/test/CodeGenCUDA/device-stub.cu b/test/CodeGenCUDA/device-stub.cu
index 3376803c5020..716381b7a826 100644
--- a/test/CodeGenCUDA/device-stub.cu
+++ b/test/CodeGenCUDA/device-stub.cu
@@ -1,33 +1,52 @@
// RUN: echo "GPU binary would be here" > %t
-// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -fcuda-include-gpubinary %t -o - | FileCheck %s
-// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -fcuda-include-gpubinary %t -o - -DNOGLOBALS \
-// RUN: | FileCheck %s -check-prefix=NOGLOBALS
-// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -o - | FileCheck %s -check-prefix=NOGPUBIN
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
+// RUN: -fcuda-include-gpubinary %t -o - \
+// RUN: | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,NORDC,CUDA,CUDANORDC
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
+// RUN: -fcuda-include-gpubinary %t -o - -DNOGLOBALS \
+// RUN: | FileCheck -allow-deprecated-dag-overlap %s -check-prefixes=NOGLOBALS,CUDANOGLOBALS
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
+// RUN: -fcuda-rdc -fcuda-include-gpubinary %t -o - \
+// RUN: | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,RDC,CUDA,CUDARDC
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -o - \
+// RUN: | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=NOGPUBIN
+
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
+// RUN: -fcuda-include-gpubinary %t -o - -x hip\
+// RUN: | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,NORDC,HIP
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
+// RUN: -fcuda-include-gpubinary %t -o - -DNOGLOBALS -x hip \
+// RUN: | FileCheck -allow-deprecated-dag-overlap %s -check-prefixes=NOGLOBALS,HIPNOGLOBALS
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
+// RUN: -fcuda-rdc -fcuda-include-gpubinary %t -o - -x hip \
+// RUN: | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,NORDC,HIP
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -o - -x hip\
+// RUN: | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=NOGPUBIN
#include "Inputs/cuda.h"
#ifndef NOGLOBALS
-// CHECK-DAG: @device_var = internal global i32
+// ALL-DAG: @device_var = internal global i32
__device__ int device_var;
-// CHECK-DAG: @constant_var = internal global i32
+// ALL-DAG: @constant_var = internal global i32
__constant__ int constant_var;
-// CHECK-DAG: @shared_var = internal global i32
+// ALL-DAG: @shared_var = internal global i32
__shared__ int shared_var;
// Make sure host globals don't get internalized...
-// CHECK-DAG: @host_var = global i32
+// ALL-DAG: @host_var = global i32
int host_var;
// ... and that extern vars remain external.
-// CHECK-DAG: @ext_host_var = external global i32
+// ALL-DAG: @ext_host_var = external global i32
extern int ext_host_var;
// Shadows for external device-side variables are *definitions* of
// those variables.
-// CHECK-DAG: @ext_device_var = internal global i32
+// ALL-DAG: @ext_device_var = internal global i32
extern __device__ int ext_device_var;
-// CHECK-DAG: @ext_device_var = internal global i32
+// ALL-DAG: @ext_device_var = internal global i32
extern __constant__ int ext_constant_var;
void use_pointers() {
@@ -43,72 +62,110 @@ void use_pointers() {
// Make sure that all parts of GPU code init/cleanup are there:
// * constant unnamed string with the kernel name
-// CHECK: private unnamed_addr constant{{.*}}kernelfunc{{.*}}\00"
+// ALL: private unnamed_addr constant{{.*}}kernelfunc{{.*}}\00"
// * constant unnamed string with GPU binary
-// CHECK: private unnamed_addr constant{{.*GPU binary would be here.*}}\00"
-// CHECK-SAME: section ".nv_fatbin", align 8
+// HIP: @[[FATBIN:__hip_fatbin]] = external constant i8, section ".hip_fatbin"
+// CUDA: @[[FATBIN:.*]] = private constant{{.*GPU binary would be here.*}}\00",
+// CUDANORDC-SAME: section ".nv_fatbin", align 8
+// CUDARDC-SAME: section "__nv_relfatbin", align 8
// * constant struct that wraps GPU binary
-// CHECK: @__cuda_fatbin_wrapper = internal constant { i32, i32, i8*, i8* }
-// CHECK-SAME: { i32 1180844977, i32 1, {{.*}}, i8* null }
-// CHECK-SAME: section ".nvFatBinSegment"
+// ALL: @__[[PREFIX:cuda|hip]]_fatbin_wrapper = internal constant
+// ALL-SAME: { i32, i32, i8*, i8* }
+// CUDA-SAME: { i32 1180844977, i32 1,
+// HIP-SAME: { i32 1212764230, i32 1,
+// CUDA-SAME: i8* getelementptr inbounds ({{.*}}@[[FATBIN]], i64 0, i64 0),
+// HIP-SAME: i8* @[[FATBIN]],
+// ALL-SAME: i8* null }
+// CUDA-SAME: section ".nvFatBinSegment"
+// HIP-SAME: section ".hipFatBinSegment"
// * variable to save GPU binary handle after initialization
-// CHECK: @__cuda_gpubin_handle = internal global i8** null
-// * Make sure our constructor/destructor was added to global ctor/dtor list.
-// CHECK: @llvm.global_ctors = appending global {{.*}}@__cuda_module_ctor
-// CHECK: @llvm.global_dtors = appending global {{.*}}@__cuda_module_dtor
+// CUDANORDC: @__[[PREFIX]]_gpubin_handle = internal global i8** null
+// HIP: @__[[PREFIX]]_gpubin_handle = linkonce global i8** null
+// * constant unnamed string with NVModuleID
+// RDC: [[MODULE_ID_GLOBAL:@.*]] = private constant
+// CUDARDC-SAME: c"[[MODULE_ID:.+]]\00", section "__nv_module_id", align 32
+// * Make sure our constructor was added to global ctor list.
+// ALL: @llvm.global_ctors = appending global {{.*}}@__[[PREFIX]]_module_ctor
+// * Alias to global symbol containing the NVModuleID.
+// RDC: @__fatbinwrap[[MODULE_ID]] = alias { i32, i32, i8*, i8* }
+// RDC-SAME: { i32, i32, i8*, i8* }* @__[[PREFIX]]_fatbin_wrapper
// Test that we build the correct number of calls to cudaSetupArgument followed
// by a call to cudaLaunch.
-// CHECK: define{{.*}}kernelfunc
-// CHECK: call{{.*}}cudaSetupArgument
-// CHECK: call{{.*}}cudaSetupArgument
-// CHECK: call{{.*}}cudaSetupArgument
-// CHECK: call{{.*}}cudaLaunch
+// ALL: define{{.*}}kernelfunc
+// ALL: call{{.*}}[[PREFIX]]SetupArgument
+// ALL: call{{.*}}[[PREFIX]]SetupArgument
+// ALL: call{{.*}}[[PREFIX]]SetupArgument
+// ALL: call{{.*}}[[PREFIX]]Launch
__global__ void kernelfunc(int i, int j, int k) {}
// Test that we've built correct kernel launch sequence.
-// CHECK: define{{.*}}hostfunc
-// CHECK: call{{.*}}cudaConfigureCall
-// CHECK: call{{.*}}kernelfunc
+// ALL: define{{.*}}hostfunc
+// ALL: call{{.*}}[[PREFIX]]ConfigureCall
+// ALL: call{{.*}}kernelfunc
void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); }
#endif
// Test that we've built a function to register kernels and global vars.
-// CHECK: define internal void @__cuda_register_globals
-// CHECK: call{{.*}}cudaRegisterFunction(i8** %0, {{.*}}kernelfunc
-// CHECK-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}device_var{{.*}}i32 0, i32 4, i32 0, i32 0
-// CHECK-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}constant_var{{.*}}i32 0, i32 4, i32 1, i32 0
-// CHECK-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}ext_device_var{{.*}}i32 1, i32 4, i32 0, i32 0
-// CHECK-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}ext_constant_var{{.*}}i32 1, i32 4, i32 1, i32 0
-// CHECK: ret void
+// ALL: define internal void @__[[PREFIX]]_register_globals
+// ALL: call{{.*}}[[PREFIX]]RegisterFunction(i8** %0, {{.*}}kernelfunc
+// ALL-DAG: call{{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}device_var{{.*}}i32 0, i32 4, i32 0, i32 0
+// ALL-DAG: call{{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}constant_var{{.*}}i32 0, i32 4, i32 1, i32 0
+// ALL-DAG: call{{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}ext_device_var{{.*}}i32 1, i32 4, i32 0, i32 0
+// ALL-DAG: call{{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}ext_constant_var{{.*}}i32 1, i32 4, i32 1, i32 0
+// ALL: ret void
+
+// Test that we've built a constructor.
+// ALL: define internal void @__[[PREFIX]]_module_ctor
+
+// In separate mode it calls __[[PREFIX]]RegisterFatBinary(&__[[PREFIX]]_fatbin_wrapper)
+// HIP only register fat binary once.
+// HIP: load i8**, i8*** @__hip_gpubin_handle
+// HIP-NEXT: icmp eq i8** {{.*}}, null
+// HIP-NEXT: br i1 {{.*}}, label %if, label %exit
+// HIP: if:
+// NORDC: call{{.*}}[[PREFIX]]RegisterFatBinary{{.*}}__[[PREFIX]]_fatbin_wrapper
+// .. stores return value in __[[PREFIX]]_gpubin_handle
+// NORDC-NEXT: store{{.*}}__[[PREFIX]]_gpubin_handle
+// .. and then calls __[[PREFIX]]_register_globals
+// HIP-NEXT: br label %exit
+// HIP: exit:
+// HIP-NEXT: load i8**, i8*** @__hip_gpubin_handle
+// NORDC-NEXT: call void @__[[PREFIX]]_register_globals
+// * In separate mode we also register a destructor.
+// NORDC-NEXT: call i32 @atexit(void (i8*)* @__[[PREFIX]]_module_dtor)
-// Test that we've built contructor..
-// CHECK: define internal void @__cuda_module_ctor
-// .. that calls __cudaRegisterFatBinary(&__cuda_fatbin_wrapper)
-// CHECK: call{{.*}}cudaRegisterFatBinary{{.*}}__cuda_fatbin_wrapper
-// .. stores return value in __cuda_gpubin_handle
-// CHECK-NEXT: store{{.*}}__cuda_gpubin_handle
-// .. and then calls __cuda_register_globals
-// CHECK-NEXT: call void @__cuda_register_globals
+// With relocatable device code we call __[[PREFIX]]RegisterLinkedBinary%NVModuleID%
+// RDC: call{{.*}}__[[PREFIX]]RegisterLinkedBinary[[MODULE_ID]](
+// RDC-SAME: __[[PREFIX]]_register_globals, {{.*}}__[[PREFIX]]_fatbin_wrapper
+// RDC-SAME: [[MODULE_ID_GLOBAL]]
// Test that we've created destructor.
-// CHECK: define internal void @__cuda_module_dtor
-// CHECK: load{{.*}}__cuda_gpubin_handle
-// CHECK-NEXT: call void @__cudaUnregisterFatBinary
+// NORDC: define internal void @__[[PREFIX]]_module_dtor
+// NORDC: load{{.*}}__[[PREFIX]]_gpubin_handle
+// CUDANORDC-NEXT: call void @__[[PREFIX]]UnregisterFatBinary
+// HIP-NEXT: icmp ne i8** {{.*}}, null
+// HIP-NEXT: br i1 {{.*}}, label %if, label %exit
+// HIP: if:
+// HIP-NEXT: call void @__[[PREFIX]]UnregisterFatBinary
+// HIP-NEXT: store i8** null, i8*** @__hip_gpubin_handle
+// HIP-NEXT: br label %exit
+// HIP: exit:
-// There should be no __cuda_register_globals if we have no
+// There should be no __[[PREFIX]]_register_globals if we have no
// device-side globals, but we still need to register GPU binary.
// Skip GPU binary string first.
-// NOGLOBALS: @0 = private unnamed_addr constant{{.*}}
-// NOGLOBALS-NOT: define internal void @__cuda_register_globals
-// NOGLOBALS: define internal void @__cuda_module_ctor
-// NOGLOBALS: call{{.*}}cudaRegisterFatBinary{{.*}}__cuda_fatbin_wrapper
-// NOGLOBALS-NOT: call void @__cuda_register_globals
-// NOGLOBALS: define internal void @__cuda_module_dtor
-// NOGLOBALS: call void @__cudaUnregisterFatBinary
+// CUDANOGLOBALS: @{{.*}} = private constant{{.*}}
+// HIPNOGLOBALS: @{{.*}} = external constant{{.*}}
+// NOGLOBALS-NOT: define internal void @__{{.*}}_register_globals
+// NOGLOBALS: define internal void @__[[PREFIX:cuda|hip]]_module_ctor
+// NOGLOBALS: call{{.*}}[[PREFIX]]RegisterFatBinary{{.*}}__[[PREFIX]]_fatbin_wrapper
+// NOGLOBALS-NOT: call void @__[[PREFIX]]_register_globals
+// NOGLOBALS: define internal void @__[[PREFIX]]_module_dtor
+// NOGLOBALS: call void @__[[PREFIX]]UnregisterFatBinary
// There should be no constructors/destructors if we have no GPU binary.
-// NOGPUBIN-NOT: define internal void @__cuda_register_globals
-// NOGPUBIN-NOT: define internal void @__cuda_module_ctor
-// NOGPUBIN-NOT: define internal void @__cuda_module_dtor
+// NOGPUBIN-NOT: define internal void @__[[PREFIX]]_register_globals
+// NOGPUBIN-NOT: define internal void @__[[PREFIX]]_module_ctor
+// NOGPUBIN-NOT: define internal void @__[[PREFIX]]_module_dtor
diff --git a/test/CodeGenCUDA/device-var-init.cu b/test/CodeGenCUDA/device-var-init.cu
index 0f4c64813406..9f788b764fbc 100644
--- a/test/CodeGenCUDA/device-var-init.cu
+++ b/test/CodeGenCUDA/device-var-init.cu
@@ -1,10 +1,14 @@
// REQUIRES: nvptx-registered-target
+// REQUIRES: amdgpu-registered-target
// Make sure we don't allow dynamic initialization for device
// variables, but accept empty constructors allowed by CUDA.
// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++11 \
-// RUN: -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck %s
+// RUN: -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck -check-prefixes=CHECK,NVPTX %s
+
+// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -std=c++11 \
+// RUN: -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck -check-prefixes=CHECK,AMDGCN %s
#ifdef __clang__
#include "Inputs/cuda.h"
@@ -105,50 +109,102 @@ __shared__ EC_I_EC s_ec_i_ec;
__constant__ EC_I_EC c_ec_i_ec;
// CHECK: @c_ec_i_ec = addrspace(4) externally_initialized global %struct.EC_I_EC zeroinitializer,
+// CHECK: @_ZZ2dfvE4s_ec = internal addrspace(3) global %struct.EC undef
+// CHECK: @_ZZ2dfvE5s_etc = internal addrspace(3) global %struct.ETC undef
+
// We should not emit global initializers for device-side variables.
// CHECK-NOT: @__cxx_global_var_init
// Make sure that initialization restrictions do not apply to local
// variables.
__device__ void df() {
+ // NVPTX: %[[ec:.*]] = alloca %struct.EC
+ // NVPTX: %[[ed:.*]] = alloca %struct.ED
+ // NVPTX: %[[ecd:.*]] = alloca %struct.ECD
+ // NVPTX: %[[etc:.*]] = alloca %struct.ETC
+ // NVPTX: %[[uc:.*]] = alloca %struct.UC
+ // NVPTX: %[[ud:.*]] = alloca %struct.UD
+ // NVPTX: %[[eci:.*]] = alloca %struct.ECI
+ // NVPTX: %[[nec:.*]] = alloca %struct.NEC
+ // NVPTX: %[[ned:.*]] = alloca %struct.NED
+ // NVPTX: %[[ncv:.*]] = alloca %struct.NCV
+ // NVPTX: %[[vd:.*]] = alloca %struct.VD
+ // NVPTX: %[[ncf:.*]] = alloca %struct.NCF
+ // NVPTX: %[[ncfs:.*]] = alloca %struct.NCFS
+ // NVPTX: %[[utc:.*]] = alloca %struct.UTC
+ // NVPTX: %[[netc:.*]] = alloca %struct.NETC
+ // NVPTX: %[[ec_i_ec:.*]] = alloca %struct.EC_I_EC
+ // NVPTX: %[[ec_i_ec1:.*]] = alloca %struct.EC_I_EC1
+ // NVPTX: %[[t_v_t:.*]] = alloca %struct.T_V_T
+ // NVPTX: %[[t_b_nec:.*]] = alloca %struct.T_B_NEC
+ // NVPTX: %[[t_f_nec:.*]] = alloca %struct.T_F_NEC
+ // NVPTX: %[[t_fa_nec:.*]] = alloca %struct.T_FA_NEC
+ // NVPTX: %[[t_b_ned:.*]] = alloca %struct.T_B_NED
+ // NVPTX: %[[t_f_ned:.*]] = alloca %struct.T_F_NED
+ // NVPTX: %[[t_fa_ned:.*]] = alloca %struct.T_FA_NED
+ // AMDGCN: %[[ec:.*]] = addrspacecast %struct.EC addrspace(5)* %ec to %struct.EC*
+ // AMDGCN: %[[ed:.*]] = addrspacecast %struct.ED addrspace(5)* %ed to %struct.ED*
+ // AMDGCN: %[[ecd:.*]] = addrspacecast %struct.ECD addrspace(5)* %ecd to %struct.ECD*
+ // AMDGCN: %[[etc:.*]] = addrspacecast %struct.ETC addrspace(5)* %etc to %struct.ETC*
+ // AMDGCN: %[[uc:.*]] = addrspacecast %struct.UC addrspace(5)* %uc to %struct.UC*
+ // AMDGCN: %[[ud:.*]] = addrspacecast %struct.UD addrspace(5)* %ud to %struct.UD*
+ // AMDGCN: %[[eci:.*]] = addrspacecast %struct.ECI addrspace(5)* %eci to %struct.ECI*
+ // AMDGCN: %[[nec:.*]] = addrspacecast %struct.NEC addrspace(5)* %nec to %struct.NEC*
+ // AMDGCN: %[[ned:.*]] = addrspacecast %struct.NED addrspace(5)* %ned to %struct.NED*
+ // AMDGCN: %[[ncv:.*]] = addrspacecast %struct.NCV addrspace(5)* %ncv to %struct.NCV*
+ // AMDGCN: %[[vd:.*]] = addrspacecast %struct.VD addrspace(5)* %vd to %struct.VD*
+ // AMDGCN: %[[ncf:.*]] = addrspacecast %struct.NCF addrspace(5)* %ncf to %struct.NCF*
+ // AMDGCN: %[[ncfs:.*]] = addrspacecast %struct.NCFS addrspace(5)* %ncfs to %struct.NCFS*
+ // AMDGCN: %[[utc:.*]] = addrspacecast %struct.UTC addrspace(5)* %utc to %struct.UTC*
+ // AMDGCN: %[[netc:.*]] = addrspacecast %struct.NETC addrspace(5)* %netc to %struct.NETC*
+ // AMDGCN: %[[ec_i_ec:.*]] = addrspacecast %struct.EC_I_EC addrspace(5)* %ec_i_ec to %struct.EC_I_EC*
+ // AMDGCN: %[[ec_i_ec1:.*]] = addrspacecast %struct.EC_I_EC1 addrspace(5)* %ec_i_ec1 to %struct.EC_I_EC1*
+ // AMDGCN: %[[t_v_t:.*]] = addrspacecast %struct.T_V_T addrspace(5)* %t_v_t to %struct.T_V_T*
+ // AMDGCN: %[[t_b_nec:.*]] = addrspacecast %struct.T_B_NEC addrspace(5)* %t_b_nec to %struct.T_B_NEC*
+ // AMDGCN: %[[t_f_nec:.*]] = addrspacecast %struct.T_F_NEC addrspace(5)* %t_f_nec to %struct.T_F_NEC*
+ // AMDGCN: %[[t_fa_nec:.*]] = addrspacecast %struct.T_FA_NEC addrspace(5)* %t_fa_nec to %struct.T_FA_NEC*
+ // AMDGCN: %[[t_b_ned:.*]] = addrspacecast %struct.T_B_NED addrspace(5)* %t_b_ned to %struct.T_B_NED*
+ // AMDGCN: %[[t_f_ned:.*]] = addrspacecast %struct.T_F_NED addrspace(5)* %t_f_ned to %struct.T_F_NED*
+ // AMDGCN: %[[t_fa_ned:.*]] = addrspacecast %struct.T_FA_NED addrspace(5)* %t_fa_ned to %struct.T_FA_NED*
+
T t;
// CHECK-NOT: call
EC ec;
- // CHECK: call void @_ZN2ECC1Ev(%struct.EC* %ec)
+ // CHECK: call void @_ZN2ECC1Ev(%struct.EC* %[[ec]])
ED ed;
// CHECK-NOT: call
ECD ecd;
- // CHECK: call void @_ZN3ECDC1Ev(%struct.ECD* %ecd)
+ // CHECK: call void @_ZN3ECDC1Ev(%struct.ECD* %[[ecd]])
ETC etc;
- // CHECK: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* %etc)
+ // CHECK: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* %[[etc]])
UC uc;
// undefined constructor -- not allowed
- // CHECK: call void @_ZN2UCC1Ev(%struct.UC* %uc)
+ // CHECK: call void @_ZN2UCC1Ev(%struct.UC* %[[uc]])
UD ud;
// undefined destructor -- not allowed
// CHECK-NOT: call
ECI eci;
// empty constructor w/ initializer list -- not allowed
- // CHECK: call void @_ZN3ECIC1Ev(%struct.ECI* %eci)
+ // CHECK: call void @_ZN3ECIC1Ev(%struct.ECI* %[[eci]])
NEC nec;
// non-empty constructor -- not allowed
- // CHECK: call void @_ZN3NECC1Ev(%struct.NEC* %nec)
+ // CHECK: call void @_ZN3NECC1Ev(%struct.NEC* %[[nec]])
// non-empty destructor -- not allowed
NED ned;
// no-constructor, virtual method -- not allowed
- // CHECK: call void @_ZN3NCVC1Ev(%struct.NCV* %ncv)
+ // CHECK: call void @_ZN3NCVC1Ev(%struct.NCV* %[[ncv]])
NCV ncv;
// CHECK-NOT: call
VD vd;
- // CHECK: call void @_ZN2VDC1Ev(%struct.VD* %vd)
+ // CHECK: call void @_ZN2VDC1Ev(%struct.VD* %[[vd]])
NCF ncf;
- // CHECK: call void @_ZN3NCFC1Ev(%struct.NCF* %ncf)
+ // CHECK: call void @_ZN3NCFC1Ev(%struct.NCF* %[[ncf]])
NCFS ncfs;
- // CHECK: call void @_ZN4NCFSC1Ev(%struct.NCFS* %ncfs)
+ // CHECK: call void @_ZN4NCFSC1Ev(%struct.NCFS* %[[ncfs]])
UTC utc;
- // CHECK: call void @_ZN3UTCC1IJEEEDpT_(%struct.UTC* %utc)
+ // CHECK: call void @_ZN3UTCC1IJEEEDpT_(%struct.UTC* %[[utc]])
NETC netc;
- // CHECK: call void @_ZN4NETCC1IJEEEDpT_(%struct.NETC* %netc)
+ // CHECK: call void @_ZN4NETCC1IJEEEDpT_(%struct.NETC* %[[netc]])
T_B_T t_b_t;
// CHECK-NOT: call
T_F_T t_f_t;
@@ -156,17 +212,17 @@ __device__ void df() {
T_FA_T t_fa_t;
// CHECK-NOT: call
EC_I_EC ec_i_ec;
- // CHECK: call void @_ZN7EC_I_ECC1Ev(%struct.EC_I_EC* %ec_i_ec)
+ // CHECK: call void @_ZN7EC_I_ECC1Ev(%struct.EC_I_EC* %[[ec_i_ec]])
EC_I_EC1 ec_i_ec1;
- // CHECK: call void @_ZN8EC_I_EC1C1Ev(%struct.EC_I_EC1* %ec_i_ec1)
+ // CHECK: call void @_ZN8EC_I_EC1C1Ev(%struct.EC_I_EC1* %[[ec_i_ec1]])
T_V_T t_v_t;
- // CHECK: call void @_ZN5T_V_TC1Ev(%struct.T_V_T* %t_v_t)
+ // CHECK: call void @_ZN5T_V_TC1Ev(%struct.T_V_T* %[[t_v_t]])
T_B_NEC t_b_nec;
- // CHECK: call void @_ZN7T_B_NECC1Ev(%struct.T_B_NEC* %t_b_nec)
+ // CHECK: call void @_ZN7T_B_NECC1Ev(%struct.T_B_NEC* %[[t_b_nec]])
T_F_NEC t_f_nec;
- // CHECK: call void @_ZN7T_F_NECC1Ev(%struct.T_F_NEC* %t_f_nec)
+ // CHECK: call void @_ZN7T_F_NECC1Ev(%struct.T_F_NEC* %[[t_f_nec]])
T_FA_NEC t_fa_nec;
- // CHECK: call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC* %t_fa_nec)
+ // CHECK: call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC* %[[t_fa_nec]])
T_B_NED t_b_ned;
// CHECK-NOT: call
T_F_NED t_f_ned;
@@ -182,14 +238,14 @@ __device__ void df() {
df(); // CHECK: call void @_Z2dfv()
// Verify that we only call non-empty destructors
- // CHECK-NEXT: call void @_ZN8T_FA_NEDD1Ev(%struct.T_FA_NED* %t_fa_ned)
- // CHECK-NEXT: call void @_ZN7T_F_NEDD1Ev(%struct.T_F_NED* %t_f_ned)
- // CHECK-NEXT: call void @_ZN7T_B_NEDD1Ev(%struct.T_B_NED* %t_b_ned)
- // CHECK-NEXT: call void @_ZN2VDD1Ev(%struct.VD* %vd)
- // CHECK-NEXT: call void @_ZN3NEDD1Ev(%struct.NED* %ned)
- // CHECK-NEXT: call void @_ZN2UDD1Ev(%struct.UD* %ud)
- // CHECK-NEXT: call void @_ZN3ECDD1Ev(%struct.ECD* %ecd)
- // CHECK-NEXT: call void @_ZN2EDD1Ev(%struct.ED* %ed)
+ // CHECK-NEXT: call void @_ZN8T_FA_NEDD1Ev(%struct.T_FA_NED* %[[t_fa_ned]])
+ // CHECK-NEXT: call void @_ZN7T_F_NEDD1Ev(%struct.T_F_NED* %[[t_f_ned]])
+ // CHECK-NEXT: call void @_ZN7T_B_NEDD1Ev(%struct.T_B_NED* %[[t_b_ned]])
+ // CHECK-NEXT: call void @_ZN2VDD1Ev(%struct.VD* %[[vd]])
+ // CHECK-NEXT: call void @_ZN3NEDD1Ev(%struct.NED* %[[ned]])
+ // CHECK-NEXT: call void @_ZN2UDD1Ev(%struct.UD* %[[ud]])
+ // CHECK-NEXT: call void @_ZN3ECDD1Ev(%struct.ECD* %[[ecd]])
+ // CHECK-NEXT: call void @_ZN2EDD1Ev(%struct.ED* %[[ed]])
// CHECK-NEXT: ret void
}
diff --git a/test/CodeGenCUDA/device-vtable.cu b/test/CodeGenCUDA/device-vtable.cu
index 9730e404caa4..a7307dcdaf9b 100644
--- a/test/CodeGenCUDA/device-vtable.cu
+++ b/test/CodeGenCUDA/device-vtable.cu
@@ -19,7 +19,9 @@ struct H {
//CHECK-HOST: @_ZTV1H =
//CHECK-HOST-SAME: @_ZN1H6methodEv
//CHECK-DEVICE-NOT: @_ZTV1H =
-
+//CHECK-DEVICE-NOT: @_ZTVN10__cxxabiv117__class_type_infoE
+//CHECK-DEVICE-NOT: @_ZTS1H
+//CHECK-DEVICE-NOT: @_ZTI1H
struct D {
__device__ virtual void method();
};
@@ -27,7 +29,9 @@ struct D {
//CHECK-DEVICE: @_ZTV1D
//CHECK-DEVICE-SAME: @_ZN1D6methodEv
//CHECK-HOST-NOT: @_ZTV1D
-
+//CHECK-DEVICE-NOT: @_ZTVN10__cxxabiv117__class_type_infoE
+//CHECK-DEVICE-NOT: @_ZTS1D
+//CHECK-DEVICE-NOT: @_ZTI1D
// This is the case with mixed host and device virtual methods. It's
// impossible to emit a valid vtable in that case because only host or
// only device methods would be available during host or device
@@ -45,6 +49,9 @@ struct HD {
// CHECK-HOST-NOT: @_ZN2HD8d_methodEv
// CHECK-HOST-SAME: null
// CHECK-BOTH-SAME: ]
+// CHECK-DEVICE-NOT: @_ZTVN10__cxxabiv117__class_type_infoE
+// CHECK-DEVICE-NOT: @_ZTS2HD
+// CHECK-DEVICE-NOT: @_ZTI2HD
void H::method() {}
//CHECK-HOST: define void @_ZN1H6methodEv
diff --git a/test/CodeGenCUDA/filter-decl.cu b/test/CodeGenCUDA/filter-decl.cu
index bc744a07a330..0f4691f7c8aa 100644
--- a/test/CodeGenCUDA/filter-decl.cu
+++ b/test/CodeGenCUDA/filter-decl.cu
@@ -10,18 +10,18 @@
__asm__("file scope asm is host only");
// CHECK-HOST: constantdata = internal global
-// CHECK-DEVICE: constantdata = externally_initialized global
+// CHECK-DEVICE: constantdata = {{(dso_local )?}}externally_initialized global
__constant__ char constantdata[256];
// CHECK-HOST: devicedata = internal global
-// CHECK-DEVICE: devicedata = externally_initialized global
+// CHECK-DEVICE: devicedata = {{(dso_local )?}}externally_initialized global
__device__ char devicedata[256];
// CHECK-HOST: shareddata = internal global
-// CHECK-DEVICE: shareddata = global
+// CHECK-DEVICE: shareddata = {{(dso_local )?}}global
__shared__ char shareddata[256];
-// CHECK-HOST: hostdata = global
+// CHECK-HOST: hostdata = {{(dso_local )?}}global
// CHECK-DEVICE-NOT: hostdata = global
char hostdata[256];
diff --git a/test/CodeGenCUDA/flush-denormals.cu b/test/CodeGenCUDA/flush-denormals.cu
index 94285f1f33ce..13f2e3b2871e 100644
--- a/test/CodeGenCUDA/flush-denormals.cu
+++ b/test/CodeGenCUDA/flush-denormals.cu
@@ -5,6 +5,13 @@
// RUN: -triple nvptx-nvidia-cuda -emit-llvm -o - %s | \
// RUN: FileCheck %s -check-prefix CHECK -check-prefix FTZ
+// RUN: %clang_cc1 -fcuda-is-device -x hip \
+// RUN: -triple amdgcn-amd-amdhsa -target-cpu gfx900 -emit-llvm -o - %s | \
+// RUN: FileCheck %s -check-prefix CHECK -check-prefix AMDNOFTZ
+// RUN: %clang_cc1 -fcuda-is-device -x hip -fcuda-flush-denormals-to-zero \
+// RUN: -triple amdgcn-amd-amdhsa -target-cpu gfx900 -emit-llvm -o - %s | \
+// RUN: FileCheck %s -check-prefix CHECK -check-prefix AMDFTZ
+
#include "Inputs/cuda.h"
// Checks that device function calls get emitted with the "ntpvx-f32ftz"
@@ -12,11 +19,19 @@
// -fcuda-flush-denormals-to-zero. Further, check that we reflect the presence
// or absence of -fcuda-flush-denormals-to-zero in a module flag.
+// AMDGCN targets always have +fp64-fp16-denormals.
+// AMDGCN targets without fast FMAF (e.g. gfx803) always have +fp32-denormals.
+// For AMDGCN target with fast FMAF (e.g. gfx900), it has +fp32-denormals
+// by default and -fp32-denormals when there is option
+// -fcuda-flush-denormals-to-zero.
+
// CHECK-LABEL: define void @foo() #0
extern "C" __device__ void foo() {}
// FTZ: attributes #0 = {{.*}} "nvptx-f32ftz"="true"
// NOFTZ-NOT: attributes #0 = {{.*}} "nvptx-f32ftz"
+// AMDNOFTZ: attributes #0 = {{.*}}+fp32-denormals{{.*}}+fp64-fp16-denormals
+// AMDFTZ: attributes #0 = {{.*}}+fp64-fp16-denormals{{.*}}-fp32-denormals
// FTZ:!llvm.module.flags = !{{{.*}}[[MODFLAG:![0-9]+]]}
// FTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 1}
diff --git a/test/CodeGenCUDA/kernel-amdgcn.cu b/test/CodeGenCUDA/kernel-amdgcn.cu
new file mode 100644
index 000000000000..ffa6c9549f07
--- /dev/null
+++ b/test/CodeGenCUDA/kernel-amdgcn.cu
@@ -0,0 +1,41 @@
+// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -emit-llvm %s -o - | FileCheck %s
+#include "Inputs/cuda.h"
+
+// CHECK: define amdgpu_kernel void @_ZN1A6kernelEv
+class A {
+public:
+ static __global__ void kernel(){}
+};
+
+// CHECK: define void @_Z10non_kernelv
+__device__ void non_kernel(){}
+
+// CHECK: define amdgpu_kernel void @_Z6kerneli
+__global__ void kernel(int x) {
+ non_kernel();
+}
+
+// CHECK: define amdgpu_kernel void @_Z11EmptyKernelIvEvv
+template <typename T>
+__global__ void EmptyKernel(void) {}
+
+struct Dummy {
+ /// Type definition of the EmptyKernel kernel entry point
+ typedef void (*EmptyKernelPtr)();
+ EmptyKernelPtr Empty() { return EmptyKernel<void>; }
+};
+
+// CHECK: define amdgpu_kernel void @_Z15template_kernelI1AEvT_
+template<class T>
+__global__ void template_kernel(T x) {}
+
+void launch(void *f);
+
+int main() {
+ Dummy D;
+ launch((void*)A::kernel);
+ launch((void*)kernel);
+ launch((void*)template_kernel<A>);
+ launch((void*)D.Empty());
+ return 0;
+}
diff --git a/test/CodeGenCUDA/kernel-args.cu b/test/CodeGenCUDA/kernel-args.cu
new file mode 100644
index 000000000000..d0986629f8d4
--- /dev/null
+++ b/test/CodeGenCUDA/kernel-args.cu
@@ -0,0 +1,39 @@
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
+// RUN: -emit-llvm %s -o - | FileCheck -check-prefix=AMDGCN %s
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda- -fcuda-is-device \
+// RUN: -emit-llvm %s -o - | FileCheck -check-prefix=NVPTX %s
+#include "Inputs/cuda.h"
+
+struct A {
+ int a[32];
+};
+
+// AMDGCN: define amdgpu_kernel void @_Z6kernel1A(%struct.A %x.coerce)
+// NVPTX: define void @_Z6kernel1A(%struct.A* byval align 4 %x)
+__global__ void kernel(A x) {
+}
+
+class Kernel {
+public:
+ // AMDGCN: define amdgpu_kernel void @_ZN6Kernel12memberKernelE1A(%struct.A %x.coerce)
+ // NVPTX: define void @_ZN6Kernel12memberKernelE1A(%struct.A* byval align 4 %x)
+ static __global__ void memberKernel(A x){}
+ template<typename T> static __global__ void templateMemberKernel(T x) {}
+};
+
+
+template <typename T>
+__global__ void templateKernel(T x) {}
+
+void launch(void*);
+
+void test() {
+ Kernel K;
+ // AMDGCN: define amdgpu_kernel void @_Z14templateKernelI1AEvT_(%struct.A %x.coerce)
+ // NVPTX: define void @_Z14templateKernelI1AEvT_(%struct.A* byval align 4 %x)
+ launch((void*)templateKernel<A>);
+
+ // AMDGCN: define amdgpu_kernel void @_ZN6Kernel20templateMemberKernelI1AEEvT_(%struct.A %x.coerce)
+ // NVPTX: define void @_ZN6Kernel20templateMemberKernelI1AEEvT_(%struct.A* byval align 4 %x)
+ launch((void*)Kernel::templateMemberKernel<A>);
+}
diff --git a/test/CodeGenCUDA/kernel-call.cu b/test/CodeGenCUDA/kernel-call.cu
index 9b849db908f8..43d08dfaf89b 100644
--- a/test/CodeGenCUDA/kernel-call.cu
+++ b/test/CodeGenCUDA/kernel-call.cu
@@ -1,11 +1,20 @@
-// RUN: %clang_cc1 -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -emit-llvm %s -o - | FileCheck %s --check-prefixes=CUDA,CHECK
+// RUN: %clang_cc1 -x hip -emit-llvm %s -o - | FileCheck %s --check-prefixes=HIP,CHECK
+
#include "Inputs/cuda.h"
+// CHECK-LABEL: define{{.*}}g1
+// HIP: call{{.*}}hipSetupArgument
+// HIP: call{{.*}}hipLaunchByPtr
+// CUDA: call{{.*}}cudaSetupArgument
+// CUDA: call{{.*}}cudaLaunch
__global__ void g1(int x) {}
+// CHECK-LABEL: define{{.*}}main
int main(void) {
- // CHECK: call{{.*}}cudaConfigureCall
+ // HIP: call{{.*}}hipConfigureCall
+ // CUDA: call{{.*}}cudaConfigureCall
// CHECK: icmp
// CHECK: br
// CHECK: call{{.*}}g1
diff --git a/test/CodeGenCUDA/library-builtin.cu b/test/CodeGenCUDA/library-builtin.cu
new file mode 100644
index 000000000000..4804c75cde38
--- /dev/null
+++ b/test/CodeGenCUDA/library-builtin.cu
@@ -0,0 +1,22 @@
+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -o - %s | \
+// RUN: FileCheck --check-prefixes=HOST,BOTH %s
+// RUN: %clang_cc1 -fcuda-is-device -triple nvptx64-nvidia-cuda \
+// RUN: -emit-llvm -o - %s | FileCheck %s --check-prefixes=DEVICE,BOTH
+
+// BOTH-LABEL: define float @logf(float
+
+// logf() should be calling itself recursively as we don't have any standard
+// library on device side.
+// DEVICE: call float @logf(float
+extern "C" __attribute__((device)) float logf(float __x) { return logf(__x); }
+
+// NOTE: this case is to illustrate the expected differences in behavior between
+// the host and device. In general we do not mess with host-side standard
+// library.
+//
+// Host is assumed to have standard library, so logf() calls LLVM intrinsic.
+// HOST: call float @llvm.log.f32(float
+extern "C" float logf(float __x) { return logf(__x); }
diff --git a/test/CodeGenCUDA/link-device-bitcode.cu b/test/CodeGenCUDA/link-device-bitcode.cu
index 869fcb1bc938..b307838ae82d 100644
--- a/test/CodeGenCUDA/link-device-bitcode.cu
+++ b/test/CodeGenCUDA/link-device-bitcode.cu
@@ -31,7 +31,7 @@
// Make sure NVVMReflect pass is enabled in NVPTX back-end.
// RUN: %clang_cc1 -triple nvptx-unknown-cuda -fcuda-is-device \
// RUN: -mlink-cuda-bitcode %t.bc -S -o /dev/null %s \
-// RUN: -backend-option -debug-pass=Structure 2>&1 \
+// RUN: -mllvm -debug-pass=Structure 2>&1 \
// RUN: | FileCheck %s -check-prefix CHECK-REFLECT
#include "Inputs/cuda.h"
diff --git a/test/CodeGenCUDA/propagate-metadata.cu b/test/CodeGenCUDA/propagate-metadata.cu
index f8db765099ef..1616cb933748 100644
--- a/test/CodeGenCUDA/propagate-metadata.cu
+++ b/test/CodeGenCUDA/propagate-metadata.cu
@@ -29,7 +29,7 @@
// RUN: -fcuda-is-device -menable-unsafe-fp-math -triple nvptx-unknown-unknown \
// RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=FAST
-// Wrap everything in extern "C" so we don't ahve to worry about name mangling
+// Wrap everything in extern "C" so we don't have to worry about name mangling
// in the IR.
extern "C" {
#ifdef LIB
generated by cgit v1.3 (git 2.53.0) at 2026-07-29 22:59:19 +0000