Skip to content

Commit 7853795

Browse files
committed
Update clang CodeGenCUDA/memcpy-libcall.cu test
Fix a clang test which failed with these new error messages, as it illegally passed in kernel args in the .local address space while using the -cuda triple.
1 parent 8bbb3f8 commit 7853795

File tree

1 file changed

+10
-8
lines changed

1 file changed

+10
-8
lines changed

clang/test/CodeGenCUDA/memcpy-libcall.cu

Lines changed: 10 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -32,18 +32,20 @@ void __global__ copy_param_to_global(S *global, S param) {
3232
// PTX: st.global.b32
3333
}
3434

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));
35+
// PTX-LABEL: .func (.param .b32 func_retval0) _Z19copy_param_to_local1Si(
36+
int __device__ copy_param_to_local(S param, int i) {
37+
S local;
38+
__builtin_memcpy(&local, &param, sizeof(S));
3939
// PTX: ld.param.b32
4040
// PTX: st.local.b32
41+
return local.data[i];
4142
}
4243

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));
44+
// PTX-LABEL: .func _Z21copy_local_to_genericP1Sii(
45+
void __device__ copy_local_to_generic(S *generic, int i, int j) {
46+
S src = {{0, i, 2*i, 3*i, 4*i, 5*i, 6*i, 7*i}};
47+
src.data[j] = src.data[j+1];
48+
__builtin_memcpy(generic, &src, sizeof(S));
4749
// PTX: ld.local.b32
4850
// PTX: st.b32
4951
}

0 commit comments

Comments
 (0)