Skip to content

Commit 70fac22

Browse files
pavankypavanky
committed
Removing math_ptx submodule as a dependency
1 parent c3e525f commit 70fac22

21 files changed

Lines changed: 628 additions & 96 deletions

‎.gitmodules‎

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,3 @@
1-
[submodule "src/backend/cuda/ptx"]
2-
path = src/backend/cuda/ptx
3-
url = https://github.com/arrayfire/math_ptx
41
[submodule "test/data"]
52
path = test/data
63
url = https://github.com/arrayfire/arrayfire_data

‎src/backend/cuda/Array.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -83,7 +83,7 @@ namespace cuda
8383
if (!node) {
8484
bool is_linear = isOwner() || (this->ndims() == 1);
8585
BufferNode<T> *buf_node = new BufferNode<T>(irname<T>(),
86-
shortname<T>(true), data,
86+
afShortName<T>(), data,
8787
*this, offset, is_linear);
8888
const_cast<Array<T> *>(this)->node = Node_ptr(reinterpret_cast<Node *>(buf_node));
8989
}

‎src/backend/cuda/CMakeLists.txt‎

Lines changed: 22 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -46,11 +46,14 @@ FILE(GLOB cuda_sources
4646
"kernel/*.cu")
4747

4848
FILE(GLOB jit_sources
49-
"jit/*.hpp")
49+
"JIT/*.hpp")
5050

5151
FILE(GLOB kernel_headers
5252
"kernel/*.hpp")
5353

54+
FILE(GLOB ptx_sources
55+
"JIT/*.cu")
56+
5457
source_group(backend\\cuda\\Headers FILES ${cuda_headers})
5558
source_group(backend\\cuda\\Sources FILES ${cuda_sources})
5659
source_group(backend\\cuda\\JIT FILES ${jit_sources})
@@ -96,22 +99,25 @@ ELSE()
9699
SET(PTX_COMPUTE ${CUDA_COMPUTE_CAPABILITY})
97100
ENDIF()
98101

99-
# Check if PTX data exists
100-
IF (EXISTS "${CMAKE_SOURCE_DIR}/src/backend/cuda/ptx/PTX64"
101-
AND IS_DIRECTORY "${CMAKE_SOURCE_DIR}/src/backend/cuda/ptx/PTX64")
102-
# PTX data is available
103-
# Do Nothing
104-
ELSE (EXISTS "${CMAKE_SOURCE_DIR}/src/backend/cuda/ptx/PTX64"
105-
AND IS_DIRECTORY "${CMAKE_SOURCE_DIR}/src/backend/cuda/ptx/PTX64")
106-
MESSAGE(WARNING "PTX data is not available. CUDA Backend will give compiler errors.")
107-
MESSAGE("Did you miss the --recursive option when cloning?")
108-
MESSAGE("Run the following commands to correct this:")
109-
MESSAGE("git submodule init")
110-
MESSAGE("git submodule update")
111-
MESSAGE("git submodule foreach git pull origin master")
112-
ENDIF()
102+
CUDA_COMPILE_PTX(ptx_files ${ptx_sources})
103+
104+
set(cuda_ptx "")
105+
foreach(ptx_src_file ${ptx_sources})
106+
107+
get_filename_component(_name "${ptx_src_file}" NAME_WE)
108+
109+
set(_gen_file_name
110+
"${CMAKE_BINARY_DIR}/src/backend/cuda/cuda_compile_ptx_generated_${_name}.cu.ptx")
111+
set(_out_file_name
112+
"${CMAKE_BINARY_DIR}/src/backend/cuda/${_name}.ptx")
113+
114+
ADD_CUSTOM_COMMAND(
115+
OUTPUT "${_out_file_name}"
116+
DEPENDS "${ptx_files}"
117+
COMMAND ${CMAKE_COMMAND} -E rename "${_gen_file_name}" "${_out_file_name}")
113118

114-
FILE(GLOB cuda_ptx "ptx/PTX64/sm_${PTX_COMPUTE}/*.ptx")
119+
list(APPEND cuda_ptx "${_out_file_name}")
120+
endforeach()
115121

116122
SET( ptx_headers
117123
"ptx_headers")

‎src/backend/cuda/JIT/BinaryNode.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -26,7 +26,7 @@ namespace JIT
2626

2727
public:
2828
BinaryNode(const char *out_type_str,
29-
const char *op_str,
29+
const std::string op_str,
3030
Node_ptr lhs, Node_ptr rhs, int op)
3131
: Node(out_type_str),
3232
m_op_str(op_str),

‎src/backend/cuda/JIT/ScalarNode.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -30,7 +30,7 @@ namespace JIT
3030
ScalarNode(T val)
3131
: Node(irname<T>()),
3232
m_val(val),
33-
m_name_str(shortname<T>(false)),
33+
m_name_str(afShortName<T>(false)),
3434
m_gen_name(false),
3535
m_set_arg(false)
3636
{

‎src/backend/cuda/JIT/UnaryNode.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -26,7 +26,7 @@ namespace JIT
2626

2727
public:
2828
UnaryNode(const char *out_type_str,
29-
const char *op_str,
29+
const std::string op_str,
3030
Node_ptr child, int op)
3131
: Node(out_type_str),
3232
m_op_str(op_str),

‎src/backend/cuda/JIT/arith.cu‎

Lines changed: 33 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,33 @@
1+
#include "types.h"
2+
3+
#define ARITH_BASIC(fn, op, T) \
4+
__device__ T ___##fn(T a, T b) \
5+
{ \
6+
return a op b; \
7+
} \
8+
9+
10+
#define ARITH(fn, op) \
11+
ARITH_BASIC(fn, op, float) \
12+
ARITH_BASIC(fn, op, double) \
13+
ARITH_BASIC(fn, op, int) \
14+
ARITH_BASIC(fn, op, uint) \
15+
ARITH_BASIC(fn, op, char) \
16+
ARITH_BASIC(fn, op, uchar) \
17+
ARITH_BASIC(fn, op, intl) \
18+
ARITH_BASIC(fn, op, uintl) \
19+
\
20+
__device__ cfloat ___##fn(cfloat a, cfloat b) \
21+
{ \
22+
return cuC##fn##f(a, b); \
23+
} \
24+
\
25+
__device__ cdouble ___##fn(cdouble a, cdouble b) \
26+
{ \
27+
return cuC##fn(a, b); \
28+
} \
29+
30+
ARITH(add, +)
31+
ARITH(sub, -)
32+
ARITH(mul, *)
33+
ARITH(div, /)

‎src/backend/cuda/JIT/cast.cu‎

Lines changed: 75 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,75 @@
1+
#include "types.h"
2+
3+
#define CAST_BASIC(FN, To, Ti) __device__ To FN(Ti in) { return (To) in; }
4+
5+
#define CAST_BASIC_BOOL(FN, To, Ti) __device__ To FN(Ti in) { return (To)(in != 0); }
6+
7+
#define CAST(T, X) \
8+
CAST_BASIC(___mk##X, T, float) \
9+
CAST_BASIC(___mk##X, T, double) \
10+
CAST_BASIC(___mk##X, T, int) \
11+
CAST_BASIC(___mk##X, T, uint) \
12+
CAST_BASIC(___mk##X, T, char) \
13+
CAST_BASIC(___mk##X, T, uchar) \
14+
CAST_BASIC(___mk##X, T, intl) \
15+
CAST_BASIC(___mk##X, T, uintl) \
16+
17+
CAST(float, S)
18+
CAST(double, D)
19+
CAST(int, I)
20+
CAST(intl, X)
21+
CAST(uint, U)
22+
CAST(uchar, V)
23+
CAST(uintl, Y)
24+
25+
CAST_BASIC_BOOL(___mkJ, char, float)
26+
CAST_BASIC_BOOL(___mkJ, char, double)
27+
CAST_BASIC_BOOL(___mkJ, char, int)
28+
CAST_BASIC_BOOL(___mkJ, char, uint)
29+
CAST_BASIC_BOOL(___mkJ, char, char)
30+
CAST_BASIC_BOOL(___mkJ, char, uchar)
31+
CAST_BASIC_BOOL(___mkJ, char, intl)
32+
CAST_BASIC_BOOL(___mkJ, char, uintl)
33+
34+
#define CPLX_BASIC(FN, To, Tr, Ti) \
35+
__device__ To FN(Ti in) \
36+
{ \
37+
To out = {(Tr)in, 0}; \
38+
return out; \
39+
} \
40+
41+
#define CPLX_CAST(T, Tr, X) \
42+
CPLX_BASIC(___mk##X, T, Tr, float) \
43+
CPLX_BASIC(___mk##X, T, Tr, double) \
44+
CPLX_BASIC(___mk##X, T, Tr, int) \
45+
CPLX_BASIC(___mk##X, T, Tr, uint) \
46+
CPLX_BASIC(___mk##X, T, Tr, char) \
47+
CPLX_BASIC(___mk##X, T, Tr, uchar) \
48+
49+
CPLX_CAST(cfloat, float, C)
50+
CPLX_CAST(cdouble, double, Z)
51+
52+
__device__ cfloat ___mkC(cdouble C)
53+
{
54+
cfloat res = {C.x, C.y};
55+
return res;
56+
}
57+
58+
__device__ cdouble ___mkZ(cfloat C)
59+
{
60+
cdouble res = {C.x, C.y};
61+
return res;
62+
}
63+
64+
__device__ float ___real(cfloat in) { return in.x; }
65+
__device__ double ___real(cdouble in) { return in.x; }
66+
67+
68+
__device__ float ___imag(cfloat in) { return in.y; }
69+
__device__ double ___imag(cdouble in) { return in.y; }
70+
71+
__device__ cfloat ___cplx(float l, float r) { cfloat out = {l, r}; return out; }
72+
__device__ cdouble ___cplx(double l, double r) { cdouble out = {l, r}; return out; }
73+
74+
__device__ cfloat ___conj(cfloat in) { return cuConjf(in); }
75+
__device__ cdouble ___conj(cdouble in) { return cuConj (in); }

‎src/backend/cuda/JIT/exp.cu‎

Lines changed: 71 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,71 @@
1+
#include "types.h"
2+
3+
#define MATH_BASIC(fn, T) \
4+
__device__ T ___##fn(T a) \
5+
{ \
6+
return fn##f((float)a); \
7+
} \
8+
9+
10+
#define MATH(fn) \
11+
MATH_BASIC(fn, float) \
12+
MATH_BASIC(fn, int) \
13+
MATH_BASIC(fn, uint) \
14+
MATH_BASIC(fn, char) \
15+
MATH_BASIC(fn, uchar) \
16+
__device__ double ___##fn(double a) \
17+
{ \
18+
return fn(a); \
19+
} \
20+
21+
22+
MATH(exp)
23+
MATH(expm1)
24+
MATH(erf)
25+
MATH(erfc)
26+
27+
MATH(log)
28+
MATH(log10)
29+
MATH(log1p)
30+
31+
MATH(sqrt)
32+
MATH(cbrt)
33+
34+
#define MATH2_BASIC(fn, T) \
35+
__device__ T ___##fn(T a, T b) \
36+
{ \
37+
return fn##f((float)a, (float)b); \
38+
} \
39+
40+
#define MATH2(fn) \
41+
MATH2_BASIC(fn, float) \
42+
MATH2_BASIC(fn, int) \
43+
MATH2_BASIC(fn, uint) \
44+
MATH2_BASIC(fn, char) \
45+
MATH2_BASIC(fn, uchar) \
46+
__device__ double ___##fn(double a, double b) \
47+
{ \
48+
return fn(a, b); \
49+
} \
50+
51+
MATH2(pow)
52+
53+
__device__ cfloat ___pow(cfloat a, float b)
54+
{
55+
float R = cuCabsf(a);
56+
float Theta = atan2(a.y, a.x);
57+
float R_b = powf(R, b);
58+
float Theta_b = Theta * b;
59+
cfloat res = {R_b * cosf(Theta_b), R_b * sinf(Theta_b)};
60+
return res;
61+
}
62+
63+
__device__ cdouble ___pow(cdouble a, float b)
64+
{
65+
float R = cuCabs(a);
66+
float Theta = atan2(a.y, a.x);
67+
float R_b = pow(R, b);
68+
float Theta_b = Theta * b;
69+
cdouble res = {R_b * cos(Theta_b), R_b * sin(Theta_b)};
70+
return res;
71+
}

‎src/backend/cuda/JIT/hyper.cu‎

Lines changed: 28 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,28 @@
1+
#include "types.h"
2+
3+
#define MATH_BASIC(fn, T) \
4+
__device__ T ___##fn(T a) \
5+
{ \
6+
return fn##f((float)a); \
7+
} \
8+
9+
10+
#define MATH(fn) \
11+
MATH_BASIC(fn, float) \
12+
MATH_BASIC(fn, int) \
13+
MATH_BASIC(fn, uint) \
14+
MATH_BASIC(fn, char) \
15+
MATH_BASIC(fn, uchar) \
16+
__device__ double ___##fn(double a) \
17+
{ \
18+
return fn(a); \
19+
} \
20+
21+
22+
MATH(sinh)
23+
MATH(cosh)
24+
MATH(tanh)
25+
26+
MATH(asinh)
27+
MATH(acosh)
28+
MATH(atanh)

0 commit comments

Comments
 (0)