Differential D7136 Diff 22751 intern/cycles/app/cycles_cubin_cc.cpp

Changeset View

Standalone View

intern/cycles/app/cycles_cubin_cc.cpp

Show All 37 Lines	template<typename T> std::string to_string(const T &n)
std::ostringstream s;		std::ostringstream s;
s << n;		s << n;
return s.str();		return s.str();
}		}
} // namespace std		} // namespace std

class CompilationSettings {		class CompilationSettings {
public:		public:
CompilationSettings() : target_arch(0), bits(64), verbose(false), fast_math(false)		CompilationSettings()
		: target_arch(0), bits(64), verbose(false), fast_math(false), ptx_only(false)
{		{
}		}

string cuda_toolkit_dir;		string cuda_toolkit_dir;
string input_file;		string input_file;
string output_file;		string output_file;
string ptx_file;		string ptx_file;
vector<string> defines;		vector<string> defines;
vector<string> includes;		vector<string> includes;
int target_arch;		int target_arch;
int bits;		int bits;
bool verbose;		bool verbose;
bool fast_math;		bool fast_math;
		bool ptx_only;
};		};

static bool compile_cuda(CompilationSettings &settings)		static bool compile_cuda(CompilationSettings &settings)
{		{
const char *headers[] = {"stdlib.h", "float.h", "math.h", "stdio.h"};		const char *headers[] = {"stdlib.h", "float.h", "math.h", "stdio.h"};
const char *header_content[] = {"\n", "\n", "\n", "\n"};		const char *header_content[] = {"\n", "\n", "\n", "\n"};

printf("Building %s\n", settings.input_file.c_str());		printf("Building %s\n", settings.input_file.c_str());
Show All 10 Lines	static bool compile_cuda(CompilationSettings &settings)
}		}

for (size_t i = 0; i < settings.defines.size(); i++) {		for (size_t i = 0; i < settings.defines.size(); i++) {
options.push_back("-D" + settings.defines[i]);		options.push_back("-D" + settings.defines[i]);
}		}
options.push_back("-D__KERNEL_CUDA_VERSION__=" + std::to_string(cuewNvrtcVersion()));		options.push_back("-D__KERNEL_CUDA_VERSION__=" + std::to_string(cuewNvrtcVersion()));
options.push_back("-arch=compute_" + std::to_string(settings.target_arch));		options.push_back("-arch=compute_" + std::to_string(settings.target_arch));
options.push_back("--device-as-default-execution-space");		options.push_back("--device-as-default-execution-space");
		options.push_back("-DCYCLES_CUBIN_CC");
if (settings.fast_math)		if (settings.fast_math)
options.push_back("--use_fast_math");		options.push_back("--use_fast_math");

nvrtcProgram prog;		nvrtcProgram prog;
nvrtcResult result = nvrtcCreateProgram(&prog,		nvrtcResult result = nvrtcCreateProgram(&prog,
code.c_str(), // buffer		code.c_str(), // buffer
NULL, // name		NULL, // name
sizeof(headers) / sizeof(void *), // numHeaders		sizeof(headers) / sizeof(void *), // numHeaders
Show All 35 Lines	static bool compile_cuda(CompilationSettings &settings)
}		}

vector<char> ptx_code(ptx_size);		vector<char> ptx_code(ptx_size);
result = nvrtcGetPTX(prog, &ptx_code[0]);		result = nvrtcGetPTX(prog, &ptx_code[0]);
if (result != NVRTC_SUCCESS) {		if (result != NVRTC_SUCCESS) {
fprintf(stderr, "Error: nvrtcGetPTX failed (%d)\n\n", (int)result);		fprintf(stderr, "Error: nvrtcGetPTX failed (%d)\n\n", (int)result);
return false;		return false;
}		}
		if (settings.ptx_only) {
		settings.ptx_file = settings.output_file;
		}
		else {
/* Write a file in the temp folder with the ptx code. */		/* Write a file in the temp folder with the ptx code. */
settings.ptx_file = OIIO::Filesystem::temp_directory_path() + "/" +		settings.ptx_file = OIIO::Filesystem::temp_directory_path() + "/" +
OIIO::Filesystem::unique_path();		OIIO::Filesystem::unique_path();
		}
FILE *f = fopen(settings.ptx_file.c_str(), "wb");		FILE *f = fopen(settings.ptx_file.c_str(), "wb");
fwrite(&ptx_code[0], 1, ptx_size, f);		fwrite(&ptx_code[0], 1, ptx_size, f);
fclose(f);		fclose(f);

return true;		return true;
}		}

static bool link_ptxas(CompilationSettings &settings)		static bool link_ptxas(CompilationSettings &settings)
▲ Show 20 Lines • Show All 95 Lines • ▼ Show 20 Lines	ap.options("Usage: cycles_cubin_cc [options]",
&settings.output_file,		&settings.output_file,
"Output cubin filename",		"Output cubin filename",
"-I %L",		"-I %L",
&settings.includes,		&settings.includes,
"Add additional includepath",		"Add additional includepath",
"-D %L",		"-D %L",
&settings.defines,		&settings.defines,
"Add additional defines",		"Add additional defines",
		"-ptx",
		&settings.ptx_only,
		"emit PTX code",
"-v",		"-v",
&settings.verbose,		&settings.verbose,
"Use verbose logging",		"Use verbose logging",
"--use_fast_math",		"--use_fast_math",
&settings.fast_math,		&settings.fast_math,
"Use fast math",		"Use fast math",
"-cuda-toolkit-dir %s",		"-cuda-toolkit-dir %s",
&settings.cuda_toolkit_dir,		&settings.cuda_toolkit_dir,
Show All 38 Lines	if (!init(settings)) {
exit(EXIT_FAILURE);		exit(EXIT_FAILURE);
}		}

if (!compile_cuda(settings)) {		if (!compile_cuda(settings)) {
fprintf(stderr, "Error: compilation error, exiting\n");		fprintf(stderr, "Error: compilation error, exiting\n");
exit(EXIT_FAILURE);		exit(EXIT_FAILURE);
}		}

		if (!settings.ptx_only) {
if (!link_ptxas(settings)) {		if (!link_ptxas(settings)) {
exit(EXIT_FAILURE);		exit(EXIT_FAILURE);
}		}
		}

return 0;		return 0;
}		}