41 lines · plain
1// New CUDA kernel launch sequence does not require explicit specification of2// size/offset for each argument, so only the old way is tested.3//4// RUN: %clang_cc1 --std=c++11 -triple x86_64-unknown-linux-gnu -emit-llvm \5// RUN: -target-sdk-version=8.0 -o - %s \6// RUN: | FileCheck -check-prefixes=HOST-OLD,CHECK %s7 8// RUN: %clang_cc1 --std=c++11 -fcuda-is-device -triple nvptx64-nvidia-cuda \9// RUN: -emit-llvm -o - %s | FileCheck -check-prefixes=DEVICE,CHECK %s10 11#include "Inputs/cuda.h"12 13struct U {14 short x;15} __attribute__((packed));16 17struct S {18 int *ptr;19 char a;20 U u;21};22 23// Clang should generate a packed LLVM struct for S (denoted by the <>s),24// otherwise this test isn't interesting.25// CHECK: %struct.S = type <{ ptr, i8, %struct.U, [5 x i8] }>26 27static_assert(alignof(S) == 8, "Unexpected alignment.");28 29// HOST-LABEL: @_Z6kernelc1SPi30// Marshalled kernel args should be:31// 1. offset 0, width 132// 2. offset 8 (because alignof(S) == 8), width 1633// 3. offset 24, width 834// HOST-OLD: call i32 @cudaSetupArgument({{[^,]*}}, i64 1, i64 0)35// HOST-OLD: call i32 @cudaSetupArgument({{[^,]*}}, i64 16, i64 8)36// HOST-OLD: call i32 @cudaSetupArgument({{[^,]*}}, i64 8, i64 24)37 38// DEVICE-LABEL: @_Z6kernelc1SPi39// DEVICE-SAME: i8{{[^,]*}}, ptr noundef byval(%struct.S) align 8{{[^,]*}}, ptr40__global__ void kernel(char a, S s, int *b) {}41