class Ikra::Translator::CommandTranslator::KernelBuilder
Builds a CUDA kernel. This class is responsible for generating the kernel function itself (not the block functions/methods though).
For example: __global__ void kernel(env_t *env, int *result, int previous_1, …) { … }
Attributes
Additional Parameters for certain commands that are attached to the kernel
A string returning the result of this kernel for one thread
An array of all blocks that should be translated
IDs of commands that whose results are kept on the GPU
A string containing the statements that execute the body of the kernel
An array of all methods that should be translated
Additional parameters that this kernel should accept (to access the result of previous kernels)
The result type of this kernel
Public Class Methods
# File lib/translator/kernel_builder.rb, line 43 def initialize @methods = [] @blocks = [] @previous_kernel_input = [] @block_invocation = nil @num_threads = nil @additional_parameters = [] @kernel_name = "kernel_" + CommandTranslator.next_unique_id.to_s @cached_results = {} @execution = "" end
Public Instance Methods
Add additional parameters to the kernel function that might be needed for some computations
# File lib/translator/kernel_builder.rb, line 72 def add_additional_parameters(parameter) @additional_parameters.push(parameter) end
Adds a block (source code string) to this builder.
# File lib/translator/kernel_builder.rb, line 63 def add_block(block) @blocks.push(block) end
Adds a result that has to be kept on GPU. Therefore additional memory allocations will be made
# File lib/translator/kernel_builder.rb, line 77 def add_cached_result(result_id, type) @cached_results[result_id] = type end
Adds one or multiple methods (source code strings) to this builder.
# File lib/translator/kernel_builder.rb, line 58 def add_methods(*method) @methods.push(*method) end
# File lib/translator/kernel_builder.rb, line 67 def add_previous_kernel_parameter(parameter) @previous_kernel_input.push(parameter) end
# File lib/translator/kernel_builder.rb, line 81 def assert_ready_to_build required_values = [:block_invocation, :result_type] for selector in required_values if send(selector) == nil raise AssertionError.new( "Not ready to build (KernelBuilder): #{selector} is not set") end end end
# File lib/translator/kernel_builder.rb, line 99 def build_blocks return @blocks.join("\n\n") end
# File lib/translator/kernel_builder.rb, line 103 def build_kernel Log.info("Building kernel (num_blocks=#{@blocks.size})") assert_ready_to_build # Build parameters p_env = Constants::ENV_TYPE + " *" + Constants::ENV_IDENTIFIER p_num_threads = Constants::NUM_THREADS_TYPE + " " + Constants::NUM_THREADS_IDENTIFIER p_result = result_type.to_c_type + " *" + Constants::RESULT_IDENTIFIER p_cached_results = cached_results.map do |result_id, type| type.to_c_type + " *" + Constants::RESULT_IDENTIFIER + result_id end cached_results.each do |result_id, type| @execution = execution + "\n" + " " + Constants::RESULT_IDENTIFIER + result_id + "[_tid_] = " + Constants::TEMP_RESULT_IDENTIFIER + result_id + ";" end previous_kernel_params = [] for var in previous_kernel_input previous_kernel_params.push(var.type.to_c_type + " *" + var.name.to_s) end parameters = ([p_env, p_num_threads, p_result] + p_cached_results + previous_kernel_params + additional_parameters).join(", ") # Build kernel return Translator.read_file(file_name: "kernel.cpp", replacements: { "block_invocation" => block_invocation, "execution" => execution, "kernel_name" => kernel_name, "parameters" => parameters, "num_threads" => Constants::NUM_THREADS_IDENTIFIER}) end
— Constructor source code —
# File lib/translator/kernel_builder.rb, line 95 def build_methods return @methods.join("\n\n") end