diff options
Diffstat (limited to 'test/CodeGenCUDA/device-stub.cu')
| -rw-r--r-- | test/CodeGenCUDA/device-stub.cu | 177 |
1 files changed, 117 insertions, 60 deletions
diff --git a/test/CodeGenCUDA/device-stub.cu b/test/CodeGenCUDA/device-stub.cu index 3376803c50200..716381b7a8264 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 - -// 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 +// 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) + +// 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 - -// There should be no __cuda_register_globals if we have no +// 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 __[[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 |