24 lines · plain
1// RUN: %clang_cc1 -fcuda-is-device -verify=nordc %s2// RUN: %clang_cc1 -fcuda-is-device -fgpu-rdc -verify=rdc %s3// RUN: %clang_cc1 -x hip -fcuda-is-device -verify=hip %s4 5// rdc-no-diagnostics6 7#include "Inputs/cuda.h"8 9__global__ void g2(int x) {}10 11// CHECK-LABEL: define{{.*}}g112__global__ void g1(void) {13 // CHECK: [[CONFIG:%.*]] = call{{.*}}_Z22cudaGetParameterBuffermm(i64{{.*}}64, i64{{.*}}4)14 // CHECK-NEXT: [[FLAG:%.*]] = icmp ne ptr [[CONFIG]], null15 // CHECK-NEXT: br i1 [[FLAG]], label %[[THEN:.*]], label %[[ENDIF:.*]]16 // CHECK: [[THEN]]:17 // CHECK-NEXT: [[PPTR:%.*]] = getelementptr{{.*}}i8, ptr [[CONFIG]], i64 018 // CHECK-NEXT: store i32 42, ptr [[PPTR]]19 // CHECK: = call{{.*}} i32 @_Z16cudaLaunchDevicePvS_4dim3S0_jP10cudaStream(ptr{{.*}} @_Z2g2i, ptr{{.*}} [[CONFIG]],20 g2<<<1, 1>>>(42);21 // nordc-error@-1 {{kernel launch from __device__ or __global__ function requires relocatable device code (i.e. requires -fgpu-rdc)}}22 // hip-error@-2 {{device-side kernel call/launch is not supported}}23}24