ceed-cuda-compile.cpp (1f70653f2b6c0ef3909d0fb12b83efaa0f81e309) ceed-cuda-compile.cpp (a4bfdec2fbbaf1320f707b4acef5a9f16daff0b4)
1// Copyright (c) 2017-2022, Lawrence Livermore National Security, LLC and other CEED contributors.
2// All Rights Reserved. See the top-level LICENSE and NOTICE files for details.
3//
4// SPDX-License-Identifier: BSD-2-Clause
5//
6// This file is part of CEED: http://github.com/ceed
7
8#include "ceed-cuda-compile.h"

--- 59 unchanged lines hidden (view full) ---

68 {
69 char *source;
70
71 CeedCallBackend(CeedLoadSourceToBuffer(ceed, jit_defs_path, &source));
72 jit_defs_source = source;
73 }
74 code << jit_defs_source;
75 code << "\n\n";
1// Copyright (c) 2017-2022, Lawrence Livermore National Security, LLC and other CEED contributors.
2// All Rights Reserved. See the top-level LICENSE and NOTICE files for details.
3//
4// SPDX-License-Identifier: BSD-2-Clause
5//
6// This file is part of CEED: http://github.com/ceed
7
8#include "ceed-cuda-compile.h"

--- 59 unchanged lines hidden (view full) ---

68 {
69 char *source;
70
71 CeedCallBackend(CeedLoadSourceToBuffer(ceed, jit_defs_path, &source));
72 jit_defs_source = source;
73 }
74 code << jit_defs_source;
75 code << "\n\n";
76 CeedCallBackend(CeedFree(&jit_defs_path));
77 CeedCallBackend(CeedFree(&jit_defs_source));
78
79 // Non-macro options
80 opts[0] = "-default-device";
81 CeedCallBackend(CeedGetData(ceed, &ceed_data));
82 CeedCallCuda(ceed, cudaGetDeviceProperties(&prop, ceed_data->device_id));
83 std::string arch_arg = "-arch=compute_" + std::to_string(prop.major) + std::to_string(prop.minor);
84 opts[1] = arch_arg.c_str();
85 opts[2] = "-Dint32_t=int";

--- 6 unchanged lines hidden (view full) ---

92
93 // Compile kernel
94 nvrtcResult result = nvrtcCompileProgram(prog, num_opts, opts);
95
96 if (result != NVRTC_SUCCESS) {
97 char *log;
98 size_t log_size;
99
76
77 // Non-macro options
78 opts[0] = "-default-device";
79 CeedCallBackend(CeedGetData(ceed, &ceed_data));
80 CeedCallCuda(ceed, cudaGetDeviceProperties(&prop, ceed_data->device_id));
81 std::string arch_arg = "-arch=compute_" + std::to_string(prop.major) + std::to_string(prop.minor);
82 opts[1] = arch_arg.c_str();
83 opts[2] = "-Dint32_t=int";

--- 6 unchanged lines hidden (view full) ---

90
91 // Compile kernel
92 nvrtcResult result = nvrtcCompileProgram(prog, num_opts, opts);
93
94 if (result != NVRTC_SUCCESS) {
95 char *log;
96 size_t log_size;
97
98 CeedDebug256(ceed, CEED_DEBUG_COLOR_ERROR, "---------- CEED JIT SOURCE FAILED TO COMPILE ----------\n");
99 CeedDebug(ceed, "File: %s\n", jit_defs_path);
100 CeedDebug(ceed, "Source:\n%s\n", jit_defs_source);
101 CeedDebug256(ceed, CEED_DEBUG_COLOR_ERROR, "---------- CEED JIT SOURCE FAILED TO COMPILE ----------\n");
102 CeedCallBackend(CeedFree(&jit_defs_path));
103 CeedCallBackend(CeedFree(&jit_defs_source));
100 CeedCallNvrtc(ceed, nvrtcGetProgramLogSize(prog, &log_size));
101 CeedCallBackend(CeedMalloc(log_size, &log));
102 CeedCallNvrtc(ceed, nvrtcGetProgramLog(prog, log));
103 return CeedError(ceed, CEED_ERROR_BACKEND, "%s\n%s", nvrtcGetErrorString(result), log);
104 }
104 CeedCallNvrtc(ceed, nvrtcGetProgramLogSize(prog, &log_size));
105 CeedCallBackend(CeedMalloc(log_size, &log));
106 CeedCallNvrtc(ceed, nvrtcGetProgramLog(prog, log));
107 return CeedError(ceed, CEED_ERROR_BACKEND, "%s\n%s", nvrtcGetErrorString(result), log);
108 }
109 CeedCallBackend(CeedFree(&jit_defs_path));
110 CeedCallBackend(CeedFree(&jit_defs_source));
105
106 CeedCallNvrtc(ceed, nvrtcGetPTXSize(prog, &ptx_size));
107 CeedCallBackend(CeedMalloc(ptx_size, &ptx));
108 CeedCallNvrtc(ceed, nvrtcGetPTX(prog, ptx));
109 CeedCallNvrtc(ceed, nvrtcDestroyProgram(&prog));
110
111 CeedCallCuda(ceed, cuModuleLoadData(module, ptx));
112 CeedCallBackend(CeedFree(&ptx));

--- 65 unchanged lines hidden --- 		111
112 CeedCallNvrtc(ceed, nvrtcGetPTXSize(prog, &ptx_size));
113 CeedCallBackend(CeedMalloc(ptx_size, &ptx));
114 CeedCallNvrtc(ceed, nvrtcGetPTX(prog, ptx));
115 CeedCallNvrtc(ceed, nvrtcDestroyProgram(&prog));
116
117 CeedCallCuda(ceed, cuModuleLoadData(module, ptx));
118 CeedCallBackend(CeedFree(&ptx));

--- 65 unchanged lines hidden ---