| 1 | // New CUDA kernel launch sequence does not require explicit specification of |
|---|
| 2 | // 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 %s |
|---|
| 7 | |
|---|
| 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 %s |
|---|
| 10 | |
|---|
| 11 | #include "Inputs/cuda.h" |
|---|
| 12 | |
|---|
| 13 | struct U { |
|---|
| 14 | short x; |
|---|
| 15 | } __attribute__((packed)); |
|---|
| 16 | |
|---|
| 17 | struct 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 <{ i32*, i8, %struct.U, [5 x i8] }> |
|---|
| 26 | |
|---|
| 27 | static_assert(alignof(S) == 8, "Unexpected alignment."); |
|---|
| 28 | |
|---|
| 29 | // HOST-LABEL: @_Z6kernelc1SPi |
|---|
| 30 | // Marshalled kernel args should be: |
|---|
| 31 | // 1. offset 0, width 1 |
|---|
| 32 | // 2. offset 8 (because alignof(S) == 8), width 16 |
|---|
| 33 | // 3. offset 24, width 8 |
|---|
| 34 | // 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: @_Z6kernelc1SPi |
|---|
| 39 | // DEVICE-SAME: i8{{[^,]*}}, %struct.S* byval align 8{{[^,]*}}, i32* |
|---|
| 40 | __global__ void kernel(char a, S s, int *b) {} |
|---|
| 41 | |
|---|