65 lines · plain
1// REQUIRES: x86-registered-target2// REQUIRES: nvptx-registered-target3 4// RUN: %clang_cc1 -x cuda -triple nvptx64-nvidia-cuda- -fcuda-is-device \5// RUN: -O3 -S %s -o - | FileCheck -check-prefix=PTX %s6// RUN: %clang_cc1 -x cuda -triple nvptx64-nvidia-cuda- -fcuda-is-device \7// RUN: -Os -S %s -o - | FileCheck -check-prefix=PTX %s8#include "Inputs/cuda.h"9 10// PTX-LABEL: .func _Z12copy_genericPvPKv(11void __device__ copy_generic(void *dest, const void *src) {12 __builtin_memcpy(dest, src, 32);13// PTX: ld.b814// PTX: st.b815}16 17// PTX-LABEL: .entry _Z11copy_globalPvS_(18void __global__ copy_global(void *dest, void * src) {19 __builtin_memcpy(dest, src, 32);20// PTX: ld.global.b821// PTX: st.global.b822}23 24struct S {25 int data[8];26};27 28// PTX-LABEL: .entry _Z20copy_param_to_globalP1SS_(29void __global__ copy_param_to_global(S *global, S param) {30 __builtin_memcpy(global, ¶m, sizeof(S));31// PTX: ld.param.b3232// PTX: st.global.b3233}34 35// PTX-LABEL: .entry _Z19copy_param_to_localPU3AS51SS_(36void __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.b3240// PTX: st.local.b3241}42 43// PTX-LABEL: .func _Z21copy_local_to_genericP1SPU3AS5S_(44void __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.b3248// PTX: st.b3249}50 51__shared__ S shared;52 53// PTX-LABEL: .entry _Z20copy_param_to_shared1S(54void __global__ copy_param_to_shared( S param) {55 __builtin_memcpy(&shared, ¶m, sizeof(S));56// PTX: ld.param.b3257// PTX: st.shared.b3258}59 60void __device__ copy_shared_to_generic(S *generic) {61 __builtin_memcpy(generic, &shared, sizeof(S));62// PTX: ld.shared.b3263// PTX: st.b3264}65