91 lines · plain
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 52// RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -foffload-via-llvm -emit-llvm -o - | FileCheck %s --check-prefix=HST3// RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -fcuda-is-device -foffload-via-llvm -emit-llvm -o - | FileCheck %s --check-prefix=DEV4 5// Check that we generate LLVM/Offload calls, including the KERNEL_LAUNCH_PARAMS argument.6 7#define __OFFLOAD_VIA_LLVM__ 18#include "Inputs/cuda.h"9 10// HST-LABEL: define dso_local ptx_kernel void @_Z18__device_stub__fooisPvS_(11// HST-SAME: i32 noundef [[TMP0:%.*]], i16 noundef signext [[TMP1:%.*]], ptr noundef [[TMP2:%.*]], ptr noundef [[TMP3:%.*]]) #[[ATTR0:[0-9]+]] {12// HST-NEXT: [[ENTRY:.*:]]13// HST-NEXT: [[DOTADDR:%.*]] = alloca i32, align 414// HST-NEXT: [[DOTADDR1:%.*]] = alloca i16, align 215// HST-NEXT: [[DOTADDR2:%.*]] = alloca ptr, align 416// HST-NEXT: [[DOTADDR3:%.*]] = alloca ptr, align 417// HST-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[TMP0]], align 1618// HST-NEXT: [[KERNEL_LAUNCH_PARAMS:%.*]] = alloca [[TMP1]], align 1619// HST-NEXT: [[GRID_DIM:%.*]] = alloca [[STRUCT_DIM3:%.*]], align 820// HST-NEXT: [[BLOCK_DIM:%.*]] = alloca [[STRUCT_DIM3]], align 821// HST-NEXT: [[SHMEM_SIZE:%.*]] = alloca i32, align 422// HST-NEXT: [[STREAM:%.*]] = alloca ptr, align 423// HST-NEXT: store i32 [[TMP0]], ptr [[DOTADDR]], align 424// HST-NEXT: store i16 [[TMP1]], ptr [[DOTADDR1]], align 225// HST-NEXT: store ptr [[TMP2]], ptr [[DOTADDR2]], align 426// HST-NEXT: store ptr [[TMP3]], ptr [[DOTADDR3]], align 427// HST-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[TMP1]], ptr [[KERNEL_LAUNCH_PARAMS]], i32 0, i32 028// HST-NEXT: store i64 16, ptr [[TMP4]], align 1629// HST-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[TMP1]], ptr [[KERNEL_LAUNCH_PARAMS]], i32 0, i32 130// HST-NEXT: store ptr [[KERNEL_ARGS]], ptr [[TMP5]], align 831// HST-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[TMP1]], ptr [[KERNEL_LAUNCH_PARAMS]], i32 0, i32 232// HST-NEXT: store ptr null, ptr [[TMP6]], align 433// HST-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTADDR]], align 434// HST-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_ARGS]], i32 0, i32 035// HST-NEXT: store i32 [[TMP7]], ptr [[TMP8]], align 1636// HST-NEXT: [[TMP9:%.*]] = load i16, ptr [[DOTADDR1]], align 237// HST-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_ARGS]], i32 0, i32 138// HST-NEXT: store i16 [[TMP9]], ptr [[TMP10]], align 439// HST-NEXT: [[TMP11:%.*]] = load ptr, ptr [[DOTADDR2]], align 440// HST-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_ARGS]], i32 0, i32 241// HST-NEXT: store ptr [[TMP11]], ptr [[TMP12]], align 842// HST-NEXT: [[TMP13:%.*]] = load ptr, ptr [[DOTADDR3]], align 443// HST-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_ARGS]], i32 0, i32 344// HST-NEXT: store ptr [[TMP13]], ptr [[TMP14]], align 445// HST-NEXT: [[TMP15:%.*]] = call i32 @__llvmPopCallConfiguration(ptr [[GRID_DIM]], ptr [[BLOCK_DIM]], ptr [[SHMEM_SIZE]], ptr [[STREAM]])46// HST-NEXT: [[TMP16:%.*]] = load i32, ptr [[SHMEM_SIZE]], align 447// HST-NEXT: [[TMP17:%.*]] = load ptr, ptr [[STREAM]], align 448// HST-NEXT: [[CALL:%.*]] = call noundef i32 @llvmLaunchKernel(ptr noundef @_Z18__device_stub__fooisPvS_, ptr noundef byval([[STRUCT_DIM3]]) align 4 [[GRID_DIM]], ptr noundef byval([[STRUCT_DIM3]]) align 4 [[BLOCK_DIM]], ptr noundef [[KERNEL_LAUNCH_PARAMS]], i32 noundef [[TMP16]], ptr noundef [[TMP17]])49// HST-NEXT: br label %[[SETUP_END:.*]]50// HST: [[SETUP_END]]:51// HST-NEXT: ret void52//53// DEV-LABEL: define dso_local ptx_kernel void @_Z3fooisPvS_(54// DEV-SAME: i32 noundef [[TMP0:%.*]], i16 noundef signext [[TMP1:%.*]], ptr noundef [[TMP2:%.*]], ptr noundef [[TMP3:%.*]]) #[[ATTR0:[0-9]+]] {55// DEV-NEXT: [[ENTRY:.*:]]56// DEV-NEXT: [[DOTADDR:%.*]] = alloca i32, align 457// DEV-NEXT: [[DOTADDR1:%.*]] = alloca i16, align 258// DEV-NEXT: [[DOTADDR2:%.*]] = alloca ptr, align 459// DEV-NEXT: [[DOTADDR3:%.*]] = alloca ptr, align 460// DEV-NEXT: store i32 [[TMP0]], ptr [[DOTADDR]], align 461// DEV-NEXT: store i16 [[TMP1]], ptr [[DOTADDR1]], align 262// DEV-NEXT: store ptr [[TMP2]], ptr [[DOTADDR2]], align 463// DEV-NEXT: store ptr [[TMP3]], ptr [[DOTADDR3]], align 464// DEV-NEXT: ret void65//66__global__ void foo(int, short, void *, void *) {}67 68// HST-LABEL: define dso_local void @_Z5test1Pv(69// HST-SAME: ptr noundef [[PTR:%.*]]) #[[ATTR1:[0-9]+]] {70// HST-NEXT: [[ENTRY:.*:]]71// HST-NEXT: [[PTR_ADDR:%.*]] = alloca ptr, align 472// HST-NEXT: [[AGG_TMP:%.*]] = alloca [[STRUCT_DIM3:%.*]], align 473// HST-NEXT: [[AGG_TMP1:%.*]] = alloca [[STRUCT_DIM3]], align 474// HST-NEXT: store ptr [[PTR]], ptr [[PTR_ADDR]], align 475// HST-NEXT: call void @_ZN4dim3C1Ejjj(ptr noundef nonnull align 4 dereferenceable(12) [[AGG_TMP]], i32 noundef 3, i32 noundef 1, i32 noundef 1)76// HST-NEXT: call void @_ZN4dim3C1Ejjj(ptr noundef nonnull align 4 dereferenceable(12) [[AGG_TMP1]], i32 noundef 7, i32 noundef 1, i32 noundef 1)77// HST-NEXT: [[CALL:%.*]] = call i32 @__llvmPushCallConfiguration(ptr noundef byval([[STRUCT_DIM3]]) align 4 [[AGG_TMP]], ptr noundef byval([[STRUCT_DIM3]]) align 4 [[AGG_TMP1]], i32 noundef 0, ptr noundef null)78// HST-NEXT: [[TOBOOL:%.*]] = icmp ne i32 [[CALL]], 079// HST-NEXT: br i1 [[TOBOOL]], label %[[KCALL_END:.*]], label %[[KCALL_CONFIGOK:.*]]80// HST: [[KCALL_CONFIGOK]]:81// HST-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 482// HST-NEXT: [[TMP1:%.*]] = load ptr, ptr [[PTR_ADDR]], align 483// HST-NEXT: call void @_Z18__device_stub__fooisPvS_(i32 noundef 13, i16 noundef signext 1, ptr noundef [[TMP0]], ptr noundef [[TMP1]]) #[[ATTR3:[0-9]+]]84// HST-NEXT: br label %[[KCALL_END]]85// HST: [[KCALL_END]]:86// HST-NEXT: ret void87//88void test1(void *Ptr) {89 foo<<<3, 7>>>(13, 1, Ptr, Ptr);90}91