1 // REQUIRES: x86-registered-target
2 // REQUIRES: nvptx-registered-target
3
4 // RUN: %clang_cc1 -x cuda -triple nvptx64-nvidia-cuda- -fcuda-is-device \
5 // RUN: -O3 -S %s -o - | FileCheck -check-prefix=PTX %s
6 // RUN: %clang_cc1 -x cuda -triple nvptx64-nvidia-cuda- -fcuda-is-device \
7 // RUN: -Os -S %s -o - | FileCheck -check-prefix=PTX %s
8 #include "Inputs/cuda.h"
9
10 // PTX-LABEL: .func _Z12copy_genericPvPKv(
copy_generic(void * dest,const void * src)11 void __device__ copy_generic(void *dest, const void *src) {
12 __builtin_memcpy(dest, src, 32);
13 // PTX: ld.u8
14 // PTX: st.u8
15 }
16
17 // PTX-LABEL: .entry _Z11copy_globalPvS_(
copy_global(void * dest,void * src)18 void __global__ copy_global(void *dest, void * src) {
19 __builtin_memcpy(dest, src, 32);
20 // PTX: ld.global.u8
21 // PTX: st.global.u8
22 }
23
24 struct S {
25 int data[8];
26 };
27
28 // PTX-LABEL: .entry _Z20copy_param_to_globalP1SS_(
copy_param_to_global(S * global,S param)29 void __global__ copy_param_to_global(S *global, S param) {
30 __builtin_memcpy(global, ¶m, sizeof(S));
31 // PTX: ld.param.u32
32 // PTX: st.global.u32
33 }
34
35 // PTX-LABEL: .entry _Z19copy_param_to_localPU3AS51SS_(
copy_param_to_local(S * local,S param)36 void __global__ copy_param_to_local(__attribute__((address_space(5))) S *local,
37 S param) {
38 __builtin_memcpy(local, ¶m, sizeof(S));
39 // PTX: ld.param.u32
40 // PTX: st.local.u32
41 }
42
43 // PTX-LABEL: .func _Z21copy_local_to_genericP1SPU3AS5S_(
copy_local_to_generic(S * generic,S * src)44 void __device__ copy_local_to_generic(S *generic,
45 __attribute__((address_space(5))) S *src) {
46 __builtin_memcpy(generic, src, sizeof(S));
47 // PTX: ld.local.u32
48 // PTX: st.u32
49 }
50
51 __shared__ S shared;
52
53 // PTX-LABEL: .entry _Z20copy_param_to_shared1S(
copy_param_to_shared(S param)54 void __global__ copy_param_to_shared( S param) {
55 __builtin_memcpy(&shared, ¶m, sizeof(S));
56 // PTX: ld.param.u32
57 // PTX: st.shared.u32
58 }
59
copy_shared_to_generic(S * generic)60 void __device__ copy_shared_to_generic(S *generic) {
61 __builtin_memcpy(generic, &shared, sizeof(S));
62 // PTX: ld.shared.u32
63 // PTX: st.u32
64 }
65