Skip to content

Commit c9c2c07

Browse files
authored
QF headers for typedefs and macros (#1036)
* jit - qf headers for typedefs and macros * jit - smaller list of permitted files * ceed - only include ceed.h in QF source
1 parent e4820e4 commit c9c2c07

File tree

115 files changed

+472
-278
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

115 files changed

+472
-278
lines changed

.gitlab-ci.yml

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -171,7 +171,7 @@ noether-float:
171171
- rm -f .SUCCESS
172172
# libCEED
173173
# Change to single precision
174-
- sed -i 's/ceed-f64/ceed-f32/1' include/ceed/ceed.h
174+
- sed -i 's/ceed-f64/ceed-f32/1' include/ceed/types.h
175175
# Build libCEED
176176
- make configure HIP_DIR=/opt/rocm OPT='-O -march=native -ffp-contract=fast'
177177
- BACKENDS_CPU=$(make info-backends-all | grep -o '/cpu[^ ]*') && BACKENDS_GPU=$(make info-backends | grep -o '/gpu[^ ]*')
@@ -198,7 +198,7 @@ noether-float:
198198
bash <(curl -s https://codecov.io/bash) -f coverage.info -t ${CODECOV_ACCESS_TOKEN} -F tests;
199199
bash <(curl -s https://codecov.io/bash) -f coverage.info -t ${CODECOV_ACCESS_TOKEN} -F examples;
200200
fi
201-
- sed -i 's/ceed-f32/ceed-f64/1' include/ceed/ceed.h
201+
- sed -i 's/ceed-f32/ceed-f64/1' include/ceed/types.h
202202
artifacts:
203203
paths:
204204
- build/*.junit

Makefile

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -224,6 +224,7 @@ opt.c := $(sort $(wildcard backends/opt/*.c))
224224
avx.c := $(sort $(wildcard backends/avx/*.c))
225225
xsmm.c := $(sort $(wildcard backends/xsmm/*.c))
226226
cuda.c := $(sort $(wildcard backends/cuda/*.c))
227+
cuda.cpp := $(sort $(wildcard backends/cuda/*.cpp))
227228
cuda-ref.c := $(sort $(wildcard backends/cuda-ref/*.c))
228229
cuda-ref.cpp := $(sort $(wildcard backends/cuda-ref/*.cpp))
229230
cuda-ref.cu := $(sort $(wildcard backends/cuda-ref/kernels/*.cu))
@@ -391,7 +392,7 @@ ifneq ($(CUDA_LIB_DIR),)
391392
LIBCEED_CONTAINS_CXX = 1
392393
libceed.c += interface/ceed-cuda.c
393394
libceed.c += $(cuda.c) $(cuda-ref.c) $(cuda-shared.c) $(cuda-gen.c)
394-
libceed.cpp += $(cuda-ref.cpp) $(cuda-gen.cpp)
395+
libceed.cpp += $(cuda.cpp) $(cuda-ref.cpp) $(cuda-gen.cpp)
395396
libceed.cu += $(cuda-ref.cu) $(cuda-shared.cu) $(cuda-gen.cu)
396397
BACKENDS_MAKE += $(CUDA_BACKENDS)
397398
endif
@@ -635,6 +636,7 @@ install : $(libceed) $(OBJDIR)/ceed.pc
635636
"$(includedir)/ceed/jit-source/cuda/" "$(includedir)/ceed/jit-source/hip/"\
636637
"$(includedir)/ceed/jit-source/gallery/" "$(libdir)" "$(pkgconfigdir)")
637638
$(INSTALL_DATA) include/ceed/ceed.h "$(DESTDIR)$(includedir)/ceed/"
639+
$(INSTALL_DATA) include/ceed/types.h "$(DESTDIR)$(includedir)/ceed/"
638640
$(INSTALL_DATA) include/ceed/ceed-f32.h "$(DESTDIR)$(includedir)/ceed/"
639641
$(INSTALL_DATA) include/ceed/ceed-f64.h "$(DESTDIR)$(includedir)/ceed/"
640642
$(INSTALL_DATA) include/ceed/fortran.h "$(DESTDIR)$(includedir)/ceed/"

backends/cuda-gen/ceed-cuda-gen-operator-build.cpp

Lines changed: 0 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -808,11 +808,6 @@ extern "C" int CeedCudaGenOperatorBuild(CeedOperator op) {
808808
string oper;
809809
oper = "CeedKernel_Cuda_gen_" + qFunctionName;
810810

811-
code << "\n#define CEED_QFUNCTION(name) inline __device__ int name\n";
812-
code << "#define CEED_QFUNCTION_HELPER inline __device__\n";
813-
code << "#define CeedPragmaSIMD\n";
814-
code << "#define CEED_ERROR_SUCCESS 0\n\n";
815-
816811
// Find dim and Q1d
817812
bool useCollograd = true;
818813
bool allCollograd = true;

backends/cuda-ref/ceed-cuda-ref-qfunction-load.cpp

Lines changed: 0 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -58,12 +58,6 @@ extern "C" int CeedCudaBuildQFunction(CeedQFunction qf) {
5858
ostringstream code;
5959

6060
// Defintions
61-
code << "\n#define CEED_QFUNCTION(name) inline __device__ int name\n";
62-
code << "#define CEED_QFUNCTION_HELPER inline __device__\n";
63-
code << "#define CeedPragmaSIMD\n";
64-
code << "#define CEED_ERROR_SUCCESS 0\n";
65-
code << "#define CEED_Q_VLA 1\n\n";
66-
code << "typedef struct { const CeedScalar* inputs[16]; CeedScalar* outputs[16]; } Fields_Cuda;\n";
6761
code << read_write;
6862
code << qfunction_source;
6963
code << "\n";

backends/cuda/ceed-cuda-common.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -18,7 +18,7 @@
1818

1919
#define CeedChk_Cu(ceed, x) \
2020
do { \
21-
CUresult cuda_result = x; \
21+
CUresult cuda_result = (CUresult)x; \
2222
if (cuda_result != CUDA_SUCCESS) { \
2323
const char *msg; \
2424
cuGetErrorName(cuda_result, &msg); \
Lines changed: 37 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -7,16 +7,19 @@
77

88
#include <ceed/ceed.h>
99
#include <ceed/backend.h>
10+
#include <ceed/jit-tools.h>
1011
#include <cuda.h>
1112
#include <cuda_runtime.h>
1213
#include <nvrtc.h>
14+
#include <sstream>
15+
#include <stdarg.h>
1316
#include <string.h>
1417
#include "ceed-cuda-common.h"
1518
#include "ceed-cuda-compile.h"
1619

1720
#define CeedChk_Nvrtc(ceed, x) \
1821
do { \
19-
nvrtcResult result = x; \
22+
nvrtcResult result = static_cast<nvrtcResult>(x); \
2023
if (result != NVRTC_SUCCESS) \
2124
return CeedError((ceed), CEED_ERROR_BACKEND, nvrtcGetErrorString(result)); \
2225
} while (0)
@@ -25,50 +28,60 @@ do { \
2528
// Compile CUDA kernel
2629
//------------------------------------------------------------------------------
2730
int CeedCompileCuda(Ceed ceed, const char *source, CUmodule *module,
28-
const CeedInt num_opts, ...) {
31+
const CeedInt num_defines, ...) {
2932
int ierr;
3033
cudaFree(0); // Make sure a Context exists for nvrtc
3134
nvrtcProgram prog;
32-
CeedChk_Nvrtc(ceed, nvrtcCreateProgram(&prog, source, NULL, 0, NULL, NULL));
35+
36+
std::ostringstream code;
3337

3438
// Get kernel specific options, such as kernel constants
35-
const int opts_len = 32;
36-
const int opts_extra = 4;
37-
const char *opts[num_opts + opts_extra];
38-
char buf[num_opts][opts_len];
39-
if (num_opts > 0) {
39+
if (num_defines > 0) {
4040
va_list args;
41-
va_start(args, num_opts);
41+
va_start(args, num_defines);
4242
char *name;
4343
int val;
44-
for (int i = 0; i < num_opts; i++) {
44+
for (int i = 0; i < num_defines; i++) {
4545
name = va_arg(args, char *);
4646
val = va_arg(args, int);
47-
snprintf(&buf[i][0], opts_len,"-D%s=%d", name, val);
48-
opts[i] = &buf[i][0];
47+
code << "#define " << name << " " << val << "\n";
4948
}
5049
va_end(args);
5150
}
5251

53-
// Standard backend options
54-
if (CEED_SCALAR_TYPE == CEED_SCALAR_FP32) {
55-
opts[num_opts] = "-DCeedScalar=float";
56-
} else {
57-
opts[num_opts] = "-DCeedScalar=double";
58-
}
59-
opts[num_opts + 1] = "-DCeedInt=int";
60-
opts[num_opts + 2] = "-default-device";
52+
// Standard libCEED definitions for CUDA backends
53+
char *jit_defs_path, *jit_defs_source;
54+
ierr = CeedGetJitAbsolutePath(ceed,
55+
"ceed/jit-source/cuda/cuda-jit.h",
56+
&jit_defs_path); CeedChkBackend(ierr);
57+
ierr = CeedLoadSourceToBuffer(ceed, jit_defs_path, &jit_defs_source);
58+
CeedChkBackend(ierr);
59+
code << jit_defs_source;
60+
code << "\n\n";
61+
ierr = CeedFree(&jit_defs_path); CeedChkBackend(ierr);
62+
ierr = CeedFree(&jit_defs_source); CeedChkBackend(ierr);
63+
64+
// Non-macro options
65+
const int num_opts = 3;
66+
const char *opts[num_opts];
67+
opts[0] = "-default-device";
6168
struct cudaDeviceProp prop;
6269
Ceed_Cuda *ceed_data;
6370
ierr = CeedGetData(ceed, &ceed_data); CeedChkBackend(ierr);
6471
ierr = cudaGetDeviceProperties(&prop, ceed_data->device_id);
6572
CeedChk_Cu(ceed, ierr);
66-
char buff[opts_len];
67-
snprintf(buff, opts_len,"-arch=compute_%d%d", prop.major, prop.minor);
68-
opts[num_opts + 3] = buff;
73+
std::string arch_arg = "-arch=compute_" + std::to_string(prop.major) + std::to_string(prop.minor);
74+
opts[1] = arch_arg.c_str();
75+
opts[2] = "-Dint32_t=int";
76+
77+
// Add string source argument provided in call
78+
code << source;
79+
80+
// Create Program
81+
CeedChk_Nvrtc(ceed, nvrtcCreateProgram(&prog, code.str().c_str(), NULL, 0, NULL, NULL));
6982

7083
// Compile kernel
71-
nvrtcResult result = nvrtcCompileProgram(prog, num_opts + opts_extra, opts);
84+
nvrtcResult result = nvrtcCompileProgram(prog, num_opts, opts);
7285
if (result != NVRTC_SUCCESS) {
7386
size_t log_size;
7487
CeedChk_Nvrtc(ceed, nvrtcGetProgramLogSize(prog, &log_size));

backends/cuda/ceed-cuda-compile.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -17,7 +17,7 @@ static inline CeedInt CeedDivUpInt(CeedInt numerator, CeedInt denominator) {
1717
}
1818

1919
CEED_INTERN int CeedCompileCuda(Ceed ceed, const char *source, CUmodule *module,
20-
const CeedInt num_opts, ...);
20+
const CeedInt num_defines, ...);
2121

2222
CEED_INTERN int CeedGetKernelCuda(Ceed ceed, CUmodule module, const char *name,
2323
CUfunction *kernel);

backends/hip-gen/ceed-hip-gen-operator-build.cpp

Lines changed: 0 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -807,11 +807,6 @@ extern "C" int CeedHipGenOperatorBuild(CeedOperator op) {
807807
string oper;
808808
oper = "CeedKernel_Hip_gen_" + qFunctionName;
809809

810-
code << "\n#define CEED_QFUNCTION(name) inline __device__ int name\n";
811-
code << "#define CEED_QFUNCTION_HELPER __device__ __forceinline__\n";
812-
code << "#define CeedPragmaSIMD\n";
813-
code << "#define CEED_ERROR_SUCCESS 0\n\n";
814-
815810
// Find dim and Q1d
816811
bool useCollograd = true;
817812
bool allCollograd = true;

backends/hip-ref/ceed-hip-ref-qfunction-load.cpp

Lines changed: 0 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -60,12 +60,6 @@ extern "C" int CeedHipBuildQFunction(CeedQFunction qf) {
6060
ostringstream code;
6161

6262
// Defintions
63-
code << "\n#define CEED_QFUNCTION(name) inline __device__ int name\n";
64-
code << "#define CEED_QFUNCTION_HELPER __device__ __forceinline__\n";
65-
code << "#define CeedPragmaSIMD\n";
66-
code << "#define CEED_ERROR_SUCCESS 0\n";
67-
code << "#define CEED_Q_VLA 1\n\n";
68-
code << "typedef struct { const CeedScalar* inputs[16]; CeedScalar* outputs[16]; } Fields_Hip;\n";
6963
code << read_write;
7064
code << qfunction_source;
7165
code << "\n";

backends/hip/ceed-hip-compile.cpp

Lines changed: 23 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -7,6 +7,7 @@
77

88
#include <ceed/ceed.h>
99
#include <ceed/backend.h>
10+
#include <ceed/jit-tools.h>
1011
#include <sstream>
1112
#include <stdarg.h>
1213
#include <string.h>
@@ -25,12 +26,13 @@ do { \
2526
// Compile HIP kernel
2627
//------------------------------------------------------------------------------
2728
int CeedCompileHip(Ceed ceed, const char *source, hipModule_t *module,
28-
const CeedInt num_opts, ...) {
29+
const CeedInt num_defines, ...) {
2930
int ierr;
30-
hipFree(0); // Make sure a Context exists for hiprtc
31+
hipFree(0); // Make sure a Context exists for hiprtc
3132
hiprtcProgram prog;
3233

3334
std::ostringstream code;
35+
3436
// Add hip runtime include statement for generation if runtime < 40400000
3537
// (implies ROCm < 4.5)
3638
int runtime_version;
@@ -46,34 +48,35 @@ int CeedCompileHip(Ceed ceed, const char *source, hipModule_t *module,
4648
code << "#define HIP_DYNAMIC_SHARED(type, var) extern __shared__ type var[];\n";
4749
}
4850

49-
// Macro definitions
50-
// Get kernel specific options, such as kernel constants
51-
const int opts_size = 3;
52-
const char *opts[opts_size];
53-
if (num_opts > 0) {
51+
// Kernel specific options, such as kernel constants
52+
if (num_defines > 0) {
5453
va_list args;
55-
va_start(args, num_opts);
54+
va_start(args, num_defines);
5655
char *name;
5756
int val;
58-
for (int i = 0; i < num_opts; i++) {
57+
for (int i = 0; i < num_defines; i++) {
5958
name = va_arg(args, char *);
6059
val = va_arg(args, int);
6160
code << "#define " << name << " " << val << "\n";
6261
}
6362
va_end(args);
6463
}
6564

66-
// Standard backend options
67-
if (CEED_SCALAR_TYPE == CEED_SCALAR_FP32) {
68-
code << "#define CeedScalar float\n";
69-
}
70-
else {
71-
code << "#define CeedScalar double\n";
72-
}
73-
code << "#define CeedInt int\n";
74-
code << "#define CEED_ERROR_SUCCESS 0\n\n";
65+
// Standard libCEED definitions for HIP backends
66+
char *jit_defs_path, *jit_defs_source;
67+
ierr = CeedGetJitAbsolutePath(ceed,
68+
"ceed/jit-source/hip/hip-jit.h",
69+
&jit_defs_path); CeedChkBackend(ierr);
70+
ierr = CeedLoadSourceToBuffer(ceed, jit_defs_path, &jit_defs_source);
71+
CeedChkBackend(ierr);
72+
code << jit_defs_source;
73+
code << "\n\n";
74+
ierr = CeedFree(&jit_defs_path); CeedChkBackend(ierr);
75+
ierr = CeedFree(&jit_defs_source); CeedChkBackend(ierr);
7576

76-
// Non-macro options
77+
// Non-macro options
78+
const int num_opts = 3;
79+
const char *opts[num_opts];
7780
opts[0] = "-default-device";
7881
struct hipDeviceProp_t prop;
7982
Ceed_Hip *ceed_data;
@@ -90,7 +93,7 @@ int CeedCompileHip(Ceed ceed, const char *source, hipModule_t *module,
9093
CeedChk_hiprtc(ceed, hiprtcCreateProgram(&prog, code.str().c_str(), NULL, 0, NULL, NULL));
9194

9295
// Compile kernel
93-
hiprtcResult result = hiprtcCompileProgram(prog, opts_size, opts);
96+
hiprtcResult result = hiprtcCompileProgram(prog, num_opts, opts);
9497
if (result != HIPRTC_SUCCESS) {
9598
size_t log_size;
9699
CeedChk_hiprtc(ceed, hiprtcGetProgramLogSize(prog, &log_size));

0 commit comments

Comments
 (0)