diff options
Diffstat (limited to 'test/CodeGenCUDA')
| -rw-r--r-- | test/CodeGenCUDA/Inputs/cuda.h | 13 | ||||
| -rw-r--r-- | test/CodeGenCUDA/amdgpu-visibility.cu | 21 | ||||
| -rw-r--r-- | test/CodeGenCUDA/debug-info-address-class.cu | 25 | ||||
| -rw-r--r-- | test/CodeGenCUDA/debug-info-template.cu | 10 | ||||
| -rw-r--r-- | test/CodeGenCUDA/device-stub.cu | 150 | ||||
| -rw-r--r-- | test/CodeGenCUDA/kernel-args-alignment.cu | 16 | ||||
| -rw-r--r-- | test/CodeGenCUDA/kernel-call.cu | 17 | ||||
| -rw-r--r-- | test/CodeGenCUDA/kernel-stub-name.cu | 20 | ||||
| -rw-r--r-- | test/CodeGenCUDA/types.cu | 10 |
9 files changed, 224 insertions, 58 deletions
diff --git a/test/CodeGenCUDA/Inputs/cuda.h b/test/CodeGenCUDA/Inputs/cuda.h index 3adbdc5b6d..0fd175765a 100644 --- a/test/CodeGenCUDA/Inputs/cuda.h +++ b/test/CodeGenCUDA/Inputs/cuda.h @@ -15,13 +15,20 @@ struct dim3 { }; typedef struct cudaStream *cudaStream_t; - +typedef enum cudaError {} cudaError_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); +extern "C" int cudaConfigureCall(dim3 gridSize, dim3 blockSize, + size_t sharedSize = 0, + cudaStream_t stream = 0); +extern "C" int __cudaPushCallConfiguration(dim3 gridSize, dim3 blockSize, + size_t sharedSize = 0, + cudaStream_t stream = 0); +extern "C" cudaError_t cudaLaunchKernel(const void *func, dim3 gridDim, + dim3 blockDim, void **args, + size_t sharedMem, cudaStream_t stream); #endif extern "C" __device__ int printf(const char*, ...); diff --git a/test/CodeGenCUDA/amdgpu-visibility.cu b/test/CodeGenCUDA/amdgpu-visibility.cu new file mode 100644 index 0000000000..9f44eb047f --- /dev/null +++ b/test/CodeGenCUDA/amdgpu-visibility.cu @@ -0,0 +1,21 @@ +// 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 + +#include "Inputs/cuda.h" + +// CHECK-DEFAULT: @c = addrspace(4) externally_initialized global +// CHECK-DEFAULT: @g = addrspace(1) externally_initialized global +// CHECK-PROTECTED: @c = protected addrspace(4) externally_initialized global +// CHECK-PROTECTED: @g = protected addrspace(1) externally_initialized global +// CHECK-HIDDEN: @c = protected addrspace(4) externally_initialized global +// CHECK-HIDDEN: @g = protected addrspace(1) externally_initialized global +__constant__ int c; +__device__ int g; + +// CHECK-DEFAULT: define amdgpu_kernel void @_Z3foov() +// CHECK-PROTECTED: define protected amdgpu_kernel void @_Z3foov() +// CHECK-HIDDEN: define protected amdgpu_kernel void @_Z3foov() +__global__ void foo() { + g = c; +} diff --git a/test/CodeGenCUDA/debug-info-address-class.cu b/test/CodeGenCUDA/debug-info-address-class.cu new file mode 100644 index 0000000000..0ba23af124 --- /dev/null +++ b/test/CodeGenCUDA/debug-info-address-class.cu @@ -0,0 +1,25 @@ +// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx-unknown-unknown -debug-info-kind=limited -dwarf-version=2 -debugger-tuning=gdb | FileCheck %s + +#include "Inputs/cuda.h" + +// CHECK-DAG: ![[FILEVAR0:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar0", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true) +// CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR0]], expr: !DIExpression()) +__device__ int FileVar0; +// CHECK-DAG: ![[FILEVAR1:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar1", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true) +// CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR1]], expr: !DIExpression(DW_OP_constu, 8, DW_OP_swap, DW_OP_xderef)) +__device__ __shared__ int FileVar1; +// CHECK-DAG: ![[FILEVAR2:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar2", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true) +// CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR2]], expr: !DIExpression(DW_OP_constu, 4, DW_OP_swap, DW_OP_xderef)) +__device__ __constant__ int FileVar2; + +__device__ void kernel1( + // CHECK-DAG: ![[ARG:[0-9]+]] = !DILocalVariable(name: "Arg", arg: {{[0-9]+}}, scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}) + // CHECK-DAG: call void @llvm.dbg.declare(metadata i32* {{.*}}, metadata ![[ARG]], metadata !DIExpression()), !dbg !{{[0-9]+}} + int Arg) { + // CHECK-DAG: ![[FUNCVAR0:[0-9]+]] = distinct !DIGlobalVariable(name: "FuncVar0", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: true, isDefinition: true) + // CHECK-DAG: !DIGlobalVariableExpression(var: ![[FUNCVAR0]], expr: !DIExpression(DW_OP_constu, 8, DW_OP_swap, DW_OP_xderef)) + __shared__ int FuncVar0; + // CHECK-DAG: ![[FUNCVAR1:[0-9]+]] = !DILocalVariable(name: "FuncVar1", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}) + // CHECK-DAG: call void @llvm.dbg.declare(metadata i32* {{.*}}, metadata ![[FUNCVAR1]], metadata !DIExpression()), !dbg !{{[0-9]+}} + int FuncVar1; +} diff --git a/test/CodeGenCUDA/debug-info-template.cu b/test/CodeGenCUDA/debug-info-template.cu new file mode 100644 index 0000000000..078e2ecff9 --- /dev/null +++ b/test/CodeGenCUDA/debug-info-template.cu @@ -0,0 +1,10 @@ +// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-unknown-linux -emit-llvm -o - -debug-info-kind=limited -dwarf-version=2 -debugger-tuning=gdb | FileCheck %s + +#include "Inputs/cuda.h" + +__device__ void f(); +template<void(*F)()> __global__ void t() { F(); } +__host__ void g() { t<f><<<1,1>>>(); } + +// Ensure the value of device-function (as value template parameter) is null. +// CHECK: !DITemplateValueParameter(name: "F", type: !{{[0-9]+}}, value: null) diff --git a/test/CodeGenCUDA/device-stub.cu b/test/CodeGenCUDA/device-stub.cu index ea45c391d2..9db5738cde 100644 --- a/test/CodeGenCUDA/device-stub.cu +++ b/test/CodeGenCUDA/device-stub.cu @@ -1,61 +1,97 @@ // RUN: echo "GPU binary would be here" > %t // 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: -target-sdk-version=8.0 -fcuda-include-gpubinary %t -o - \ +// RUN: | FileCheck -allow-deprecated-dag-overlap %s \ +// RUN: --check-prefixes=ALL,LNX,NORDC,CUDA,CUDANORDC,CUDA-OLD // 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: -target-sdk-version=8.0 -fcuda-include-gpubinary %t \ +// RUN: -o - -DNOGLOBALS \ +// RUN: | FileCheck -allow-deprecated-dag-overlap %s \ +// RUN: -check-prefixes=NOGLOBALS,CUDANOGLOBALS // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \ -// RUN: -fgpu-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: -target-sdk-version=8.0 -fgpu-rdc -fcuda-include-gpubinary %t \ +// RUN: -o - \ +// RUN: | FileCheck -allow-deprecated-dag-overlap %s \ +// RUN: --check-prefixes=ALL,LNX,RDC,CUDA,CUDARDC,CUDA-OLD +// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \ +// RUN: -target-sdk-version=8.0 -o - \ +// RUN: | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=NOGPUBIN + +// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \ +// RUN: -target-sdk-version=9.2 -fcuda-include-gpubinary %t -o - \ +// RUN: | FileCheck %s -allow-deprecated-dag-overlap \ +// RUN: --check-prefixes=ALL,LNX,NORDC,CUDA,CUDANORDC,CUDA-NEW +// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \ +// RUN: -target-sdk-version=9.2 -fcuda-include-gpubinary %t -o - -DNOGLOBALS \ +// RUN: | FileCheck -allow-deprecated-dag-overlap %s \ +// RUN: --check-prefixes=NOGLOBALS,CUDANOGLOBALS +// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \ +// RUN: -target-sdk-version=9.2 -fgpu-rdc -fcuda-include-gpubinary %t -o - \ +// RUN: | FileCheck %s -allow-deprecated-dag-overlap \ +// RUN: --check-prefixes=ALL,LNX,RDC,CUDA,CUDARDC,CUDA_NEW +// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \ +// RUN: -target-sdk-version=9.2 -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,HIPEF +// RUN: | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,LNX,NORDC,HIP,HIPEF // 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: -fgpu-rdc -fcuda-include-gpubinary %t -o - -x hip \ -// RUN: | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,NORDC,HIP,HIPEF +// RUN: | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,LNX,NORDC,HIP,HIPEF // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -o - -x hip\ -// RUN: | FileCheck -allow-deprecated-dag-overlap %s -check-prefixes=ALL,NORDC,HIP,HIPNEF +// RUN: | FileCheck -allow-deprecated-dag-overlap %s -check-prefixes=ALL,LNX,NORDC,HIP,HIPNEF + +// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -aux-triple amdgcn -emit-llvm %s \ +// RUN: -fcuda-include-gpubinary %t -o - -x hip\ +// RUN: | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,WIN #include "Inputs/cuda.h" #ifndef NOGLOBALS -// ALL-DAG: @device_var = internal global i32 +// LNX-DAG: @device_var = internal global i32 +// WIN-DAG: @"?device_var@@3HA" = internal global i32 __device__ int device_var; -// ALL-DAG: @constant_var = internal global i32 +// LNX-DAG: @constant_var = internal global i32 +// WIN-DAG: @"?constant_var@@3HA" = internal global i32 __constant__ int constant_var; -// ALL-DAG: @shared_var = internal global i32 +// LNX-DAG: @shared_var = internal global i32 +// WIN-DAG: @"?shared_var@@3HA" = internal global i32 __shared__ int shared_var; // Make sure host globals don't get internalized... -// ALL-DAG: @host_var = global i32 +// LNX-DAG: @host_var = global i32 +// WIN-DAG: @"?host_var@@3HA" = dso_local global i32 int host_var; // ... and that extern vars remain external. -// ALL-DAG: @ext_host_var = external global i32 +// LNX-DAG: @ext_host_var = external global i32 +// WIN-DAG: @"?ext_host_var@@3HA" = external dso_local global i32 extern int ext_host_var; // external device-side variables -> extern references to their shadows. -// ALL-DAG: @ext_device_var = external global i32 +// LNX-DAG: @ext_device_var = external global i32 +// WIN-DAG: @"?ext_device_var@@3HA" = external dso_local global i32 extern __device__ int ext_device_var; -// ALL-DAG: @ext_device_var = external global i32 +// LNX-DAG: @ext_device_var = external global i32 +// WIN-DAG: @"?ext_constant_var@@3HA" = external dso_local global i32 extern __constant__ int ext_constant_var; // external device-side variables with definitions should generate // definitions for the shadows. -// ALL-DAG: @ext_device_var_def = internal global i32 undef, +// LNX-DAG: @ext_device_var_def = internal global i32 undef, +// WIN-DAG: @"?ext_device_var_def@@3HA" = internal global i32 undef extern __device__ int ext_device_var_def; __device__ int ext_device_var_def = 1; -// ALL-DAG: @ext_device_var_def = internal global i32 undef, +// LNX-DAG: @ext_device_var_def = internal global i32 undef, +// WIN-DAG: @"?ext_constant_var_def@@3HA" = internal global i32 undef __constant__ int ext_constant_var_def = 2; + void use_pointers() { int *p; p = &device_var; @@ -68,8 +104,15 @@ void use_pointers() { } // Make sure that all parts of GPU code init/cleanup are there: -// * constant unnamed string with the kernel name -// ALL: private unnamed_addr constant{{.*}}kernelfunc{{.*}}\00" +// * constant unnamed string with the device-side kernel name to be passed to +// __hipRegisterFunction/__cudaRegisterFunction. +// ALL: @0 = private unnamed_addr constant [18 x i8] c"_Z10kernelfunciii\00" +// * constant unnamed string with the device-side kernel name to be passed to +// __hipRegisterVar/__cudaRegisterVar. +// ALL: @1 = private unnamed_addr constant [11 x i8] c"device_var\00" +// ALL: @2 = private unnamed_addr constant [13 x i8] c"constant_var\00" +// ALL: @3 = private unnamed_addr constant [19 x i8] c"ext_device_var_def\00" +// ALL: @4 = private unnamed_addr constant [21 x i8] c"ext_constant_var_def\00" // * constant unnamed string with GPU binary // CUDA: @[[FATBIN:.*]] = private constant{{.*GPU binary would be here.*}}\00", // HIPEF: @[[FATBIN:.*]] = private constant{{.*GPU binary would be here.*}}\00", @@ -78,13 +121,13 @@ void use_pointers() { // CUDARDC-SAME: section "__nv_relfatbin", align 8 // * constant struct that wraps GPU binary // ALL: @__[[PREFIX:cuda|hip]]_fatbin_wrapper = internal constant -// ALL-SAME: { i32, i32, i8*, i8* } +// LNX-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), // HIPEF-SAME: i8* getelementptr inbounds ({{.*}}@[[FATBIN]], i64 0, i64 0), // HIPNEF-SAME: i8* @[[FATBIN]], -// ALL-SAME: i8* null } +// LNX-SAME: i8* null } // CUDA-SAME: section ".nvFatBinSegment" // HIP-SAME: section ".hipFatBinSegment" // * variable to save GPU binary handle after initialization @@ -94,7 +137,7 @@ void use_pointers() { // 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 +// LNX: @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 @@ -102,31 +145,50 @@ void use_pointers() { // Test that we build the correct number of calls to cudaSetupArgument followed // by a call to cudaLaunch. -// ALL: define{{.*}}kernelfunc -// ALL: call{{.*}}[[PREFIX]]SetupArgument -// ALL: call{{.*}}[[PREFIX]]SetupArgument -// ALL: call{{.*}}[[PREFIX]]SetupArgument -// ALL: call{{.*}}[[PREFIX]]Launch +// LNX: define{{.*}}kernelfunc + +// New launch sequence stores arguments into local buffer and passes array of +// pointers to them directly to cudaLaunchKernel +// CUDA-NEW: alloca +// CUDA-NEW: store +// CUDA-NEW: store +// CUDA-NEW: store +// CUDA-NEW: call{{.*}}__cudaPopCallConfiguration +// CUDA-NEW: call{{.*}}cudaLaunchKernel + +// Legacy style launch sequence sets up arguments by passing them to +// [cuda|hip]SetupArgument. +// CUDA-OLD: call{{.*}}[[PREFIX]]SetupArgument +// CUDA-OLD: call{{.*}}[[PREFIX]]SetupArgument +// CUDA-OLD: call{{.*}}[[PREFIX]]SetupArgument +// CUDA-OLD: call{{.*}}[[PREFIX]]Launch + +// HIP: call{{.*}}[[PREFIX]]SetupArgument +// HIP: call{{.*}}[[PREFIX]]SetupArgument +// HIP: call{{.*}}[[PREFIX]]SetupArgument +// HIP: call{{.*}}[[PREFIX]]Launch __global__ void kernelfunc(int i, int j, int k) {} // Test that we've built correct kernel launch sequence. -// ALL: define{{.*}}hostfunc -// ALL: call{{.*}}[[PREFIX]]ConfigureCall -// ALL: call{{.*}}kernelfunc +// LNX: define{{.*}}hostfunc +// CUDA-OLD: call{{.*}}[[PREFIX]]ConfigureCall +// CUDA-NEW: call{{.*}}__cudaPushCallConfiguration +// HIP: call{{.*}}[[PREFIX]]ConfigureCall +// LNX: 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. // 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_def{{.*}}i32 0, i32 4, i32 0, i32 0 -// ALL-DAG: call{{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}ext_constant_var_def{{.*}}i32 0, i32 4, i32 1, i32 0 +// ALL: call{{.*}}[[PREFIX]]RegisterFunction(i8** %0, {{.*}}kernelfunc{{[^,]*}}, {{[^@]*}}@0 +// ALL-DAG: call{{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}device_var{{[^,]*}}, {{[^@]*}}@1, {{.*}}i32 0, i32 4, i32 0, i32 0 +// ALL-DAG: call{{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}constant_var{{[^,]*}}, {{[^@]*}}@2, {{.*}}i32 0, i32 4, i32 1, i32 0 +// ALL-DAG: call{{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}ext_device_var_def{{[^,]*}}, {{[^@]*}}@3, {{.*}}i32 0, i32 4, i32 0, i32 0 +// ALL-DAG: call{{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}ext_constant_var_def{{[^,]*}}, {{[^@]*}}@4, {{.*}}i32 0, i32 4, i32 1, i32 0 // ALL: ret void // Test that we've built a constructor. -// ALL: define internal void @__[[PREFIX]]_module_ctor +// LNX: define internal void @__[[PREFIX]]_module_ctor // In separate mode it calls __[[PREFIX]]RegisterFatBinary(&__[[PREFIX]]_fatbin_wrapper) // HIP only register fat binary once. @@ -165,14 +227,14 @@ void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); } // 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. -// CUDANOGLOBALS: @{{.*}} = private constant{{.*}} -// HIPNOGLOBALS: @{{.*}} = internal constant{{.*}} +// CUDANOGLOBALS-NOT: @{{.*}} = private constant{{.*}} +// HIPNOGLOBALS-NOT: @{{.*}} = internal 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: define internal void @__[[PREFIX:cuda|hip]]_module_ctor +// NOGLOBALS-NOT: 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 +// NOGLOBALS-NOT: define internal void @__[[PREFIX]]_module_dtor +// NOGLOBALS-NOT: call void @__[[PREFIX]]UnregisterFatBinary // There should be no constructors/destructors if we have no GPU binary. // NOGPUBIN-NOT: define internal void @__[[PREFIX]]_register_globals diff --git a/test/CodeGenCUDA/kernel-args-alignment.cu b/test/CodeGenCUDA/kernel-args-alignment.cu index 4bd5eb1bb1..653f3eb23d 100644 --- a/test/CodeGenCUDA/kernel-args-alignment.cu +++ b/test/CodeGenCUDA/kernel-args-alignment.cu @@ -1,8 +1,12 @@ -// RUN: %clang_cc1 --std=c++11 -triple x86_64-unknown-linux-gnu -emit-llvm -o - %s | \ -// RUN: FileCheck -check-prefix HOST -check-prefix CHECK %s +// New CUDA kernel launch sequence does not require explicit specification of +// size/offset for each argument, so only the old way is tested. +// +// RUN: %clang_cc1 --std=c++11 -triple x86_64-unknown-linux-gnu -emit-llvm \ +// RUN: -target-sdk-version=8.0 -o - %s \ +// RUN: | FileCheck -check-prefixes=HOST-OLD,CHECK %s // RUN: %clang_cc1 --std=c++11 -fcuda-is-device -triple nvptx64-nvidia-cuda \ -// RUN: -emit-llvm -o - %s | FileCheck -check-prefix DEVICE -check-prefix CHECK %s +// RUN: -emit-llvm -o - %s | FileCheck -check-prefixes=DEVICE,CHECK %s #include "Inputs/cuda.h" @@ -27,9 +31,9 @@ static_assert(alignof(S) == 8, "Unexpected alignment."); // 1. offset 0, width 1 // 2. offset 8 (because alignof(S) == 8), width 16 // 3. offset 24, width 8 -// HOST: call i32 @cudaSetupArgument({{[^,]*}}, i64 1, i64 0) -// HOST: call i32 @cudaSetupArgument({{[^,]*}}, i64 16, i64 8) -// HOST: call i32 @cudaSetupArgument({{[^,]*}}, i64 8, i64 24) +// HOST-OLD: call i32 @cudaSetupArgument({{[^,]*}}, i64 1, i64 0) +// HOST-OLD: call i32 @cudaSetupArgument({{[^,]*}}, i64 16, i64 8) +// HOST-OLD: call i32 @cudaSetupArgument({{[^,]*}}, i64 8, i64 24) // DEVICE-LABEL: @_Z6kernelc1SPi // DEVICE-SAME: i8{{[^,]*}}, %struct.S* byval align 8{{[^,]*}}, i32* diff --git a/test/CodeGenCUDA/kernel-call.cu b/test/CodeGenCUDA/kernel-call.cu index 43d08dfaf8..ed48a6cc81 100644 --- a/test/CodeGenCUDA/kernel-call.cu +++ b/test/CodeGenCUDA/kernel-call.cu @@ -1,5 +1,9 @@ -// 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 +// RUN: %clang_cc1 -target-sdk-version=8.0 -emit-llvm %s -o - \ +// RUN: | FileCheck %s --check-prefixes=CUDA-OLD,CHECK +// RUN: %clang_cc1 -target-sdk-version=9.2 -emit-llvm %s -o - \ +// RUN: | FileCheck %s --check-prefixes=CUDA-NEW,CHECK +// RUN: %clang_cc1 -x hip -emit-llvm %s -o - \ +// RUN: | FileCheck %s --check-prefixes=HIP,CHECK #include "Inputs/cuda.h" @@ -7,14 +11,17 @@ // CHECK-LABEL: define{{.*}}g1 // HIP: call{{.*}}hipSetupArgument // HIP: call{{.*}}hipLaunchByPtr -// CUDA: call{{.*}}cudaSetupArgument -// CUDA: call{{.*}}cudaLaunch +// CUDA-OLD: call{{.*}}cudaSetupArgument +// CUDA-OLD: call{{.*}}cudaLaunch +// CUDA-NEW: call{{.*}}__cudaPopCallConfiguration +// CUDA-NEW: call{{.*}}cudaLaunchKernel __global__ void g1(int x) {} // CHECK-LABEL: define{{.*}}main int main(void) { // HIP: call{{.*}}hipConfigureCall - // CUDA: call{{.*}}cudaConfigureCall + // CUDA-OLD: call{{.*}}cudaConfigureCall + // CUDA-NEW: call{{.*}}__cudaPushCallConfiguration // CHECK: icmp // CHECK: br // CHECK: call{{.*}}g1 diff --git a/test/CodeGenCUDA/kernel-stub-name.cu b/test/CodeGenCUDA/kernel-stub-name.cu new file mode 100644 index 0000000000..539d7eec1b --- /dev/null +++ b/test/CodeGenCUDA/kernel-stub-name.cu @@ -0,0 +1,20 @@ +// RUN: echo "GPU binary would be here" > %t + +// 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=CHECK + +#include "Inputs/cuda.h" + +template<class T> +__global__ void kernelfunc() {} + +// CHECK-LABEL: define{{.*}}@_Z8hostfuncv() +// CHECK: call void @[[STUB:_Z10kernelfuncIiEvv.stub]]() +void hostfunc(void) { kernelfunc<int><<<1, 1>>>(); } + +// CHECK: define{{.*}}@[[STUB]] +// CHECK: call{{.*}}@hipLaunchByPtr{{.*}}@[[STUB]] + +// CHECK-LABEL: define{{.*}}@__hip_register_globals +// CHECK: call{{.*}}@__hipRegisterFunction{{.*}}@[[STUB]] diff --git a/test/CodeGenCUDA/types.cu b/test/CodeGenCUDA/types.cu new file mode 100644 index 0000000000..43c2b85ea7 --- /dev/null +++ b/test/CodeGenCUDA/types.cu @@ -0,0 +1,10 @@ +// 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 x86_64 -aux-triple amdgcn -emit-llvm %s -o - | FileCheck -check-prefix=HOST %s + +#include "Inputs/cuda.h" + +// HOST: @ld_host = global x86_fp80 0xK00000000000000000000 +long double ld_host; + +// DEV: @ld_device = addrspace(1) externally_initialized global double 0.000000e+00 +__device__ long double ld_device; |