clang/test/CodeGenCUDA/memcpy-libcall.cu

   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(
  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_(
  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_(
  29 void __global__ copy_param_to_global(S *global, S param) {
  30   __builtin_memcpy(global, &param, sizeof(S));
  31 // PTX:        ld.param.u32
  32 // PTX:        st.global.u32
  33 }
  34
  35 // PTX-LABEL: .entry _Z19copy_param_to_localPU3AS51SS_(
  36 void __global__ copy_param_to_local(__attribute__((address_space(5))) S *local,
  37                                     S param) {
  38   __builtin_memcpy(local, &param, sizeof(S));
  39 // PTX:        ld.param.u32
  40 // PTX:        st.local.u32
  41 }
  42
  43 // PTX-LABEL: .func _Z21copy_local_to_genericP1SPU3AS5S_(
  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(
  54 void __global__ copy_param_to_shared( S param) {
  55   __builtin_memcpy(&shared, &param, sizeof(S));
  56 // PTX:        ld.param.u32
  57 // PTX:        st.shared.u32
  58 }
  59
  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 }