Differential D2593 Diff 8523 intern/cycles/device/opencl/opencl.h

Changeset View

Standalone View

intern/cycles/device/opencl/opencl.h

Show All 11 Lines
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.		* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and		* See the License for the specific language governing permissions and
* limitations under the License.		* limitations under the License.
*/		*/

#ifdef WITH_OPENCL		#ifdef WITH_OPENCL

#include "device/device.h"		#include "device/device.h"
		#include "device/device_denoising.h"

#include "util/util_map.h"		#include "util/util_map.h"
#include "util/util_param.h"		#include "util/util_param.h"
#include "util/util_string.h"		#include "util/util_string.h"

#include "clew.h"		#include "clew.h"

CCL_NAMESPACE_BEGIN		CCL_NAMESPACE_BEGIN
▲ Show 20 Lines • Show All 99 Lines • ▼ Show 20 Lines	public:

static string get_device_name(cl_device_id device_id);		static string get_device_name(cl_device_id device_id);

static bool get_device_type(cl_device_id device_id,		static bool get_device_type(cl_device_id device_id,
cl_device_type *device_type,		cl_device_type *device_type,
cl_int* error = NULL);		cl_int* error = NULL);
static cl_device_type get_device_type(cl_device_id device_id);		static cl_device_type get_device_type(cl_device_id device_id);

		static int get_base_align_bytes(cl_device_id device_id);

/* Get somewhat more readable device name.		/* Get somewhat more readable device name.
* Main difference is AMD OpenCL here which only gives code name		* Main difference is AMD OpenCL here which only gives code name
* for the regular device name. This will give more sane device		* for the regular device name. This will give more sane device
* name using some extensions.		* name using some extensions.
*/		*/
static string get_readable_device_name(cl_device_id device_id);		static string get_readable_device_name(cl_device_id device_id);
};		};

▲ Show 20 Lines • Show All 137 Lines • ▼ Show 20 Lines	private:

bool use_stdout;		bool use_stdout;
string log, error_msg;		string log, error_msg;
string compile_output;		string compile_output;

map<ustring, cl_kernel> kernels;		map<ustring, cl_kernel> kernels;
};		};

OpenCLProgram base_program;		OpenCLProgram base_program, denoising_program;

typedef map<string, device_vector<uchar>*> ConstMemMap;		typedef map<string, device_vector<uchar>*> ConstMemMap;
typedef map<string, device_ptr> MemMap;		typedef map<string, device_ptr> MemMap;

ConstMemMap const_mem_map;		ConstMemMap const_mem_map;
MemMap mem_map;		MemMap mem_map;
device_ptr null_mem;		device_ptr null_mem;

Show All 21 Lines	public:
virtual bool load_kernels(const DeviceRequestedFeatures& requested_features,		virtual bool load_kernels(const DeviceRequestedFeatures& requested_features,
vector<OpenCLProgram*> &programs) = 0;		vector<OpenCLProgram*> &programs) = 0;

void mem_alloc(const char *name, device_memory& mem, MemoryType type);		void mem_alloc(const char *name, device_memory& mem, MemoryType type);
void mem_copy_to(device_memory& mem);		void mem_copy_to(device_memory& mem);
void mem_copy_from(device_memory& mem, int y, int w, int h, int elem);		void mem_copy_from(device_memory& mem, int y, int w, int h, int elem);
void mem_zero(device_memory& mem);		void mem_zero(device_memory& mem);
void mem_free(device_memory& mem);		void mem_free(device_memory& mem);
		int mem_get_offset_alignment();
		device_ptr mem_get_offset_ptr(device_memory& mem, int offset, int size, MemoryType type);
void const_copy_to(const char name, void host, size_t size);		void const_copy_to(const char name, void host, size_t size);
void tex_alloc(const char *name,		void tex_alloc(const char *name,
device_memory& mem,		device_memory& mem,
InterpolationType /interpolation/,		InterpolationType /interpolation/,
ExtensionType /extension/);		ExtensionType /extension/);
void tex_free(device_memory& mem);		void tex_free(device_memory& mem);

size_t global_size_round_up(int group_size, int global_size);		size_t global_size_round_up(int group_size, int global_size);
void enqueue_kernel(cl_kernel kernel, size_t w, size_t h);		void enqueue_kernel(cl_kernel kernel, size_t w, size_t h, size_t max_workgroup_size = -1);
void set_kernel_arg_mem(cl_kernel kernel, cl_uint narg, const char name);		void set_kernel_arg_mem(cl_kernel kernel, cl_uint narg, const char name);

void film_convert(DeviceTask& task, device_ptr buffer, device_ptr rgba_byte, device_ptr rgba_half);		void film_convert(DeviceTask& task, device_ptr buffer, device_ptr rgba_byte, device_ptr rgba_half);
void shader(DeviceTask& task);		void shader(DeviceTask& task);

		void denoise(RenderTile& tile, const DeviceTask& task);

class OpenCLDeviceTask : public DeviceTask {		class OpenCLDeviceTask : public DeviceTask {
public:		public:
OpenCLDeviceTask(OpenCLDeviceBase *device, DeviceTask& task)		OpenCLDeviceTask(OpenCLDeviceBase *device, DeviceTask& task)
: DeviceTask(task)		: DeviceTask(task)
{		{
run = function_bind(&OpenCLDeviceBase::thread_run,		run = function_bind(&OpenCLDeviceBase::thread_run,
device,		device,
this);		this);
Show All 17 Lines	public:

void task_cancel()		void task_cancel()
{		{
task_pool.cancel();		task_pool.cancel();
}		}

virtual void thread_run(DeviceTask * /task/) = 0;		virtual void thread_run(DeviceTask * /task/) = 0;

		virtual bool is_split_kernel() = 0;

protected:		protected:
string kernel_build_options(const string *debug_src = NULL);		string kernel_build_options(const string *debug_src = NULL);

void mem_zero_kernel(device_ptr ptr, size_t size);		void mem_zero_kernel(device_ptr ptr, size_t size);

		bool denoising_non_local_means(device_ptr image_ptr,
		device_ptr guide_ptr,
		device_ptr variance_ptr,
		device_ptr out_ptr,
		DenoisingTask *task);
		bool denoising_construct_transform(DenoisingTask *task);
		bool denoising_reconstruct(device_ptr color_ptr,
		device_ptr color_variance_ptr,
		device_ptr guide_ptr,
		device_ptr guide_variance_ptr,
		device_ptr output_ptr,
		DenoisingTask *task);
		bool denoising_combine_halves(device_ptr a_ptr,
		device_ptr b_ptr,
		device_ptr mean_ptr,
		device_ptr variance_ptr,
		int r, int4 rect,
		DenoisingTask *task);
		bool denoising_divide_shadow(device_ptr a_ptr,
		device_ptr b_ptr,
		device_ptr sample_variance_ptr,
		device_ptr sv_variance_ptr,
		device_ptr buffer_variance_ptr,
		DenoisingTask *task);
		bool denoising_get_feature(int mean_offset,
		int variance_offset,
		device_ptr mean_ptr,
		device_ptr variance_ptr,
		DenoisingTask *task);
		bool denoising_set_tiles(device_ptr *buffers,
		DenoisingTask *task);

class ArgumentWrapper {		class ArgumentWrapper {
public:		public:
ArgumentWrapper() : size(0), pointer(NULL)		ArgumentWrapper() : size(0), pointer(NULL)
{		{
}		}

ArgumentWrapper(device_memory& argument) : size(sizeof(void*)),		ArgumentWrapper(device_memory& argument) : size(sizeof(void*)),
pointer((void*)(&argument.device_pointer))		pointer((void*)(&argument.device_pointer))
▲ Show 20 Lines • Show All 97 Lines • Show Last 20 Lines