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