AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[HIPIFY][perl][fix] Treat ::device_function as a device function

Browse Source 
+ Do not treat somenamespace::device_function_name as a device function
+ Fix generation of warnUnsupportedDeviceFunctions function in hipify-clang
+ Update hipify-perl based on hipify-clang -perl generation
+ Update device test math_functions.cu for hipify-perl

[Restrictions]
- hipify-perl is yet unable to handle function declarations in user namespaces
- hipify-perl is yet unable to handle using directive
This commit is contained in:
Evgeny Mankov
2019-09-16 17:36:55 +03:00
parent 3f66e7b0b9
commit 4f59ec25fe
3 changed files with 27 additions and 11 deletions
Download Patch File Download Diff File Expand all files Collapse all files
tests/hipify-clang/unit_tests/device/math_functions.cu
+15 -3
View File
@@ -1,5 +1,5 @@
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// Test to warn only on device functions umin and umax as unsupported, but not on user defined ones.
// Synthetic test to warn only on device functions umin and umax as unsupported, but not on user defined ones.
// ToDo: change lit testing in order to parse the output.
#define LEN 1024
@@ -8,9 +8,11 @@
#include <algorithm>
namespace my {
// user defined function
unsigned int umin(unsigned int arg1, unsigned int arg2) {
return (arg1 < arg2) ? arg1 : arg2;
}
// user defined function
unsigned int umax(unsigned int arg1, unsigned int arg2) {
return (arg1 > arg2) ? arg1 : arg2;
}
@@ -18,8 +20,16 @@ namespace my {
__global__ void uint_arithm(float* A, float* B, float* C, unsigned int u1, unsigned int u2)
{
unsigned int _umin = umin(u1, u2);
unsigned int _umax = umax(u1, u2);
// device function call (warn if unsupported)
unsigned int _umin = umin ( u1, u2 );
// device function call (warn if unsupported)
unsigned int _umax = umax ( u1, u2 );
// device function call (warn if unsupported)
unsigned int _umin_global = ::umin ( u1, u2 );
// device function call (warn if unsupported)
unsigned int _umax_global = ::umax(u1, u2);
if (_umin != _umin_global) return;
if (_umax != _umax_global) return;
int i = threadIdx.x;
A[i] = i + _umin;
B[i] = i + _umax;
@@ -29,7 +39,9 @@ __global__ void uint_arithm(float* A, float* B, float* C, unsigned int u1, unsig
int main() {
unsigned int u1 = 33;
unsigned int u2 = 34;
// user defined function call
unsigned int _min = my::umin(u1, u2);
// user defined function call
unsigned int _max = my::umax(u1, u2);
float *A, *B, *C;
// CHECK: hipMalloc((void**)&A, SIZE);