2016-06-08 09:56:42 +08:00
// REQUIRES: amdgpu-registered-target
2016-11-13 10:37:05 +08:00
// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s
2016-06-08 09:56:42 +08:00
2016-08-05 23:38:46 +08:00
# pragma OPENCL EXTENSION cl_khr_fp64 : enable
2016-06-08 09:56:42 +08:00
typedef unsigned long ulong ;
2016-08-05 23:38:46 +08:00
typedef unsigned int uint ;
2016-06-08 09:56:42 +08:00
2017-03-10 09:30:46 +08:00
// To get all errors for feature checking we need to put them in one function
// since Clang will stop codegen for the next function if it finds error during
// codegen of the previous function.
void test_target_builtin ( global int* out, int a )
2016-06-08 09:56:42 +08:00
{
2017-03-10 09:30:46 +08:00
__builtin_amdgcn_s_memrealtime ( ) ; // expected-error {{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}}
*out = __builtin_amdgcn_mov_dpp ( a, 0 , 0 , 0 , false ) ; // expected-error {{'__builtin_amdgcn_mov_dpp' needs target feature dpp}}
2016-06-08 09:56:42 +08:00
}
void test_s_sleep ( int x )
{
__builtin_amdgcn_s_sleep ( x ) ; // expected-error {{argument to '__builtin_amdgcn_s_sleep' must be a constant integer}}
}
2017-02-25 12:20:20 +08:00
void test_s_waitcnt ( int x )
{
__builtin_amdgcn_s_waitcnt ( x ) ; // expected-error {{argument to '__builtin_amdgcn_s_waitcnt' must be a constant integer}}
}
2017-02-25 12:20:24 +08:00
void test_s_sendmsg ( int in )
{
__builtin_amdgcn_s_sendmsg ( in, 1 ) ; // expected-error {{argument to '__builtin_amdgcn_s_sendmsg' must be a constant integer}}
}
void test_s_sendmsg_var ( int in1, int in2 )
{
__builtin_amdgcn_s_sendmsg ( in1, in2 ) ; // expected-error {{argument to '__builtin_amdgcn_s_sendmsg' must be a constant integer}}
}
void test_s_sendmsghalt ( int in )
{
__builtin_amdgcn_s_sendmsghalt ( in, 1 ) ; // expected-error {{argument to '__builtin_amdgcn_s_sendmsghalt' must be a constant integer}}
}
void test_s_sendmsghalt_var ( int in1, int in2 )
{
__builtin_amdgcn_s_sendmsghalt ( in1, in2 ) ; // expected-error {{argument to '__builtin_amdgcn_s_sendmsghalt' must be a constant integer}}
}
2016-08-19 20:54:31 +08:00
void test_s_incperflevel ( int x )
{
__builtin_amdgcn_s_incperflevel ( x ) ; // expected-error {{argument to '__builtin_amdgcn_s_incperflevel' must be a constant integer}}
}
void test_s_decperflevel ( int x )
{
__builtin_amdgcn_s_decperflevel ( x ) ; // expected-error {{argument to '__builtin_amdgcn_s_decperflevel' must be a constant integer}}
}
2016-08-05 23:38:46 +08:00
void test_sicmp_i32 ( global ulong* out, int a, int b, uint c )
{
*out = __builtin_amdgcn_sicmp ( a, b, c ) ; // expected-error {{argument to '__builtin_amdgcn_sicmp' must be a constant integer}}
}
void test_uicmp_i32 ( global ulong* out, uint a, uint b, uint c )
{
*out = __builtin_amdgcn_uicmp ( a, b, c ) ; // expected-error {{argument to '__builtin_amdgcn_uicmp' must be a constant integer}}
}
void test_sicmp_i64 ( global ulong* out, long a, long b, uint c )
{
*out = __builtin_amdgcn_sicmpl ( a, b, c ) ; // expected-error {{argument to '__builtin_amdgcn_sicmpl' must be a constant integer}}
}
void test_uicmp_i64 ( global ulong* out, ulong a, ulong b, uint c )
{
*out = __builtin_amdgcn_uicmpl ( a, b, c ) ; // expected-error {{argument to '__builtin_amdgcn_uicmpl' must be a constant integer}}
}
void test_fcmp_f32 ( global ulong* out, float a, float b, uint c )
{
*out = __builtin_amdgcn_fcmpf ( a, b, c ) ; // expected-error {{argument to '__builtin_amdgcn_fcmpf' must be a constant integer}}
}
void test_fcmp_f64 ( global ulong* out, double a, double b, uint c )
{
*out = __builtin_amdgcn_fcmp ( a, b, c ) ; // expected-error {{argument to '__builtin_amdgcn_fcmp' must be a constant integer}}
}
2016-08-19 06:04:54 +08:00
void test_ds_swizzle ( global int* out, int a, int b )
{
*out = __builtin_amdgcn_ds_swizzle ( a, b ) ; // expected-error {{argument to '__builtin_amdgcn_ds_swizzle' must be a constant integer}}
}
2017-01-21 03:24:22 +08:00
void test_s_getreg ( global int* out, int a )
{
*out = __builtin_amdgcn_s_getreg ( a ) ; // expected-error {{argument to '__builtin_amdgcn_s_getreg' must be a constant integer}}
}
2017-03-10 09:30:46 +08:00
void test_mov_dpp2 ( global int* out, int a, int b, int c, int d, bool e )
{
*out = __builtin_amdgcn_mov_dpp ( a, b, 0 , 0 , false ) ; // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
*out = __builtin_amdgcn_mov_dpp ( a, 0 , c, 0 , false ) ; // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
*out = __builtin_amdgcn_mov_dpp ( a, 0 , 0 , d, false ) ; // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
*out = __builtin_amdgcn_mov_dpp ( a, 0 , 0 , 0 , e ) ; // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
}
2018-10-17 10:32:26 +08:00
void test_update_dpp2 ( global int* out, int a, int b, int c, int d, int e, bool f )
{
*out = __builtin_amdgcn_update_dpp ( a, b, 0 , 0 , 0 , false ) ;
*out = __builtin_amdgcn_update_dpp ( a, 0 , c, 0 , 0 , false ) ; // expected-error {{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
*out = __builtin_amdgcn_update_dpp ( a, 0 , 0 , d, 0 , false ) ; // expected-error {{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
*out = __builtin_amdgcn_update_dpp ( a, 0 , 0 , 0 , e, false ) ; // expected-error {{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
*out = __builtin_amdgcn_update_dpp ( a, 0 , 0 , 0 , 0 , f ) ; // expected-error {{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
}
2018-05-22 00:18:07 +08:00
void test_ds_faddf ( local float *out, float src, int a ) {
*out = __builtin_amdgcn_ds_faddf ( out, src, a, 0 , false ) ; // expected-error {{argument to '__builtin_amdgcn_ds_faddf' must be a constant integer}}
*out = __builtin_amdgcn_ds_faddf ( out, src, 0 , a, false ) ; // expected-error {{argument to '__builtin_amdgcn_ds_faddf' must be a constant integer}}
*out = __builtin_amdgcn_ds_faddf ( out, src, 0 , 0 , a ) ; // expected-error {{argument to '__builtin_amdgcn_ds_faddf' must be a constant integer}}
}
void test_ds_fminf ( local float *out, float src, int a ) {
*out = __builtin_amdgcn_ds_fminf ( out, src, a, 0 , false ) ; // expected-error {{argument to '__builtin_amdgcn_ds_fminf' must be a constant integer}}
*out = __builtin_amdgcn_ds_fminf ( out, src, 0 , a, false ) ; // expected-error {{argument to '__builtin_amdgcn_ds_fminf' must be a constant integer}}
*out = __builtin_amdgcn_ds_fminf ( out, src, 0 , 0 , a ) ; // expected-error {{argument to '__builtin_amdgcn_ds_fminf' must be a constant integer}}
}
void test_ds_fmaxf ( local float *out, float src, int a ) {
*out = __builtin_amdgcn_ds_fmaxf ( out, src, a, 0 , false ) ; // expected-error {{argument to '__builtin_amdgcn_ds_fmaxf' must be a constant integer}}
*out = __builtin_amdgcn_ds_fmaxf ( out, src, 0 , a, false ) ; // expected-error {{argument to '__builtin_amdgcn_ds_fmaxf' must be a constant integer}}
*out = __builtin_amdgcn_ds_fmaxf ( out, src, 0 , 0 , a ) ; // expected-error {{argument to '__builtin_amdgcn_ds_fmaxf' must be a constant integer}}
}
2020-04-27 10:56:03 +08:00
void test_fence ( ) {
__builtin_amdgcn_fence ( __ATOMIC_SEQ_CST + 1 , "workgroup" ) ; // expected-warning {{memory order argument to atomic operation is invalid}}
__builtin_amdgcn_fence ( __ATOMIC_ACQUIRE - 1 , "workgroup" ) ; // expected-warning {{memory order argument to atomic operation is invalid}}
__builtin_amdgcn_fence ( 4 ) ; // expected-error {{too few arguments to function call, expected 2}}
__builtin_amdgcn_fence ( 4 , 4 , 4 ) ; // expected-error {{too many arguments to function call, expected 2}}
__builtin_amdgcn_fence ( 3.14 , "" ) ; // expected-warning {{implicit conversion from 'double' to 'unsigned int' changes value from 3.14 to 3}}
__builtin_amdgcn_fence ( __ATOMIC_ACQUIRE, 5 ) ; // expected-warning {{incompatible integer to pointer conversion passing 'int' to parameter of type 'const char *'}}
const char ptr[] = "workgroup" ;
__builtin_amdgcn_fence ( __ATOMIC_ACQUIRE, ptr ) ; // expected-error {{expression is not a string literal}}
}
2020-05-07 04:43:33 +08:00
void test_s_setreg ( int x, int y ) {
__builtin_amdgcn_s_setreg ( x, 0 ) ; // expected-error {{argument to '__builtin_amdgcn_s_setreg' must be a constant integer}}
__builtin_amdgcn_s_setreg ( x, y ) ; // expected-error {{argument to '__builtin_amdgcn_s_setreg' must be a constant integer}}
}
[AMDGPU] Introduce Clang builtins to be mapped to AMDGCN atomic inc/dec intrinsics
Summary:
__builtin_amdgcn_atomic_inc32(int *Ptr, int Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_inc64(int64_t *Ptr, int64_t Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_dec32(int *Ptr, int Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_dec64(int64_t *Ptr, int64_t Val, unsigned MemoryOrdering, const char *SyncScope)
First and second arguments gets transparently passed to the amdgcn atomic
inc/dec intrinsic. Fifth argument of the intrinsic is set as true if the
first argument of the builtin is a volatile pointer. The third argument of
this builtin is one of the memory-ordering specifiers ATOMIC_ACQUIRE,
ATOMIC_RELEASE, ATOMIC_ACQ_REL, or ATOMIC_SEQ_CST following C++11 memory
model semantics. This is mapped to corresponding LLVM atomic memory ordering
for the atomic inc/dec instruction using CLANG atomic C ABI. The fourth
argument is an AMDGPU-specific synchronization scope defined as string.
Reviewers: arsenm, sameerds, JonChesterfield, jdoerfert
Reviewed By: arsenm, sameerds
Subscribers: kzhuravl, jvesely, wdng, nhaehnle, yaxunl, dstuttard, tpr, t-tye, jfb, kerbowa, cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D80804
2020-05-29 22:16:07 +08:00
void test_atomic_inc32 ( ) {
[AMDGPU] Change Clang AMDGCN atomic inc/dec builtins to take unsigned values
builtin_amdgcn_atomic_inc32(uint *Ptr, uint Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_inc64(uint64_t *Ptr, uint64_t Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_dec32(uint *Ptr, uint Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_dec64(uint64_t *Ptr, uint64_t Val, unsigned MemoryOrdering, const char *SyncScope)
As AMDGCN IR instrinsic for atomic inc/dec does unsigned comparison,
these clang builtins should also take unsigned types instead of signed
int types.
Reviewed By: JonChesterfield
Differential Revision: https://reviews.llvm.org/D83121
2020-07-07 14:13:43 +08:00
uint val = 17 ;
[AMDGPU] Introduce Clang builtins to be mapped to AMDGCN atomic inc/dec intrinsics
Summary:
__builtin_amdgcn_atomic_inc32(int *Ptr, int Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_inc64(int64_t *Ptr, int64_t Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_dec32(int *Ptr, int Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_dec64(int64_t *Ptr, int64_t Val, unsigned MemoryOrdering, const char *SyncScope)
First and second arguments gets transparently passed to the amdgcn atomic
inc/dec intrinsic. Fifth argument of the intrinsic is set as true if the
first argument of the builtin is a volatile pointer. The third argument of
this builtin is one of the memory-ordering specifiers ATOMIC_ACQUIRE,
ATOMIC_RELEASE, ATOMIC_ACQ_REL, or ATOMIC_SEQ_CST following C++11 memory
model semantics. This is mapped to corresponding LLVM atomic memory ordering
for the atomic inc/dec instruction using CLANG atomic C ABI. The fourth
argument is an AMDGPU-specific synchronization scope defined as string.
Reviewers: arsenm, sameerds, JonChesterfield, jdoerfert
Reviewed By: arsenm, sameerds
Subscribers: kzhuravl, jvesely, wdng, nhaehnle, yaxunl, dstuttard, tpr, t-tye, jfb, kerbowa, cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D80804
2020-05-29 22:16:07 +08:00
val = __builtin_amdgcn_atomic_inc32 ( &val, val, __ATOMIC_SEQ_CST + 1 , "workgroup" ) ; // expected-warning {{memory order argument to atomic operation is invalid}}
val = __builtin_amdgcn_atomic_inc32 ( &val, val, __ATOMIC_ACQUIRE - 1 , "workgroup" ) ; // expected-warning {{memory order argument to atomic operation is invalid}}
val = __builtin_amdgcn_atomic_inc32 ( 4 ) ; // expected-error {{too few arguments to function call, expected 4}}
val = __builtin_amdgcn_atomic_inc32 ( &val, val, 4 , 4 , 4 , 4 ) ; // expected-error {{too many arguments to function call, expected 4}}
val = __builtin_amdgcn_atomic_inc32 ( &val, val, 3.14 , "" ) ; // expected-warning {{implicit conversion from 'double' to 'unsigned int' changes value from 3.14 to 3}}
val = __builtin_amdgcn_atomic_inc32 ( &val, val, __ATOMIC_ACQUIRE, 5 ) ; // expected-warning {{incompatible integer to pointer conversion passing 'int' to parameter of type 'const char *'}}
const char ptr[] = "workgroup" ;
val = __builtin_amdgcn_atomic_inc32 ( &val, val, __ATOMIC_ACQUIRE, ptr ) ; // expected-error {{expression is not a string literal}}
[AMDGPU] Change Clang AMDGCN atomic inc/dec builtins to take unsigned values
builtin_amdgcn_atomic_inc32(uint *Ptr, uint Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_inc64(uint64_t *Ptr, uint64_t Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_dec32(uint *Ptr, uint Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_dec64(uint64_t *Ptr, uint64_t Val, unsigned MemoryOrdering, const char *SyncScope)
As AMDGCN IR instrinsic for atomic inc/dec does unsigned comparison,
these clang builtins should also take unsigned types instead of signed
int types.
Reviewed By: JonChesterfield
Differential Revision: https://reviews.llvm.org/D83121
2020-07-07 14:13:43 +08:00
int signedVal = 15 ;
signedVal = __builtin_amdgcn_atomic_inc32 ( &signedVal, signedVal, __ATOMIC_ACQUIRE, "" ) ; // expected-warning {{passing '__private int *' to parameter of type 'volatile __private unsigned int *' converts between pointers to integer types with different sign}}
[AMDGPU] Introduce Clang builtins to be mapped to AMDGCN atomic inc/dec intrinsics
Summary:
__builtin_amdgcn_atomic_inc32(int *Ptr, int Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_inc64(int64_t *Ptr, int64_t Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_dec32(int *Ptr, int Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_dec64(int64_t *Ptr, int64_t Val, unsigned MemoryOrdering, const char *SyncScope)
First and second arguments gets transparently passed to the amdgcn atomic
inc/dec intrinsic. Fifth argument of the intrinsic is set as true if the
first argument of the builtin is a volatile pointer. The third argument of
this builtin is one of the memory-ordering specifiers ATOMIC_ACQUIRE,
ATOMIC_RELEASE, ATOMIC_ACQ_REL, or ATOMIC_SEQ_CST following C++11 memory
model semantics. This is mapped to corresponding LLVM atomic memory ordering
for the atomic inc/dec instruction using CLANG atomic C ABI. The fourth
argument is an AMDGPU-specific synchronization scope defined as string.
Reviewers: arsenm, sameerds, JonChesterfield, jdoerfert
Reviewed By: arsenm, sameerds
Subscribers: kzhuravl, jvesely, wdng, nhaehnle, yaxunl, dstuttard, tpr, t-tye, jfb, kerbowa, cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D80804
2020-05-29 22:16:07 +08:00
}
void test_atomic_inc64 ( ) {
[AMDGPU] Change Clang AMDGCN atomic inc/dec builtins to take unsigned values
builtin_amdgcn_atomic_inc32(uint *Ptr, uint Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_inc64(uint64_t *Ptr, uint64_t Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_dec32(uint *Ptr, uint Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_dec64(uint64_t *Ptr, uint64_t Val, unsigned MemoryOrdering, const char *SyncScope)
As AMDGCN IR instrinsic for atomic inc/dec does unsigned comparison,
these clang builtins should also take unsigned types instead of signed
int types.
Reviewed By: JonChesterfield
Differential Revision: https://reviews.llvm.org/D83121
2020-07-07 14:13:43 +08:00
__UINT64_TYPE__ val = 17 ;
[AMDGPU] Introduce Clang builtins to be mapped to AMDGCN atomic inc/dec intrinsics
Summary:
__builtin_amdgcn_atomic_inc32(int *Ptr, int Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_inc64(int64_t *Ptr, int64_t Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_dec32(int *Ptr, int Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_dec64(int64_t *Ptr, int64_t Val, unsigned MemoryOrdering, const char *SyncScope)
First and second arguments gets transparently passed to the amdgcn atomic
inc/dec intrinsic. Fifth argument of the intrinsic is set as true if the
first argument of the builtin is a volatile pointer. The third argument of
this builtin is one of the memory-ordering specifiers ATOMIC_ACQUIRE,
ATOMIC_RELEASE, ATOMIC_ACQ_REL, or ATOMIC_SEQ_CST following C++11 memory
model semantics. This is mapped to corresponding LLVM atomic memory ordering
for the atomic inc/dec instruction using CLANG atomic C ABI. The fourth
argument is an AMDGPU-specific synchronization scope defined as string.
Reviewers: arsenm, sameerds, JonChesterfield, jdoerfert
Reviewed By: arsenm, sameerds
Subscribers: kzhuravl, jvesely, wdng, nhaehnle, yaxunl, dstuttard, tpr, t-tye, jfb, kerbowa, cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D80804
2020-05-29 22:16:07 +08:00
val = __builtin_amdgcn_atomic_inc64 ( &val, val, __ATOMIC_SEQ_CST + 1 , "workgroup" ) ; // expected-warning {{memory order argument to atomic operation is invalid}}
val = __builtin_amdgcn_atomic_inc64 ( &val, val, __ATOMIC_ACQUIRE - 1 , "workgroup" ) ; // expected-warning {{memory order argument to atomic operation is invalid}}
val = __builtin_amdgcn_atomic_inc64 ( 4 ) ; // expected-error {{too few arguments to function call, expected 4}}
val = __builtin_amdgcn_atomic_inc64 ( &val, val, 4 , 4 , 4 , 4 ) ; // expected-error {{too many arguments to function call, expected 4}}
val = __builtin_amdgcn_atomic_inc64 ( &val, val, 3.14 , "" ) ; // expected-warning {{implicit conversion from 'double' to 'unsigned int' changes value from 3.14 to 3}}
val = __builtin_amdgcn_atomic_inc64 ( &val, val, __ATOMIC_ACQUIRE, 5 ) ; // expected-warning {{incompatible integer to pointer conversion passing 'int' to parameter of type 'const char *'}}
const char ptr[] = "workgroup" ;
val = __builtin_amdgcn_atomic_inc64 ( &val, val, __ATOMIC_ACQUIRE, ptr ) ; // expected-error {{expression is not a string literal}}
[AMDGPU] Change Clang AMDGCN atomic inc/dec builtins to take unsigned values
builtin_amdgcn_atomic_inc32(uint *Ptr, uint Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_inc64(uint64_t *Ptr, uint64_t Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_dec32(uint *Ptr, uint Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_dec64(uint64_t *Ptr, uint64_t Val, unsigned MemoryOrdering, const char *SyncScope)
As AMDGCN IR instrinsic for atomic inc/dec does unsigned comparison,
these clang builtins should also take unsigned types instead of signed
int types.
Reviewed By: JonChesterfield
Differential Revision: https://reviews.llvm.org/D83121
2020-07-07 14:13:43 +08:00
__INT64_TYPE__ signedVal = 15 ;
signedVal = __builtin_amdgcn_atomic_inc64 ( &signedVal, signedVal, __ATOMIC_ACQUIRE, "" ) ; // expected-warning {{passing '__private long *' to parameter of type 'volatile __private unsigned long *' converts between pointers to integer types with different sign}}
[AMDGPU] Introduce Clang builtins to be mapped to AMDGCN atomic inc/dec intrinsics
Summary:
__builtin_amdgcn_atomic_inc32(int *Ptr, int Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_inc64(int64_t *Ptr, int64_t Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_dec32(int *Ptr, int Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_dec64(int64_t *Ptr, int64_t Val, unsigned MemoryOrdering, const char *SyncScope)
First and second arguments gets transparently passed to the amdgcn atomic
inc/dec intrinsic. Fifth argument of the intrinsic is set as true if the
first argument of the builtin is a volatile pointer. The third argument of
this builtin is one of the memory-ordering specifiers ATOMIC_ACQUIRE,
ATOMIC_RELEASE, ATOMIC_ACQ_REL, or ATOMIC_SEQ_CST following C++11 memory
model semantics. This is mapped to corresponding LLVM atomic memory ordering
for the atomic inc/dec instruction using CLANG atomic C ABI. The fourth
argument is an AMDGPU-specific synchronization scope defined as string.
Reviewers: arsenm, sameerds, JonChesterfield, jdoerfert
Reviewed By: arsenm, sameerds
Subscribers: kzhuravl, jvesely, wdng, nhaehnle, yaxunl, dstuttard, tpr, t-tye, jfb, kerbowa, cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D80804
2020-05-29 22:16:07 +08:00
}
void test_atomic_dec32 ( ) {
[AMDGPU] Change Clang AMDGCN atomic inc/dec builtins to take unsigned values
builtin_amdgcn_atomic_inc32(uint *Ptr, uint Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_inc64(uint64_t *Ptr, uint64_t Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_dec32(uint *Ptr, uint Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_dec64(uint64_t *Ptr, uint64_t Val, unsigned MemoryOrdering, const char *SyncScope)
As AMDGCN IR instrinsic for atomic inc/dec does unsigned comparison,
these clang builtins should also take unsigned types instead of signed
int types.
Reviewed By: JonChesterfield
Differential Revision: https://reviews.llvm.org/D83121
2020-07-07 14:13:43 +08:00
uint val = 17 ;
[AMDGPU] Introduce Clang builtins to be mapped to AMDGCN atomic inc/dec intrinsics
Summary:
__builtin_amdgcn_atomic_inc32(int *Ptr, int Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_inc64(int64_t *Ptr, int64_t Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_dec32(int *Ptr, int Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_dec64(int64_t *Ptr, int64_t Val, unsigned MemoryOrdering, const char *SyncScope)
First and second arguments gets transparently passed to the amdgcn atomic
inc/dec intrinsic. Fifth argument of the intrinsic is set as true if the
first argument of the builtin is a volatile pointer. The third argument of
this builtin is one of the memory-ordering specifiers ATOMIC_ACQUIRE,
ATOMIC_RELEASE, ATOMIC_ACQ_REL, or ATOMIC_SEQ_CST following C++11 memory
model semantics. This is mapped to corresponding LLVM atomic memory ordering
for the atomic inc/dec instruction using CLANG atomic C ABI. The fourth
argument is an AMDGPU-specific synchronization scope defined as string.
Reviewers: arsenm, sameerds, JonChesterfield, jdoerfert
Reviewed By: arsenm, sameerds
Subscribers: kzhuravl, jvesely, wdng, nhaehnle, yaxunl, dstuttard, tpr, t-tye, jfb, kerbowa, cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D80804
2020-05-29 22:16:07 +08:00
val = __builtin_amdgcn_atomic_dec32 ( &val, val, __ATOMIC_SEQ_CST + 1 , "workgroup" ) ; // expected-warning {{memory order argument to atomic operation is invalid}}
val = __builtin_amdgcn_atomic_dec32 ( &val, val, __ATOMIC_ACQUIRE - 1 , "workgroup" ) ; // expected-warning {{memory order argument to atomic operation is invalid}}
val = __builtin_amdgcn_atomic_dec32 ( 4 ) ; // expected-error {{too few arguments to function call, expected 4}}
val = __builtin_amdgcn_atomic_dec32 ( &val, val, 4 , 4 , 4 , 4 ) ; // expected-error {{too many arguments to function call, expected 4}}
val = __builtin_amdgcn_atomic_dec32 ( &val, val, 3.14 , "" ) ; // expected-warning {{implicit conversion from 'double' to 'unsigned int' changes value from 3.14 to 3}}
val = __builtin_amdgcn_atomic_dec32 ( &val, val, __ATOMIC_ACQUIRE, 5 ) ; // expected-warning {{incompatible integer to pointer conversion passing 'int' to parameter of type 'const char *'}}
const char ptr[] = "workgroup" ;
val = __builtin_amdgcn_atomic_dec32 ( &val, val, __ATOMIC_ACQUIRE, ptr ) ; // expected-error {{expression is not a string literal}}
[AMDGPU] Change Clang AMDGCN atomic inc/dec builtins to take unsigned values
builtin_amdgcn_atomic_inc32(uint *Ptr, uint Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_inc64(uint64_t *Ptr, uint64_t Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_dec32(uint *Ptr, uint Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_dec64(uint64_t *Ptr, uint64_t Val, unsigned MemoryOrdering, const char *SyncScope)
As AMDGCN IR instrinsic for atomic inc/dec does unsigned comparison,
these clang builtins should also take unsigned types instead of signed
int types.
Reviewed By: JonChesterfield
Differential Revision: https://reviews.llvm.org/D83121
2020-07-07 14:13:43 +08:00
int signedVal = 15 ;
signedVal = __builtin_amdgcn_atomic_dec32 ( &signedVal, signedVal, __ATOMIC_ACQUIRE, "" ) ; // expected-warning {{passing '__private int *' to parameter of type 'volatile __private unsigned int *' converts between pointers to integer types with different sign}}
[AMDGPU] Introduce Clang builtins to be mapped to AMDGCN atomic inc/dec intrinsics
Summary:
__builtin_amdgcn_atomic_inc32(int *Ptr, int Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_inc64(int64_t *Ptr, int64_t Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_dec32(int *Ptr, int Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_dec64(int64_t *Ptr, int64_t Val, unsigned MemoryOrdering, const char *SyncScope)
First and second arguments gets transparently passed to the amdgcn atomic
inc/dec intrinsic. Fifth argument of the intrinsic is set as true if the
first argument of the builtin is a volatile pointer. The third argument of
this builtin is one of the memory-ordering specifiers ATOMIC_ACQUIRE,
ATOMIC_RELEASE, ATOMIC_ACQ_REL, or ATOMIC_SEQ_CST following C++11 memory
model semantics. This is mapped to corresponding LLVM atomic memory ordering
for the atomic inc/dec instruction using CLANG atomic C ABI. The fourth
argument is an AMDGPU-specific synchronization scope defined as string.
Reviewers: arsenm, sameerds, JonChesterfield, jdoerfert
Reviewed By: arsenm, sameerds
Subscribers: kzhuravl, jvesely, wdng, nhaehnle, yaxunl, dstuttard, tpr, t-tye, jfb, kerbowa, cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D80804
2020-05-29 22:16:07 +08:00
}
void test_atomic_dec64 ( ) {
[AMDGPU] Change Clang AMDGCN atomic inc/dec builtins to take unsigned values
builtin_amdgcn_atomic_inc32(uint *Ptr, uint Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_inc64(uint64_t *Ptr, uint64_t Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_dec32(uint *Ptr, uint Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_dec64(uint64_t *Ptr, uint64_t Val, unsigned MemoryOrdering, const char *SyncScope)
As AMDGCN IR instrinsic for atomic inc/dec does unsigned comparison,
these clang builtins should also take unsigned types instead of signed
int types.
Reviewed By: JonChesterfield
Differential Revision: https://reviews.llvm.org/D83121
2020-07-07 14:13:43 +08:00
__UINT64_TYPE__ val = 17 ;
[AMDGPU] Introduce Clang builtins to be mapped to AMDGCN atomic inc/dec intrinsics
Summary:
__builtin_amdgcn_atomic_inc32(int *Ptr, int Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_inc64(int64_t *Ptr, int64_t Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_dec32(int *Ptr, int Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_dec64(int64_t *Ptr, int64_t Val, unsigned MemoryOrdering, const char *SyncScope)
First and second arguments gets transparently passed to the amdgcn atomic
inc/dec intrinsic. Fifth argument of the intrinsic is set as true if the
first argument of the builtin is a volatile pointer. The third argument of
this builtin is one of the memory-ordering specifiers ATOMIC_ACQUIRE,
ATOMIC_RELEASE, ATOMIC_ACQ_REL, or ATOMIC_SEQ_CST following C++11 memory
model semantics. This is mapped to corresponding LLVM atomic memory ordering
for the atomic inc/dec instruction using CLANG atomic C ABI. The fourth
argument is an AMDGPU-specific synchronization scope defined as string.
Reviewers: arsenm, sameerds, JonChesterfield, jdoerfert
Reviewed By: arsenm, sameerds
Subscribers: kzhuravl, jvesely, wdng, nhaehnle, yaxunl, dstuttard, tpr, t-tye, jfb, kerbowa, cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D80804
2020-05-29 22:16:07 +08:00
val = __builtin_amdgcn_atomic_dec64 ( &val, val, __ATOMIC_SEQ_CST + 1 , "workgroup" ) ; // expected-warning {{memory order argument to atomic operation is invalid}}
val = __builtin_amdgcn_atomic_dec64 ( &val, val, __ATOMIC_ACQUIRE - 1 , "workgroup" ) ; // expected-warning {{memory order argument to atomic operation is invalid}}
val = __builtin_amdgcn_atomic_dec64 ( 4 ) ; // expected-error {{too few arguments to function call, expected 4}}
val = __builtin_amdgcn_atomic_dec64 ( &val, val, 4 , 4 , 4 , 4 ) ; // expected-error {{too many arguments to function call, expected 4}}
val = __builtin_amdgcn_atomic_dec64 ( &val, val, 3.14 , "" ) ; // expected-warning {{implicit conversion from 'double' to 'unsigned int' changes value from 3.14 to 3}}
val = __builtin_amdgcn_atomic_dec64 ( &val, val, __ATOMIC_ACQUIRE, 5 ) ; // expected-warning {{incompatible integer to pointer conversion passing 'int' to parameter of type 'const char *'}}
const char ptr[] = "workgroup" ;
val = __builtin_amdgcn_atomic_dec64 ( &val, val, __ATOMIC_ACQUIRE, ptr ) ; // expected-error {{expression is not a string literal}}
[AMDGPU] Change Clang AMDGCN atomic inc/dec builtins to take unsigned values
builtin_amdgcn_atomic_inc32(uint *Ptr, uint Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_inc64(uint64_t *Ptr, uint64_t Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_dec32(uint *Ptr, uint Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_dec64(uint64_t *Ptr, uint64_t Val, unsigned MemoryOrdering, const char *SyncScope)
As AMDGCN IR instrinsic for atomic inc/dec does unsigned comparison,
these clang builtins should also take unsigned types instead of signed
int types.
Reviewed By: JonChesterfield
Differential Revision: https://reviews.llvm.org/D83121
2020-07-07 14:13:43 +08:00
__INT64_TYPE__ signedVal = 15 ;
signedVal = __builtin_amdgcn_atomic_dec64 ( &signedVal, signedVal, __ATOMIC_ACQUIRE, "" ) ; // expected-warning {{passing '__private long *' to parameter of type 'volatile __private unsigned long *' converts between pointers to integer types with different sign}}
[AMDGPU] Introduce Clang builtins to be mapped to AMDGCN atomic inc/dec intrinsics
Summary:
__builtin_amdgcn_atomic_inc32(int *Ptr, int Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_inc64(int64_t *Ptr, int64_t Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_dec32(int *Ptr, int Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_dec64(int64_t *Ptr, int64_t Val, unsigned MemoryOrdering, const char *SyncScope)
First and second arguments gets transparently passed to the amdgcn atomic
inc/dec intrinsic. Fifth argument of the intrinsic is set as true if the
first argument of the builtin is a volatile pointer. The third argument of
this builtin is one of the memory-ordering specifiers ATOMIC_ACQUIRE,
ATOMIC_RELEASE, ATOMIC_ACQ_REL, or ATOMIC_SEQ_CST following C++11 memory
model semantics. This is mapped to corresponding LLVM atomic memory ordering
for the atomic inc/dec instruction using CLANG atomic C ABI. The fourth
argument is an AMDGPU-specific synchronization scope defined as string.
Reviewers: arsenm, sameerds, JonChesterfield, jdoerfert
Reviewed By: arsenm, sameerds
Subscribers: kzhuravl, jvesely, wdng, nhaehnle, yaxunl, dstuttard, tpr, t-tye, jfb, kerbowa, cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D80804
2020-05-29 22:16:07 +08:00
}