@@ -183,19 +183,19 @@ void test_ds_fminf(__attribute__((address_space(3))) float *out, float src) {
183
183
* out = __builtin_amdgcn_ds_fminf (out , src , 0 , 0 , true);
184
184
185
185
// Test all orders.
186
- * out = __builtin_amdgcn_ds_fminf (out , src , 1 , 0 , false);
187
- * out = __builtin_amdgcn_ds_fminf (out , src , 2 , 0 , false);
188
- * out = __builtin_amdgcn_ds_fminf (out , src , 3 , 0 , false);
189
- * out = __builtin_amdgcn_ds_fminf (out , src , 4 , 0 , false);
190
- * out = __builtin_amdgcn_ds_fminf (out , src , 5 , 0 , false);
191
- * out = __builtin_amdgcn_ds_fminf (out , src , 5 , 0 , false); // invalid
186
+ * out = __builtin_amdgcn_ds_fminf (out , src , __ATOMIC_CONSUME , __MEMORY_SCOPE_SYSTEM , false);
187
+ * out = __builtin_amdgcn_ds_fminf (out , src , __ATOMIC_ACQUIRE , __MEMORY_SCOPE_SYSTEM , false);
188
+ * out = __builtin_amdgcn_ds_fminf (out , src , __ATOMIC_RELEASE , __MEMORY_SCOPE_SYSTEM , false);
189
+ * out = __builtin_amdgcn_ds_fminf (out , src , __ATOMIC_ACQ_REL , __MEMORY_SCOPE_SYSTEM , false);
190
+ * out = __builtin_amdgcn_ds_fminf (out , src , __ATOMIC_SEQ_CST , __MEMORY_SCOPE_SYSTEM , false);
191
+ * out = __builtin_amdgcn_ds_fminf (out , src , __ATOMIC_SEQ_CST , __MEMORY_SCOPE_SYSTEM , false); // invalid
192
192
193
193
// Test all syncscopes.
194
- * out = __builtin_amdgcn_ds_fminf (out , src , 0 , 1 , false);
195
- * out = __builtin_amdgcn_ds_fminf (out , src , 0 , 2 , false);
196
- * out = __builtin_amdgcn_ds_fminf (out , src , 0 , 3 , false);
197
- * out = __builtin_amdgcn_ds_fminf (out , src , 0 , 4 , false);
198
- * out = __builtin_amdgcn_ds_fminf (out , src , 0 , 5 , false); // invalid
194
+ * out = __builtin_amdgcn_ds_fminf (out , src , __ATOMIC_RELAXED , __MEMORY_SCOPE_DEVICE , false);
195
+ * out = __builtin_amdgcn_ds_fminf (out , src , __ATOMIC_RELAXED , __MEMORY_SCOPE_WRKGRP , false);
196
+ * out = __builtin_amdgcn_ds_fminf (out , src , __ATOMIC_RELAXED , __MEMORY_SCOPE_WVFRNT , false);
197
+ * out = __builtin_amdgcn_ds_fminf (out , src , __ATOMIC_RELAXED , __MEMORY_SCOPE_SINGLE , false);
198
+ * out = __builtin_amdgcn_ds_fminf (out , src , __ATOMIC_RELAXED , 5 , false); // invalid
199
199
}
200
200
201
201
// CHECK-LABEL: @test_ds_fmax
@@ -224,19 +224,19 @@ void test_ds_fmaxf(__attribute__((address_space(3))) float *out, float src) {
224
224
* out = __builtin_amdgcn_ds_fmaxf (out , src , 0 , 0 , true);
225
225
226
226
// Test all orders.
227
- * out = __builtin_amdgcn_ds_fmaxf (out , src , 1 , 0 , false);
228
- * out = __builtin_amdgcn_ds_fmaxf (out , src , 2 , 0 , false);
229
- * out = __builtin_amdgcn_ds_fmaxf (out , src , 3 , 0 , false);
230
- * out = __builtin_amdgcn_ds_fmaxf (out , src , 4 , 0 , false);
231
- * out = __builtin_amdgcn_ds_fmaxf (out , src , 5 , 0 , false);
232
- * out = __builtin_amdgcn_ds_fmaxf (out , src , 5 , 0 , false); // invalid
227
+ * out = __builtin_amdgcn_ds_fmaxf (out , src , __ATOMIC_CONSUME , __MEMORY_SCOPE_SYSTEM , false);
228
+ * out = __builtin_amdgcn_ds_fmaxf (out , src , __ATOMIC_ACQUIRE , __MEMORY_SCOPE_SYSTEM , false);
229
+ * out = __builtin_amdgcn_ds_fmaxf (out , src , __ATOMIC_RELEASE , __MEMORY_SCOPE_SYSTEM , false);
230
+ * out = __builtin_amdgcn_ds_fmaxf (out , src , __ATOMIC_ACQ_REL , __MEMORY_SCOPE_SYSTEM , false);
231
+ * out = __builtin_amdgcn_ds_fmaxf (out , src , __ATOMIC_SEQ_CST , __MEMORY_SCOPE_SYSTEM , false);
232
+ * out = __builtin_amdgcn_ds_fmaxf (out , src , __ATOMIC_SEQ_CST , __MEMORY_SCOPE_SYSTEM , false); // invalid
233
233
234
234
// Test all syncscopes.
235
- * out = __builtin_amdgcn_ds_fmaxf (out , src , 0 , 1 , false);
236
- * out = __builtin_amdgcn_ds_fmaxf (out , src , 0 , 2 , false);
237
- * out = __builtin_amdgcn_ds_fmaxf (out , src , 0 , 3 , false);
238
- * out = __builtin_amdgcn_ds_fmaxf (out , src , 0 , 4 , false);
239
- * out = __builtin_amdgcn_ds_fmaxf (out , src , 0 , 5 , false); // invalid
235
+ * out = __builtin_amdgcn_ds_fmaxf (out , src , __ATOMIC_RELAXED , __MEMORY_SCOPE_DEVICE , false);
236
+ * out = __builtin_amdgcn_ds_fmaxf (out , src , __ATOMIC_RELAXED , __MEMORY_SCOPE_WRKGRP , false);
237
+ * out = __builtin_amdgcn_ds_fmaxf (out , src , __ATOMIC_RELAXED , __MEMORY_SCOPE_WVFRNT , false);
238
+ * out = __builtin_amdgcn_ds_fmaxf (out , src , __ATOMIC_RELAXED , __MEMORY_SCOPE_SINGLE , false);
239
+ * out = __builtin_amdgcn_ds_fmaxf (out , src , __ATOMIC_RELAXED , 5 , false); // invalid
240
240
}
241
241
242
242
// CHECK-LABEL: @test_s_memtime
0 commit comments