ilusha/blender
1
0
Fork 0
mirror of https://github.com/blender/blender synced 2026-09-29 04:37:17 +03:00
Code Issues Projects Releases Packages Wiki Activity Actions

Refactor: Cycles: More consistent naming of image functions and structs

Browse source 
Previously there was a mix of "image" and "texture" to refer to the same
thing, use "image" when possible now. An exception is MEM_IMAGE_TEXTURE
to avoid conflicts with the MEM_IMAGE macro on Windows.

Pull Request: https://projects.blender.org/blender/blender/pulls/152665
This commit is contained in:
Brecht Van Lommel 2026-01-10 22:45:36 +01:00
parent 850b986d89
commit 527f9ea306
50 changed files with 479 additions and 481 deletions
Download patch file Download diff file Expand all files Collapse all files

49
intern/cycles/device/cpu/device_impl.cpp
View file

@ -36,11 +36,12 @@
#include "util/log.h"
#include "util/progress.h"
#include "util/task.h"
#include "util/types_image.h"
CCL_NAMESPACE_BEGIN
CPUDevice::CPUDevice(const DeviceInfo &info_, Stats &stats_, Profiler &profiler_, bool headless_)
: Device(info_, stats_, profiler_, headless_), texture_info(this, "texture_info", MEM_GLOBAL)
: Device(info_, stats_, profiler_, headless_), image_info(this, "image_info", MEM_GLOBAL)
{
/* Pick any kernel, all of them are supposed to have same level of microarchitecture
* optimization. */
@ -54,7 +55,7 @@ CPUDevice::CPUDevice(const DeviceInfo &info_, Stats &stats_, Profiler &profiler_
#ifdef WITH_EMBREE
embree_device = rtcNewDevice("verbose=0");
#endif
need_texture_info = false;
need_image_info = false;
}
CPUDevice::~CPUDevice()
@ -63,7 +64,7 @@ CPUDevice::~CPUDevice()
rtcReleaseDevice(embree_device);
#endif
texture_info.free();
image_info.free();
}
BVHLayoutMask CPUDevice::get_bvh_layout_mask(uint /*kernel_features*/) const
@ -75,22 +76,22 @@ BVHLayoutMask CPUDevice::get_bvh_layout_mask(uint /*kernel_features*/) const
return bvh_layout_mask;
}
bool CPUDevice::load_texture_info()
bool CPUDevice::load_image_info()
{
if (!need_texture_info) {
if (!need_image_info) {
return false;
}
texture_info.copy_to_device();
need_texture_info = false;
image_info.copy_to_device();
need_image_info = false;
return true;
}
void CPUDevice::mem_alloc(device_memory &mem)
{
if (mem.type == MEM_TEXTURE) {
assert(!"mem_alloc not supported for textures.");
if (mem.type == MEM_IMAGE_TEXTURE) {
assert(!"mem_alloc not supported for images.");
}
else if (mem.type == MEM_GLOBAL) {
assert(!"mem_alloc not supported for global memory.");
@ -123,9 +124,9 @@ void CPUDevice::mem_copy_to(device_memory &mem)
global_free(mem);
global_alloc(mem);
}
else if (mem.type == MEM_TEXTURE) {
tex_free((device_texture &)mem);
tex_alloc((device_texture &)mem);
else if (mem.type == MEM_IMAGE_TEXTURE) {
image_free((device_image &)mem);
image_alloc((device_image &)mem);
}
else {
if (!mem.device_pointer) {
@ -163,8 +164,8 @@ void CPUDevice::mem_free(device_memory &mem)
if (mem.type == MEM_GLOBAL) {
global_free(mem);
}
else if (mem.type == MEM_TEXTURE) {
tex_free((device_texture &)mem);
else if (mem.type == MEM_IMAGE_TEXTURE) {
image_free((device_image &)mem);
}
else if (mem.device_pointer) {
if (mem.type == MEM_DEVICE_ONLY) {
@ -219,7 +220,7 @@ void CPUDevice::global_free(device_memory &mem)
}
}
void CPUDevice::tex_alloc(device_texture &mem)
void CPUDevice::image_alloc(device_image &mem)
{
LOG_DEBUG << "Texture allocate: " << mem.name << ", "
<< string_human_readable_number(mem.memory_size()) << " bytes. ("
@ -230,23 +231,23 @@ void CPUDevice::tex_alloc(device_texture &mem)
stats.mem_alloc(mem.device_size);
const uint slot = mem.slot;
if (slot >= texture_info.size()) {
if (slot >= image_info.size()) {
/* Allocate some slots in advance, to reduce amount of re-allocations. */
texture_info.resize(slot + 128);
image_info.resize(slot + 128);
}
texture_info[slot] = mem.info;
texture_info[slot].data = (uint64_t)mem.host_pointer;
need_texture_info = true;
image_info[slot] = mem.info;
image_info[slot].data = (uint64_t)mem.host_pointer;
need_image_info = true;
}
void CPUDevice::tex_free(device_texture &mem)
void CPUDevice::image_free(device_image &mem)
{
if (mem.device_pointer) {
mem.device_pointer = 0;
stats.mem_free(mem.device_size);
mem.device_size = 0;
need_texture_info = true;
need_image_info = true;
}
}
@ -302,8 +303,8 @@ void *CPUDevice::get_guiding_device() const
void CPUDevice::get_cpu_kernel_thread_globals(
vector<ThreadKernelGlobalsCPU> &kernel_thread_globals)
{
/* Ensure latest texture info is loaded into kernel globals before returning. */
load_texture_info();
/* Ensure latest image info is loaded into kernel globals before returning. */
load_image_info();
kernel_thread_globals.clear();
OSLGlobals *osl_globals = get_cpu_osl_memory();

12
intern/cycles/device/cpu/device_impl.h
View file

@ -38,8 +38,8 @@ class CPUDevice : public Device {
public:
KernelGlobalsCPU kernel_globals;
device_vector<TextureInfo> texture_info;
bool need_texture_info;
device_vector<KernelImageInfo> image_info;
bool need_image_info;
#ifdef WITH_OSL
OSLGlobals osl_globals;
@ -61,9 +61,9 @@ class CPUDevice : public Device {
BVHLayoutMask get_bvh_layout_mask(uint /*kernel_features*/) const override;
/* Returns true if the texture info was copied to the device (meaning, some more
/* Returns true if the image info was copied to the device (meaning, some more
* re-initialization might be needed). */
bool load_texture_info();
bool load_image_info();
void mem_alloc(device_memory &mem) override;
void mem_copy_to(device_memory &mem) override;
@ -79,8 +79,8 @@ class CPUDevice : public Device {
void global_alloc(device_memory &mem);
void global_free(device_memory &mem);
void tex_alloc(device_texture &mem);
void tex_free(device_texture &mem);
void image_alloc(device_image &mem);
void image_free(device_image &mem);
void build_bvh(BVH *bvh, Progress &progress, bool refit) override;

92
intern/cycles/device/cuda/device_impl.cpp
View file

@ -17,9 +17,9 @@
# include "util/path.h"
# include "util/string.h"
# include "util/system.h"
# include "util/texture.h"
# include "util/time.h"
# include "util/types.h"
# include "util/types_image.h"
# ifdef _WIN32
# include "util/windows.h"
@ -70,7 +70,7 @@ CUDADevice::CUDADevice(const DeviceInfo &info, Stats &stats, Profiler &profiler,
cuModule = nullptr;
need_texture_info = false;
need_image_info = false;
pitch_alignment = 0;
@ -133,7 +133,7 @@ CUDADevice::CUDADevice(const DeviceInfo &info, Stats &stats, Profiler &profiler,
CUDADevice::~CUDADevice()
{
texture_info.free();
image_info.free();
if (cuModule) {
cuda_assert(cuModuleUnload(cuModule));
}
@ -173,7 +173,7 @@ bool CUDADevice::check_peer_access(Device *peer_device)
return false;
}
// Ensure array access over the link is possible as well (for 3D textures)
// Ensure array access over the link is possible as well (for 3D images)
cuda_assert(cuDeviceGetP2PAttribute(&can_access,
CU_DEVICE_P2P_ATTRIBUTE_CUDA_ARRAY_ACCESS_SUPPORTED,
cuDevice,
@ -567,8 +567,8 @@ void CUDADevice::copy_host_to_device(void *device_pointer, void *host_pointer, c
void CUDADevice::mem_alloc(device_memory &mem)
{
if (mem.type == MEM_TEXTURE) {
assert(!"mem_alloc not supported for textures.");
if (mem.type == MEM_IMAGE_TEXTURE) {
assert(!"mem_alloc not supported for images.");
}
else if (mem.type == MEM_GLOBAL) {
assert(!"mem_alloc not supported for global memory.");
@ -583,8 +583,8 @@ void CUDADevice::mem_copy_to(device_memory &mem)
if (mem.type == MEM_GLOBAL) {
global_copy_to(mem);
}
else if (mem.type == MEM_TEXTURE) {
tex_copy_to((device_texture &)mem);
else if (mem.type == MEM_IMAGE_TEXTURE) {
image_copy_to((device_image &)mem);
}
else {
if (!mem.device_pointer) {
@ -603,20 +603,20 @@ void CUDADevice::mem_move_to_host(device_memory &mem)
global_free(mem);
global_alloc(mem);
}
else if (mem.type == MEM_TEXTURE) {
tex_free((device_texture &)mem);
tex_alloc((device_texture &)mem);
else if (mem.type == MEM_IMAGE_TEXTURE) {
image_free((device_image &)mem);
image_alloc((device_image &)mem);
}
else {
assert(!"mem_move_to_host only supported for texture and global memory");
assert(!"mem_move_to_host only supported for image and global memory");
}
}
void CUDADevice::mem_copy_from(
device_memory &mem, const size_t y, size_t w, const size_t h, size_t elem)
{
if (mem.type == MEM_TEXTURE || mem.type == MEM_GLOBAL) {
assert(!"mem_copy_from not supported for textures.");
if (mem.type == MEM_IMAGE_TEXTURE) {
assert(!"mem_copy_from not supported for images.");
}
else if (mem.host_pointer) {
const size_t size = elem * w * h;
@ -656,8 +656,8 @@ void CUDADevice::mem_free(device_memory &mem)
if (mem.type == MEM_GLOBAL) {
global_free(mem);
}
else if (mem.type == MEM_TEXTURE) {
tex_free((device_texture &)mem);
else if (mem.type == MEM_IMAGE_TEXTURE) {
image_free((device_image &)mem);
}
else {
generic_free(mem);
@ -720,14 +720,14 @@ void CUDADevice::global_free(device_memory &mem)
}
}
static size_t tex_src_pitch(const device_texture &mem)
static size_t tex_src_pitch(const device_image &mem)
{
return mem.data_width * datatype_size(mem.data_type) * mem.data_elements;
}
static CUDA_MEMCPY2D tex_2d_copy_param(const device_texture &mem, const int pitch_alignment)
static CUDA_MEMCPY2D tex_2d_copy_param(const device_image &mem, const int pitch_alignment)
{
/* 2D texture using pitch aligned linear memory. */
/* 2D image using pitch aligned linear memory. */
const size_t src_pitch = tex_src_pitch(mem);
const size_t dst_pitch = align_up(src_pitch, pitch_alignment);
@ -745,7 +745,7 @@ static CUDA_MEMCPY2D tex_2d_copy_param(const device_texture &mem, const int pitc
return param;
}
void CUDADevice::tex_alloc(device_texture &mem)
void CUDADevice::image_alloc(device_image &mem)
{
CUDAContextScope scope(this);
@ -776,13 +776,15 @@ void CUDADevice::tex_alloc(device_texture &mem)
filter_mode = CU_TR_FILTER_MODE_LINEAR;
}
/* Image Texture Storage */
/* Cycles expects to read all texture data as normalized float values in
/* Image Texture Storage
*
* Cycles expects to read all image data as normalized float values in
* kernel/device/gpu/image.h. But storing all data as floats would be very inefficient due to the
* huge size of float textures. So in the code below, we define different texture types including
* huge size of float image. So in the code below, we define different texture types including
* integer types, with the aim of using CUDA's default promotion behavior of integer data to
* floating point data in the range [0, 1], as noted in the CUDA documentation on
* cuTexObjectCreate API Call.
*
* Note that 32-bit integers are not supported by this promotion behavior and cannot be used
* with Cycles's current implementation in kernel/device/gpu/image.h.
*/
@ -813,7 +815,7 @@ void CUDADevice::tex_alloc(device_texture &mem)
cmem->texobject = 0;
}
else if (mem.data_height > 0) {
/* 2D texture, using pitch aligned linear memory. */
/* 2D image, using pitch aligned linear memory. */
const size_t dst_pitch = align_up(tex_src_pitch(mem), pitch_alignment);
const size_t dst_size = dst_pitch * mem.data_height;
@ -826,7 +828,7 @@ void CUDADevice::tex_alloc(device_texture &mem)
cuda_assert(cuMemcpy2DUnaligned(&param));
}
else {
/* 1D texture, using linear memory. */
/* 1D image, using linear memory. */
cmem = generic_alloc(mem);
if (!cmem) {
return;
@ -836,7 +838,7 @@ void CUDADevice::tex_alloc(device_texture &mem)
}
/* Set Mapping and tag that we need to (re-)upload to device */
TextureInfo tex_info = mem.info;
KernelImageInfo tex_info = mem.info;
if (!is_nanovdb_type(mem.info.data_type)) {
CUDA_RESOURCE_DESC resDesc;
@ -868,7 +870,7 @@ void CUDADevice::tex_alloc(device_texture &mem)
texDesc.addressMode[2] = address_mode;
texDesc.filterMode = filter_mode;
/* CUDA's flag CU_TRSF_READ_AS_INTEGER is intentionally not used and it is
* significant, see above an explanation about how Blender treat textures. */
* significant, see above an explanation about how Blender treat images. */
texDesc.flags = CU_TRSF_NORMALIZED_COORDINATES;
thread_scoped_lock lock(device_mem_map_mutex);
@ -883,33 +885,33 @@ void CUDADevice::tex_alloc(device_texture &mem)
}
{
/* Update texture info. */
thread_scoped_lock lock(texture_info_mutex);
/* Update image info. */
thread_scoped_lock lock(image_info_mutex);
const uint slot = mem.slot;
if (slot >= texture_info.size()) {
if (slot >= image_info.size()) {
/* Allocate some slots in advance, to reduce amount of re-allocations. */
texture_info.resize(slot + 128);
image_info.resize(slot + 128);
}
texture_info[slot] = tex_info;
need_texture_info = true;
image_info[slot] = tex_info;
need_image_info = true;
}
}
void CUDADevice::tex_copy_to(device_texture &mem)
void CUDADevice::image_copy_to(device_image &mem)
{
if (!mem.device_pointer) {
/* Not yet allocated on device. */
tex_alloc(mem);
image_alloc(mem);
}
else if (!mem.is_resident(this)) {
/* Peering with another device, may still need to create texture info and object. */
bool texture_allocated = false;
/* Peering with another device, may still need to create image info and object. */
bool image_allocated = false;
{
thread_scoped_lock lock(texture_info_mutex);
texture_allocated = mem.slot < texture_info.size() && texture_info[mem.slot].data != 0;
thread_scoped_lock lock(image_info_mutex);
image_allocated = mem.slot < image_info.size() && image_info[mem.slot].data != 0;
}
if (!texture_allocated) {
tex_alloc(mem);
if (!image_allocated) {
image_alloc(mem);
}
}
else {
@ -925,7 +927,7 @@ void CUDADevice::tex_copy_to(device_texture &mem)
}
}
void CUDADevice::tex_free(device_texture &mem)
void CUDADevice::image_free(device_image &mem)
{
CUDAContextScope scope(this);
thread_scoped_lock lock(device_mem_map_mutex);
@ -938,10 +940,10 @@ void CUDADevice::tex_free(device_texture &mem)
const Mem &cmem = it->second;
/* Always clear texture info and texture object, regardless of residency. */
/* Always clear image info and image object, regardless of residency. */
{
thread_scoped_lock lock(texture_info_mutex);
texture_info[mem.slot] = TextureInfo();
thread_scoped_lock lock(image_info_mutex);
image_info[mem.slot] = KernelImageInfo();
}
if (cmem.texobject) {

8
intern/cycles/device/cuda/device_impl.h
View file

@ -77,10 +77,10 @@ class CUDADevice : public GPUDevice {
void global_copy_to(device_memory &mem);
void global_free(device_memory &mem);
/* Texture memory. */
void tex_alloc(device_texture &mem);
void tex_copy_to(device_texture &mem);
void tex_free(device_texture &mem);
/* Image memory. */
void image_alloc(device_image &mem);
void image_copy_to(device_image &mem);
void image_free(device_image &mem);
/* Device side memory. */
void get_device_memory_info(size_t &total, size_t &free) override;

12
intern/cycles/device/cuda/queue.cpp
View file

@ -66,7 +66,7 @@ void CUDADeviceQueue::init_execution()
{
/* Synchronize all textures and memory copies before executing task. */
CUDAContextScope scope(cuda_device_);
cuda_device_->load_texture_info();
cuda_device_->load_image_info();
cuda_device_assert(cuda_device_, cuCtxSynchronize());
debug_init_execution();
@ -84,8 +84,8 @@ bool CUDADeviceQueue::enqueue(DeviceKernel kernel,
const CUDAContextScope scope(cuda_device_);
/* Update texture info in case integrator memory alloc caused texture to move to host. */
if (cuda_device_->load_texture_info()) {
/* Update image info in case integrator memory alloc caused texture to move to host. */
if (cuda_device_->load_image_info()) {
cuda_device_assert(cuda_device_, cuCtxSynchronize());
if (cuda_device_->have_error()) {
return false;
@ -151,7 +151,7 @@ bool CUDADeviceQueue::synchronize()
void CUDADeviceQueue::zero_to_device(device_memory &mem)
{
assert(mem.type != MEM_GLOBAL && mem.type != MEM_TEXTURE);
assert(mem.type != MEM_GLOBAL && mem.type != MEM_IMAGE_TEXTURE);
if (mem.memory_size() == 0) {
return;
@ -173,7 +173,7 @@ void CUDADeviceQueue::zero_to_device(device_memory &mem)
void CUDADeviceQueue::copy_to_device(device_memory &mem)
{
assert(mem.type != MEM_GLOBAL && mem.type != MEM_TEXTURE);
assert(mem.type != MEM_GLOBAL && mem.type != MEM_IMAGE_TEXTURE);
if (mem.memory_size() == 0) {
return;
@ -197,7 +197,7 @@ void CUDADeviceQueue::copy_to_device(device_memory &mem)
void CUDADeviceQueue::copy_from_device(device_memory &mem)
{
assert(mem.type != MEM_GLOBAL && mem.type != MEM_TEXTURE);
assert(mem.type != MEM_GLOBAL && mem.type != MEM_IMAGE_TEXTURE);
if (mem.memory_size() == 0) {
return;

30
intern/cycles/device/device.cpp
View file

@ -523,13 +523,13 @@ void Device::host_free(const MemoryType /*type*/, void *host_pointer, const size
GPUDevice::~GPUDevice() noexcept(false) = default;
bool GPUDevice::load_texture_info()
bool GPUDevice::load_image_info()
{
/* Note texture_info is never host mapped, and load_texture_info() should only
/* Note image_info is never host mapped, and load_image_info() should only
* be called right before kernel enqueue when all memory operations have completed. */
if (need_texture_info) {
texture_info.copy_to_device();
need_texture_info = false;
if (need_image_info) {
image_info.copy_to_device();
need_image_info = false;
return true;
}
return false;
@ -563,8 +563,8 @@ void GPUDevice::init_host_memory(const size_t preferred_texture_headroom,
* is space left for it. */
device_working_headroom = preferred_working_headroom > 0 ? preferred_working_headroom :
32 * 1024 * 1024LL; // 32MB
device_texture_headroom = preferred_texture_headroom > 0 ? preferred_texture_headroom :
128 * 1024 * 1024LL; // 128MB
device_image_headroom = preferred_texture_headroom > 0 ? preferred_texture_headroom :
128 * 1024 * 1024LL; // 128MB
LOG_INFO << "Mapped host memory limit set to " << string_human_readable_number(map_host_limit)
<< " bytes. (" << string_human_readable_size(map_host_limit) << ")";
@ -601,8 +601,8 @@ void GPUDevice::move_textures_to_host(size_t size, const size_t headroom, const
continue;
}
const bool is_texture = (mem.type == MEM_TEXTURE || mem.type == MEM_GLOBAL) &&
(&mem != &texture_info);
const bool is_texture = (mem.type == MEM_IMAGE_TEXTURE || mem.type == MEM_GLOBAL) &&
(&mem != &image_info);
const bool is_image = is_texture && (mem.data_height > 1);
/* Can't move this type of memory. */
@ -642,8 +642,8 @@ void GPUDevice::move_textures_to_host(size_t size, const size_t headroom, const
max_mem->move_to_host = false;
size = (max_size >= size) ? 0 : size - max_size;
/* Tag texture info update for new pointers. */
need_texture_info = true;
/* Tag image info update for new pointers. */
need_image_info = true;
}
else {
break;
@ -660,17 +660,17 @@ GPUDevice::Mem *GPUDevice::generic_alloc(device_memory &mem, const size_t pitch_
const char *status = "";
/* First try allocating in device memory, respecting headroom. We make
* an exception for texture info. It is small and frequently accessed,
* an exception for image info. It is small and frequently accessed,
* so treat it as working memory.
*
* If there is not enough room for working memory, we will try to move
* textures to host memory, assuming the performance impact would have
* been worse for working memory. */
const bool is_texture = (mem.type == MEM_TEXTURE || mem.type == MEM_GLOBAL) &&
(&mem != &texture_info);
const bool is_texture = (mem.type == MEM_IMAGE_TEXTURE || mem.type == MEM_GLOBAL) &&
(&mem != &image_info);
const bool is_image = is_texture && (mem.data_height > 1);
const size_t headroom = (is_texture) ? device_texture_headroom : device_working_headroom;
const size_t headroom = (is_texture) ? device_image_headroom : device_working_headroom;
/* Move textures to host memory if needed. */
if (!mem.move_to_host && !is_image && can_map_host) {

16
intern/cycles/device/device.h
View file

@ -15,9 +15,9 @@
#include "util/profiling.h"
#include "util/stats.h"
#include "util/string.h"
#include "util/texture.h"
#include "util/thread.h"
#include "util/types.h"
#include "util/types_image.h"
#include "util/unique_ptr.h"
#include "util/vector.h"
@ -333,7 +333,7 @@ class Device {
class GPUDevice : public Device {
protected:
GPUDevice(const DeviceInfo &info_, Stats &stats_, Profiler &profiler_, bool headless_)
: Device(info_, stats_, profiler_, headless_), texture_info(this, "texture_info", MEM_GLOBAL)
: Device(info_, stats_, profiler_, headless_), image_info(this, "image_info", MEM_GLOBAL)
{
}
@ -341,12 +341,12 @@ class GPUDevice : public Device {
~GPUDevice() noexcept(false) override;
/* For GPUs that can use bindless textures in some way or another. */
device_vector<TextureInfo> texture_info;
thread_mutex texture_info_mutex;
bool need_texture_info = false;
/* Returns true if the texture info was copied to the device (meaning, some more
device_vector<KernelImageInfo> image_info;
thread_mutex image_info_mutex;
bool need_image_info = false;
/* Returns true if the image info was copied to the device (meaning, some more
* re-initialization might be needed). */
virtual bool load_texture_info();
virtual bool load_image_info();
protected:
/* Memory allocation, only accessed through device_memory. */
@ -355,7 +355,7 @@ class GPUDevice : public Device {
bool can_map_host = false;
size_t map_host_used = 0;
size_t map_host_limit = 0;
size_t device_texture_headroom = 0;
size_t device_image_headroom = 0;
size_t device_working_headroom = 0;
using texMemObject = unsigned long long;
using arrayMemObject = unsigned long long;

84
intern/cycles/device/hip/device_impl.cpp
View file

@ -69,7 +69,7 @@ HIPDevice::HIPDevice(const DeviceInfo &info, Stats &stats, Profiler &profiler, b
hipModule = nullptr;
need_texture_info = false;
need_image_info = false;
pitch_alignment = 0;
@ -126,7 +126,7 @@ HIPDevice::HIPDevice(const DeviceInfo &info, Stats &stats, Profiler &profiler, b
HIPDevice::~HIPDevice()
{
texture_info.free();
image_info.free();
if (hipModule) {
hip_assert(hipModuleUnload(hipModule));
}
@ -164,7 +164,7 @@ bool HIPDevice::check_peer_access(Device *peer_device)
return false;
}
// Ensure array access over the link is possible as well (for 3D textures)
// Ensure array access over the link is possible as well (for 3D images)
hip_assert(hipDeviceGetP2PAttribute(
&can_access, hipDevP2PAttrHipArrayAccessSupported, hipDevice, peer_device_hip->hipDevice));
if (can_access == 0) {
@ -526,8 +526,8 @@ void HIPDevice::copy_host_to_device(void *device_pointer, void *host_pointer, co
void HIPDevice::mem_alloc(device_memory &mem)
{
if (mem.type == MEM_TEXTURE) {
assert(!"mem_alloc not supported for textures.");
if (mem.type == MEM_IMAGE_TEXTURE) {
assert(!"mem_alloc not supported for images.");
}
else if (mem.type == MEM_GLOBAL) {
assert(!"mem_alloc not supported for global memory.");
@ -542,8 +542,8 @@ void HIPDevice::mem_copy_to(device_memory &mem)
if (mem.type == MEM_GLOBAL) {
global_copy_to(mem);
}
else if (mem.type == MEM_TEXTURE) {
tex_copy_to((device_texture &)mem);
else if (mem.type == MEM_IMAGE_TEXTURE) {
image_copy_to((device_image &)mem);
}
else {
if (!mem.device_pointer) {
@ -562,20 +562,20 @@ void HIPDevice::mem_move_to_host(device_memory &mem)
global_free(mem);
global_alloc(mem);
}
else if (mem.type == MEM_TEXTURE) {
tex_free((device_texture &)mem);
tex_alloc((device_texture &)mem);
else if (mem.type == MEM_IMAGE_TEXTURE) {
image_free((device_image &)mem);
image_alloc((device_image &)mem);
}
else {
assert(!"mem_move_to_host only supported for texture and global memory");
assert(!"mem_move_to_host only supported for image and global memory");
}
}
void HIPDevice::mem_copy_from(
device_memory &mem, const size_t y, size_t w, const size_t h, size_t elem)
{
if (mem.type == MEM_TEXTURE || mem.type == MEM_GLOBAL) {
assert(!"mem_copy_from not supported for textures.");
if (mem.type == MEM_IMAGE_TEXTURE) {
assert(!"mem_copy_from not supported for images.");
}
else if (mem.host_pointer) {
const size_t size = elem * w * h;
@ -615,8 +615,8 @@ void HIPDevice::mem_free(device_memory &mem)
if (mem.type == MEM_GLOBAL) {
global_free(mem);
}
else if (mem.type == MEM_TEXTURE) {
tex_free((device_texture &)mem);
else if (mem.type == MEM_IMAGE_TEXTURE) {
image_free((device_image &)mem);
}
else {
generic_free(mem);
@ -679,14 +679,14 @@ void HIPDevice::global_free(device_memory &mem)
}
}
static size_t tex_src_pitch(const device_texture &mem)
static size_t tex_src_pitch(const device_image &mem)
{
return mem.data_width * datatype_size(mem.data_type) * mem.data_elements;
}
static hip_Memcpy2D tex_2d_copy_param(const device_texture &mem, const int pitch_alignment)
static hip_Memcpy2D tex_2d_copy_param(const device_image &mem, const int pitch_alignment)
{
/* 2D texture using pitch aligned linear memory. */
/* 2D image using pitch aligned linear memory. */
const size_t src_pitch = tex_src_pitch(mem);
const size_t dst_pitch = align_up(src_pitch, pitch_alignment);
@ -704,7 +704,7 @@ static hip_Memcpy2D tex_2d_copy_param(const device_texture &mem, const int pitch
return param;
}
void HIPDevice::tex_alloc(device_texture &mem)
void HIPDevice::image_alloc(device_image &mem)
{
HIPContextScope scope(this);
@ -735,7 +735,7 @@ void HIPDevice::tex_alloc(device_texture &mem)
filter_mode = hipFilterModeLinear;
}
/* Image Texture Storage */
/* Image image Storage */
hipArray_Format format;
switch (mem.data_type) {
case TYPE_UCHAR:
@ -769,7 +769,7 @@ void HIPDevice::tex_alloc(device_texture &mem)
cmem->texobject = 0;
}
else if (mem.data_height > 0) {
/* 2D texture, using pitch aligned linear memory. */
/* 2D image, using pitch aligned linear memory. */
const size_t dst_pitch = align_up(tex_src_pitch(mem), pitch_alignment);
const size_t dst_size = dst_pitch * mem.data_height;
@ -782,7 +782,7 @@ void HIPDevice::tex_alloc(device_texture &mem)
hip_assert(hipDrvMemcpy2DUnaligned(&param));
}
else {
/* 1D texture, using linear memory. */
/* 1D image, using linear memory. */
cmem = generic_alloc(mem);
if (!cmem) {
return;
@ -792,7 +792,7 @@ void HIPDevice::tex_alloc(device_texture &mem)
}
/* Set Mapping and tag that we need to (re-)upload to device */
TextureInfo tex_info = mem.info;
KernelImageInfo tex_info = mem.info;
if (!is_nanovdb_type(mem.info.data_type)) {
/* Bindless textures. */
@ -831,7 +831,7 @@ void HIPDevice::tex_alloc(device_texture &mem)
if (hipTexObjectCreate(&cmem->texobject, &resDesc, &texDesc, nullptr) != hipSuccess) {
set_error(
"Failed to create texture. Maximum GPU texture size or available GPU memory was likely "
"Failed to create image. Maximum GPU image size or available GPU memory was likely "
"exceeded.");
}
@ -842,33 +842,33 @@ void HIPDevice::tex_alloc(device_texture &mem)
}
{
/* Update texture info. */
thread_scoped_lock lock(texture_info_mutex);
/* Update image info. */
thread_scoped_lock lock(image_info_mutex);
const uint slot = mem.slot;
if (slot >= texture_info.size()) {
if (slot >= image_info.size()) {
/* Allocate some slots in advance, to reduce amount of re-allocations. */
texture_info.resize(slot + 128);
image_info.resize(slot + 128);
}
texture_info[slot] = tex_info;
need_texture_info = true;
image_info[slot] = tex_info;
need_image_info = true;
}
}
void HIPDevice::tex_copy_to(device_texture &mem)
void HIPDevice::image_copy_to(device_image &mem)
{
if (!mem.device_pointer) {
/* Not yet allocated on device. */
tex_alloc(mem);
image_alloc(mem);
}
else if (!mem.is_resident(this)) {
/* Peering with another device, may still need to create texture info and object. */
bool texture_allocated = false;
/* Peering with another device, may still need to create image info and object. */
bool image_allocated = false;
{
thread_scoped_lock lock(texture_info_mutex);
texture_allocated = mem.slot < texture_info.size() && texture_info[mem.slot].data != 0;
thread_scoped_lock lock(image_info_mutex);
image_allocated = mem.slot < image_info.size() && image_info[mem.slot].data != 0;
}
if (!texture_allocated) {
tex_alloc(mem);
if (!image_allocated) {
image_alloc(mem);
}
}
else {
@ -884,7 +884,7 @@ void HIPDevice::tex_copy_to(device_texture &mem)
}
}
void HIPDevice::tex_free(device_texture &mem)
void HIPDevice::image_free(device_image &mem)
{
HIPContextScope scope(this);
thread_scoped_lock lock(device_mem_map_mutex);
@ -897,10 +897,10 @@ void HIPDevice::tex_free(device_texture &mem)
const Mem &cmem = it->second;
/* Always clear texture info and texture object, regardless of residency. */
/* Always clear image info and texture object, regardless of residency. */
{
thread_scoped_lock lock(texture_info_mutex);
texture_info[mem.slot] = TextureInfo();
thread_scoped_lock lock(image_info_mutex);
image_info[mem.slot] = KernelImageInfo();
}
if (cmem.texobject) {

8
intern/cycles/device/hip/device_impl.h
View file

@ -74,10 +74,10 @@ class HIPDevice : public GPUDevice {
void global_copy_to(device_memory &mem);
void global_free(device_memory &mem);
/* Texture memory. */
void tex_alloc(device_texture &mem);
void tex_copy_to(device_texture &mem);
void tex_free(device_texture &mem);
/* Image memory. */
void image_alloc(device_image &mem);
void image_copy_to(device_image &mem);
void image_free(device_image &mem);
/* Device side memory. */
void get_device_memory_info(size_t &total, size_t &free) override;

12
intern/cycles/device/hip/queue.cpp
View file

@ -66,7 +66,7 @@ void HIPDeviceQueue::init_execution()
{
/* Synchronize all textures and memory copies before executing task. */
HIPContextScope scope(hip_device_);
hip_device_->load_texture_info();
hip_device_->load_image_info();
hip_device_assert(hip_device_, hipDeviceSynchronize());
debug_init_execution();
@ -84,8 +84,8 @@ bool HIPDeviceQueue::enqueue(DeviceKernel kernel,
const HIPContextScope scope(hip_device_);
/* Update texture info in case memory moved to host. */
if (hip_device_->load_texture_info()) {
/* Update image info in case memory moved to host. */
if (hip_device_->load_image_info()) {
hip_device_assert(hip_device_, hipDeviceSynchronize());
if (hip_device_->have_error()) {
return false;
@ -149,7 +149,7 @@ bool HIPDeviceQueue::synchronize()
void HIPDeviceQueue::zero_to_device(device_memory &mem)
{
assert(mem.type != MEM_GLOBAL && mem.type != MEM_TEXTURE);
assert(mem.type != MEM_GLOBAL && mem.type != MEM_IMAGE_TEXTURE);
if (mem.memory_size() == 0) {
return;
@ -171,7 +171,7 @@ void HIPDeviceQueue::zero_to_device(device_memory &mem)
void HIPDeviceQueue::copy_to_device(device_memory &mem)
{
assert(mem.type != MEM_GLOBAL && mem.type != MEM_TEXTURE);
assert(mem.type != MEM_GLOBAL && mem.type != MEM_IMAGE_TEXTURE);
if (mem.memory_size() == 0) {
return;
@ -195,7 +195,7 @@ void HIPDeviceQueue::copy_to_device(device_memory &mem)
void HIPDeviceQueue::copy_from_device(device_memory &mem)
{
assert(mem.type != MEM_GLOBAL && mem.type != MEM_TEXTURE);
assert(mem.type != MEM_GLOBAL && mem.type != MEM_IMAGE_TEXTURE);
if (mem.memory_size() == 0) {
return;

24
intern/cycles/device/memory.cpp
View file

@ -73,7 +73,7 @@ void device_memory::host_and_device_free()
void device_memory::device_alloc()
{
assert(!device_pointer && type != MEM_TEXTURE && type != MEM_GLOBAL);
assert(!device_pointer && type != MEM_IMAGE_TEXTURE && type != MEM_GLOBAL);
device->mem_alloc(*this);
}
@ -93,7 +93,7 @@ void device_memory::device_move_to_host()
void device_memory::device_copy_from(const size_t y, const size_t w, size_t h, const size_t elem)
{
assert(type != MEM_TEXTURE && type != MEM_READ_ONLY && type != MEM_GLOBAL);
assert(type != MEM_IMAGE_TEXTURE && type != MEM_READ_ONLY);
device->mem_copy_from(*this, y, w, h, elem);
}
@ -154,13 +154,13 @@ device_sub_ptr::~device_sub_ptr()
/* Device Texture */
device_texture::device_texture(Device *device,
const char *name,
const uint slot,
ImageDataType image_data_type,
InterpolationType interpolation,
ExtensionType extension)
: device_memory(device, name, MEM_TEXTURE), slot(slot)
device_image::device_image(Device *device,
const char *name,
const uint slot,
ImageDataType image_data_type,
InterpolationType interpolation,
ExtensionType extension)
: device_memory(device, name, MEM_IMAGE_TEXTURE), slot(slot)
{
switch (image_data_type) {
case IMAGE_DATA_TYPE_FLOAT4:
@ -211,13 +211,13 @@ device_texture::device_texture(Device *device,
info.extension = extension;
}
device_texture::~device_texture()
device_image::~device_image()
{
host_and_device_free();
}
/* Host memory allocation. */
void *device_texture::alloc(const size_t width, const size_t height)
void *device_image::alloc(const size_t width, const size_t height)
{
const size_t new_size = size(width, height);
@ -237,7 +237,7 @@ void *device_texture::alloc(const size_t width, const size_t height)
return host_pointer;
}
void device_texture::copy_to_device()
void device_image::copy_to_device()
{
device_copy_to();
}

30
intern/cycles/device/memory.h
View file

@ -11,8 +11,8 @@
#include "util/array.h"
#include "util/half.h"
#include "util/string.h"
#include "util/texture.h"
#include "util/types.h"
#include "util/types_image.h"
CCL_NAMESPACE_BEGIN
@ -30,7 +30,7 @@ enum MemoryType {
MEM_READ_WRITE,
MEM_DEVICE_ONLY,
MEM_GLOBAL,
MEM_TEXTURE,
MEM_IMAGE_TEXTURE,
};
/* Supported Data Types */
@ -581,25 +581,31 @@ class device_sub_ptr {
device_ptr ptr;
};
/* Device Texture
/* Device Image
*
* 2D or 3D image texture memory. */
class device_texture : public device_memory {
class device_image : public device_memory {
public:
device_texture(Device *device,
const char *name,
const uint slot,
ImageDataType image_data_type,
InterpolationType interpolation,
ExtensionType extension);
~device_texture() override;
device_image(Device *device,
const char *name,
const uint slot,
ImageDataType image_data_type,
InterpolationType interpolation,
ExtensionType extension);
~device_image() override;
void *alloc(const size_t width, const size_t height);
template<typename T = void> T *data()
{
return reinterpret_cast<T *>(host_pointer);
}
void copy_to_device();
uint slot = 0;
TextureInfo info;
KernelImageInfo info;
protected:
size_t size(const size_t width, const size_t height)

18
intern/cycles/device/metal/device_impl.h
View file

@ -77,10 +77,10 @@ class MetalDevice : public Device {
std::recursive_mutex metal_mem_map_mutex;
/* Bindless Textures */
bool is_texture(const TextureInfo &tex);
device_vector<TextureInfo> texture_info;
id<MTLBuffer> texture_bindings = nil;
std::vector<id<MTLResource>> texture_slot_map;
bool is_texture(const KernelImageInfo &info);
device_vector<KernelImageInfo> image_info;
id<MTLBuffer> image_bindings = nil;
std::vector<id<MTLResource>> image_slot_map;
MetalPipelineType kernel_specialization_level = PSO_GENERIC;
@ -124,7 +124,7 @@ class MetalDevice : public Device {
bool load_kernels(const uint kernel_features) override;
void load_texture_info();
void load_image_info();
void erase_allocation(device_memory &mem);
@ -178,10 +178,10 @@ class MetalDevice : public Device {
void global_alloc(device_memory &mem);
void global_free(device_memory &mem);
void tex_alloc(device_texture &mem);
void tex_alloc_as_buffer(device_texture &mem);
void tex_copy_to(device_texture &mem);
void tex_free(device_texture &mem);
void image_alloc(device_image &mem);
void image_alloc_as_buffer(device_image &mem);
void image_copy_to(device_image &mem);
void image_free(device_image &mem);
void flush_delayed_free_list();

99
intern/cycles/device/metal/device_impl.mm
View file

@ -67,7 +67,7 @@ void MetalDevice::set_error(const string &error)
}
MetalDevice::MetalDevice(const DeviceInfo &info, Stats &stats, Profiler &profiler, bool headless)
: Device(info, stats, profiler, headless), texture_info(this, "texture_info", MEM_GLOBAL)
: Device(info, stats, profiler, headless), image_info(this, "image_info", MEM_GLOBAL)
{
@autoreleasepool {
{
@ -163,8 +163,8 @@ MetalDevice::MetalDevice(const DeviceInfo &info, Stats &stats, Profiler &profile
kernel_type_as_string(
(MetalPipelineType)min((int)kernel_specialization_level, (int)PSO_NUM - 1)));
texture_bindings = [mtlDevice newBufferWithLength:8192 options:MTLResourceStorageModeShared];
stats.mem_alloc(texture_bindings.allocatedSize);
image_bindings = [mtlDevice newBufferWithLength:8192 options:MTLResourceStorageModeShared];
stats.mem_alloc(image_bindings.allocatedSize);
launch_params_buffer = [mtlDevice newBufferWithLength:sizeof(KernelParamsMetal)
options:MTLResourceStorageModeShared];
@ -194,9 +194,9 @@ MetalDevice::~MetalDevice()
thread_scoped_lock lock(existing_devices_mutex);
/* Release textures that weren't already freed by tex_free. */
for (int res = 0; res < texture_info.size(); res++) {
[texture_slot_map[res] release];
texture_slot_map[res] = nil;
for (int res = 0; res < image_info.size(); res++) {
[image_slot_map[res] release];
image_slot_map[res] = nil;
}
free_bvh();
@ -205,8 +205,8 @@ MetalDevice::~MetalDevice()
stats.mem_free(sizeof(KernelParamsMetal));
[launch_params_buffer release];
stats.mem_free(texture_bindings.allocatedSize);
[texture_bindings release];
stats.mem_free(image_bindings.allocatedSize);
[image_bindings release];
[mtlComputeCommandQueue release];
[mtlGeneralCommandQueue release];
@ -215,7 +215,7 @@ MetalDevice::~MetalDevice()
}
[mtlDevice release];
texture_info.free();
image_info.free();
}
bool MetalDevice::support_device(const uint /*kernel_features*/)
@ -274,7 +274,7 @@ string MetalDevice::preprocess_source(MetalPipelineType pso_type,
}
# ifdef WITH_NANOVDB
/* Compiling in NanoVDB results in a marginal drop in render performance,
* so disable it for specialized PSOs when no textures are using it. */
* so disable it for specialized PSOs when no images are using it. */
if ((pso_type == PSO_GENERIC || using_nanovdb) && DebugFlags().metal.use_nanovdb) {
global_defines += "#define WITH_NANOVDB\n";
}
@ -548,13 +548,11 @@ void MetalDevice::compile_and_load(const int device_id, MetalPipelineType pso_ty
}
}
bool MetalDevice::is_texture(const TextureInfo &tex)
bool MetalDevice::is_texture(const KernelImageInfo &info)
{
return tex.height > 0;
return info.height > 0;
}
void MetalDevice::load_texture_info() {}
void MetalDevice::erase_allocation(device_memory &mem)
{
stats.mem_free(mem.device_size);
@ -660,7 +658,7 @@ MetalDevice::MetalMem *MetalDevice::generic_alloc(device_memory &mem)
}
}
void MetalDevice::generic_copy_to(device_memory &)
void MetalDevice::generic_copy_to(device_memory & /*mem*/)
{
/* No need to copy - Apple Silicon has Unified Memory Architecture. */
}
@ -711,8 +709,8 @@ void MetalDevice::generic_free(device_memory &mem)
void MetalDevice::mem_alloc(device_memory &mem)
{
if (mem.type == MEM_TEXTURE) {
assert(!"mem_alloc not supported for textures.");
if (mem.type == MEM_IMAGE_TEXTURE) {
assert(!"mem_alloc not supported for images.");
}
else if (mem.type == MEM_GLOBAL) {
generic_alloc(mem);
@ -728,8 +726,8 @@ void MetalDevice::mem_copy_to(device_memory &mem)
if (mem.type == MEM_GLOBAL) {
global_alloc(mem);
}
else if (mem.type == MEM_TEXTURE) {
tex_alloc((device_texture &)mem);
else if (mem.type == MEM_IMAGE_TEXTURE) {
image_alloc((device_image &)mem);
}
else {
generic_alloc(mem);
@ -740,8 +738,8 @@ void MetalDevice::mem_copy_to(device_memory &mem)
if (mem.type == MEM_GLOBAL) {
generic_copy_to(mem);
}
else if (mem.type == MEM_TEXTURE) {
tex_copy_to((device_texture &)mem);
else if (mem.type == MEM_IMAGE_TEXTURE) {
image_copy_to((device_image &)mem);
}
else {
generic_copy_to(mem);
@ -755,7 +753,8 @@ void MetalDevice::mem_move_to_host(device_memory & /*mem*/)
assert(!"Metal does not support mem_move_to_host");
}
void MetalDevice::mem_copy_from(device_memory &, const size_t, size_t, const size_t, size_t)
void MetalDevice::mem_copy_from(
device_memory & /*mem*/, const size_t /*y*/, size_t /*w*/, const size_t /*h*/, size_t /*elem*/)
{
/* No need to copy - Apple Silicon has Unified Memory Architecture. */
}
@ -774,8 +773,8 @@ void MetalDevice::mem_free(device_memory &mem)
if (mem.type == MEM_GLOBAL) {
global_free(mem);
}
else if (mem.type == MEM_TEXTURE) {
tex_free((device_texture &)mem);
else if (mem.type == MEM_IMAGE_TEXTURE) {
image_free((device_image &)mem);
}
else {
generic_free(mem);
@ -974,29 +973,29 @@ void MetalDevice::global_free(device_memory &mem)
}
}
void MetalDevice::tex_alloc_as_buffer(device_texture &mem)
void MetalDevice::image_alloc_as_buffer(device_image &mem)
{
MetalDevice::MetalMem *mmem = generic_alloc(mem);
generic_copy_to(mem);
/* Resize once */
const uint slot = mem.slot;
if (slot >= texture_info.size()) {
if (slot >= image_info.size()) {
/* Allocate some slots in advance, to reduce amount
* of re-allocations. */
texture_info.resize(round_up(slot + 1, 128));
texture_slot_map.resize(round_up(slot + 1, 128));
image_info.resize(round_up(slot + 1, 128));
image_slot_map.resize(round_up(slot + 1, 128));
}
texture_info[slot] = mem.info;
texture_slot_map[slot] = mmem->mtlBuffer;
image_info[slot] = mem.info;
image_slot_map[slot] = mmem->mtlBuffer;
if (is_nanovdb_type(mem.info.data_type)) {
using_nanovdb = true;
}
}
void MetalDevice::tex_alloc(device_texture &mem)
void MetalDevice::image_alloc(device_image &mem)
{
@autoreleasepool {
/* Check that dimensions fit within maximum allowable size.
@ -1109,8 +1108,8 @@ void MetalDevice::tex_alloc(device_texture &mem)
bytesPerRow:src_pitch];
}
else {
/* 1D texture, using linear memory. */
tex_alloc_as_buffer(mem);
/* 1D image, using linear memory. */
image_alloc_as_buffer(mem);
return;
}
@ -1126,29 +1125,29 @@ void MetalDevice::tex_alloc(device_texture &mem)
/* Resize once */
const uint slot = mem.slot;
if (slot >= texture_info.size()) {
if (slot >= image_info.size()) {
/* Allocate some slots in advance, to reduce amount
* of re-allocations. */
texture_info.resize(slot + 128);
texture_slot_map.resize(slot + 128);
image_info.resize(slot + 128);
image_slot_map.resize(slot + 128);
ssize_t min_buffer_length = sizeof(void *) * texture_info.size();
if (!texture_bindings || (texture_bindings.length < min_buffer_length)) {
if (texture_bindings) {
delayed_free_list.push_back(texture_bindings);
stats.mem_free(texture_bindings.allocatedSize);
ssize_t min_buffer_length = sizeof(void *) * image_info.size();
if (!image_bindings || (image_bindings.length < min_buffer_length)) {
if (image_bindings) {
delayed_free_list.push_back(image_bindings);
stats.mem_free(image_bindings.allocatedSize);
}
texture_bindings = [mtlDevice newBufferWithLength:min_buffer_length
options:MTLResourceStorageModeShared];
image_bindings = [mtlDevice newBufferWithLength:min_buffer_length
options:MTLResourceStorageModeShared];
stats.mem_alloc(texture_bindings.allocatedSize);
stats.mem_alloc(image_bindings.allocatedSize);
}
}
/* Set Mapping. */
texture_slot_map[slot] = mtlTexture;
texture_info[slot] = mem.info;
texture_info[slot].data = uint64_t(slot) | (sampler_index << 32);
image_slot_map[slot] = mtlTexture;
image_info[slot] = mem.info;
image_info[slot].data = uint64_t(slot) | (sampler_index << 32);
if (max_working_set_exceeded()) {
set_error("System is out of GPU memory");
@ -1156,7 +1155,7 @@ void MetalDevice::tex_alloc(device_texture &mem)
}
}
void MetalDevice::tex_copy_to(device_texture &mem)
void MetalDevice::image_copy_to(device_image &mem)
{
if (mem.is_resident(this)) {
const size_t src_pitch = mem.data_width * datatype_size(mem.data_type) * mem.data_elements;
@ -1178,7 +1177,7 @@ void MetalDevice::tex_copy_to(device_texture &mem)
}
}
void MetalDevice::tex_free(device_texture &mem)
void MetalDevice::image_free(device_image &mem)
{
int slot = mem.slot;
if (mem.data_height == 0) {
@ -1193,7 +1192,7 @@ void MetalDevice::tex_free(device_texture &mem)
mmem.mtlTexture = nil;
erase_allocation(mem);
}
texture_slot_map[slot] = nil;
image_slot_map[slot] = nil;
}
unique_ptr<DeviceQueue> MetalDevice::gpu_queue_create()

32
intern/cycles/device/metal/queue.mm
View file

@ -359,24 +359,24 @@ void MetalDeviceQueue::init_execution()
write_resource(blas_array, metal_device_->blas_array[slot], slot);
}
device_vector<TextureInfo> &texture_info = metal_device_->texture_info;
id<MTLBuffer> &texture_bindings = metal_device_->texture_bindings;
std::vector<id<MTLResource>> &texture_slot_map = metal_device_->texture_slot_map;
device_vector<KernelImageInfo> &image_info = metal_device_->image_info;
id<MTLBuffer> &image_bindings = metal_device_->image_bindings;
std::vector<id<MTLResource>> &image_slot_map = metal_device_->image_slot_map;
/* Ensure texture_info is allocated before populating. */
texture_info.copy_to_device();
/* Ensure image_info is allocated before populating. */
image_info.copy_to_device();
/* Populate texture bindings. */
uint64_t *bindings = (uint64_t *)texture_bindings.contents;
memset(bindings, 0, texture_bindings.length);
for (int slot = 0; slot < texture_info.size(); ++slot) {
if (texture_slot_map[slot]) {
if (metal_device_->is_texture(texture_info[slot])) {
write_resource(bindings, id<MTLTexture>(texture_slot_map[slot]), slot);
uint64_t *bindings = (uint64_t *)image_bindings.contents;
memset(bindings, 0, image_bindings.length);
for (int slot = 0; slot < image_info.size(); ++slot) {
if (image_slot_map[slot]) {
if (metal_device_->is_texture(image_info[slot])) {
write_resource(bindings, id<MTLTexture>(image_slot_map[slot]), slot);
}
else {
/* The GPU address of a 1D buffer texture is written into the slot data field. */
write_resource(&texture_info[slot].data, id<MTLBuffer>(texture_slot_map[slot]), 0);
write_resource(&image_info[slot].data, id<MTLBuffer>(image_slot_map[slot]), 0);
}
}
}
@ -450,7 +450,7 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
/* Encode ancillaries */
int ancillary_index = 0;
write_resource(ancillary_args, metal_device_->texture_bindings, ancillary_index++);
write_resource(ancillary_args, metal_device_->image_bindings, ancillary_index++);
if (metal_device_->use_metalrt) {
write_resource(ancillary_args, metal_device_->accel_struct, ancillary_index++);
@ -673,7 +673,7 @@ void MetalDeviceQueue::zero_to_device(device_memory &mem)
return;
}
assert(mem.type != MEM_GLOBAL && mem.type != MEM_TEXTURE);
assert(mem.type != MEM_GLOBAL && mem.type != MEM_IMAGE_TEXTURE);
if (mem.memory_size() == 0) {
return;
@ -735,7 +735,7 @@ void MetalDeviceQueue::prepare_resources(DeviceKernel /*kernel*/)
device_memory *mem = it.first;
MTLResourceUsage usage = MTLResourceUsageRead;
if (mem->type != MEM_GLOBAL && mem->type != MEM_READ_ONLY && mem->type != MEM_TEXTURE) {
if (mem->type != MEM_GLOBAL && mem->type != MEM_READ_ONLY && mem->type != MEM_IMAGE_TEXTURE) {
usage |= MTLResourceUsageWrite;
}
@ -750,7 +750,7 @@ void MetalDeviceQueue::prepare_resources(DeviceKernel /*kernel*/)
}
/* ancillaries */
[mtlComputeEncoder_ useResource:metal_device_->texture_bindings usage:MTLResourceUsageRead];
[mtlComputeEncoder_ useResource:metal_device_->image_bindings usage:MTLResourceUsageRead];
}
id<MTLComputeCommandEncoder> MetalDeviceQueue::get_compute_encoder(DeviceKernel kernel)

7
intern/cycles/device/multi/device.cpp
View file

@ -15,6 +15,7 @@
#include "util/list.h"
#include "util/map.h"
#include "util/types_image.h"
CCL_NAMESPACE_BEGIN
@ -402,7 +403,7 @@ class MultiDevice : public Device {
owner_sub->device->mem_copy_to(mem);
owner_sub->ptr_map[key] = mem.device_pointer;
if (mem.type == MEM_GLOBAL || mem.type == MEM_TEXTURE) {
if (mem.type == MEM_GLOBAL || mem.type == MEM_IMAGE_TEXTURE) {
/* Need to create texture objects and update pointer in kernel globals on all devices */
for (SubDevice *island_sub : island) {
if (island_sub != owner_sub) {
@ -419,7 +420,7 @@ class MultiDevice : public Device {
void mem_move_to_host(device_memory &mem) override
{
assert(mem.type == MEM_GLOBAL || mem.type == MEM_TEXTURE);
assert(mem.type == MEM_GLOBAL || mem.type == MEM_IMAGE_TEXTURE);
device_ptr existing_key = mem.device_pointer;
device_ptr key = (existing_key) ? existing_key : unique_key++;
@ -526,7 +527,7 @@ class MultiDevice : public Device {
owner_sub->device->mem_free(mem);
owner_sub->ptr_map.erase(owner_sub->ptr_map.find(key));
if (mem.type == MEM_TEXTURE) {
if (mem.type == MEM_IMAGE_TEXTURE) {
/* Free texture objects on all devices */
for (SubDevice *island_sub : island) {
if (island_sub != owner_sub) {

81
intern/cycles/device/oneapi/device_impl.cpp
View file

@ -57,7 +57,7 @@ OneapiDevice::OneapiDevice(const DeviceInfo &info, Stats &stats, Profiler &profi
static_assert(sizeof(arrayMemObject) ==
sizeof(sycl::ext::oneapi::experimental::image_mem_handle));
need_texture_info = false;
need_image_info = false;
use_hardware_raytracing = info.use_hardware_raytracing;
oneapi_set_error_cb(queue_error_cb, &oneapi_error_string_);
@ -116,7 +116,7 @@ OneapiDevice::OneapiDevice(const DeviceInfo &info, Stats &stats, Profiler &profi
if (headroom_str != nullptr) {
const long long override_headroom = (float)atoll(headroom_str);
device_working_headroom = override_headroom;
device_texture_headroom = override_headroom;
device_image_headroom = override_headroom;
}
LOG_TRACE << "oneAPI memory headroom size: "
<< string_human_readable_size(device_working_headroom);
@ -130,7 +130,7 @@ OneapiDevice::~OneapiDevice()
}
# endif
texture_info.free();
image_info.free();
usm_free(device_queue_, kg_memory_);
usm_free(device_queue_, kg_memory_device_);
@ -421,8 +421,8 @@ void OneapiDevice::host_free(const MemoryType type, void *host_pointer, const si
void OneapiDevice::mem_alloc(device_memory &mem)
{
if (mem.type == MEM_TEXTURE) {
assert(!"mem_alloc not supported for textures.");
if (mem.type == MEM_IMAGE_TEXTURE) {
assert(!"mem_alloc not supported for images.");
}
else if (mem.type == MEM_GLOBAL) {
assert(!"mem_alloc not supported for global memory.");
@ -454,8 +454,8 @@ void OneapiDevice::mem_copy_to(device_memory &mem)
if (mem.type == MEM_GLOBAL) {
global_copy_to(mem);
}
else if (mem.type == MEM_TEXTURE) {
tex_copy_to((device_texture &)mem);
else if (mem.type == MEM_IMAGE_TEXTURE) {
image_copy_to((device_image &)mem);
}
else {
if (!mem.device_pointer) {
@ -483,9 +483,9 @@ void OneapiDevice::mem_move_to_host(device_memory &mem)
global_free(mem);
global_alloc(mem);
}
else if (mem.type == MEM_TEXTURE) {
tex_free((device_texture &)mem);
tex_alloc((device_texture &)mem);
else if (mem.type == MEM_IMAGE_TEXTURE) {
image_free((device_image &)mem);
image_alloc((device_image &)mem);
}
else {
assert(0);
@ -495,8 +495,8 @@ void OneapiDevice::mem_move_to_host(device_memory &mem)
void OneapiDevice::mem_copy_from(
device_memory &mem, const size_t y, size_t w, const size_t h, size_t elem)
{
if (mem.type == MEM_TEXTURE || mem.type == MEM_GLOBAL) {
assert(!"mem_copy_from not supported for textures.");
if (mem.type == MEM_IMAGE_TEXTURE) {
assert(!"mem_copy_from not supported for images.");
}
else if (mem.host_pointer) {
const size_t size = (w > 0 || h > 0 || elem > 0) ? (elem * w * h) : mem.memory_size();
@ -571,8 +571,8 @@ void OneapiDevice::mem_free(device_memory &mem)
if (mem.type == MEM_GLOBAL) {
global_free(mem);
}
else if (mem.type == MEM_TEXTURE) {
tex_free((device_texture &)mem);
else if (mem.type == MEM_IMAGE_TEXTURE) {
image_free((device_image &)mem);
}
else {
generic_free(mem);
@ -670,7 +670,7 @@ void OneapiDevice::global_free(device_memory &mem)
}
}
static sycl::ext::oneapi::experimental::image_descriptor image_desc(const device_texture &mem)
static sycl::ext::oneapi::experimental::image_descriptor image_desc(const device_image &mem)
{
/* Image Texture Storage */
sycl::image_channel_type channel_type;
@ -703,7 +703,7 @@ static sycl::ext::oneapi::experimental::image_descriptor image_desc(const device
return param;
}
void OneapiDevice::tex_alloc(device_texture &mem)
void OneapiDevice::image_alloc(device_image &mem)
{
assert(device_queue_);
@ -765,6 +765,7 @@ void OneapiDevice::tex_alloc(device_texture &mem)
sycl::ext::oneapi::experimental::image_descriptor desc{};
if (mem.data_height > 0) {
/* 2D/3D image -- Tile optimized */
const sycl::device &device = reinterpret_cast<sycl::queue *>(queue)->get_device();
const size_t max_width = device.get_info<sycl::info::device::image2d_max_width>();
const size_t max_height = device.get_info<sycl::info::device::image2d_max_height>();
@ -790,11 +791,11 @@ void OneapiDevice::tex_alloc(device_texture &mem)
sycl::ext::oneapi::experimental::image_mem_handle memHandle =
sycl::ext::oneapi::experimental::alloc_image_mem(desc, *queue);
if (!memHandle.raw_handle) {
set_error("GPU texture allocation failed: Raw handle is null");
set_error("GPU image allocation failed: Raw handle is null");
return;
}
/* Copy data from host to the texture properly based on the texture description */
/* Copy data from host to the image properly based on the image description */
queue->ext_oneapi_copy(mem.host_pointer, memHandle, desc);
mem.device_pointer = (device_ptr)memHandle.raw_handle;
@ -807,7 +808,7 @@ void OneapiDevice::tex_alloc(device_texture &mem)
cmem->array = (arrayMemObject)(memHandle.raw_handle);
}
else {
/* 1D texture -- Linear memory */
/* 1D image -- Linear memory */
desc = sycl::ext::oneapi::experimental::image_descriptor(
{mem.data_width}, mem.data_elements, channel_type);
cmem = generic_alloc(mem);
@ -821,7 +822,7 @@ void OneapiDevice::tex_alloc(device_texture &mem)
queue->wait_and_throw();
/* Set Mapping and tag that we need to (re-)upload to device */
TextureInfo tex_info = mem.info;
KernelImageInfo tex_info = mem.info;
sycl::ext::oneapi::experimental::bindless_image_sampler samp(
address_mode, sycl::coordinate_normalization_mode::normalized, filter_mode);
@ -830,11 +831,11 @@ void OneapiDevice::tex_alloc(device_texture &mem)
sycl::ext::oneapi::experimental::sampled_image_handle imgHandle;
if (memHandle.raw_handle) {
/* Create 2D/3D texture handle */
/* Create 2D/3D image handle */
imgHandle = sycl::ext::oneapi::experimental::create_image(memHandle, samp, desc, *queue);
}
else {
/* Create 1D texture */
/* Create 1D image */
imgHandle = sycl::ext::oneapi::experimental::create_image(
(void *)mem.device_pointer, 0, samp, desc, *queue);
}
@ -850,36 +851,36 @@ void OneapiDevice::tex_alloc(device_texture &mem)
}
{
/* Update texture info. */
thread_scoped_lock lock(texture_info_mutex);
/* Update image info. */
thread_scoped_lock lock(image_info_mutex);
const uint slot = mem.slot;
if (slot >= texture_info.size()) {
if (slot >= image_info.size()) {
/* Allocate some slots in advance, to reduce amount of re-allocations. */
texture_info.resize(slot + 128);
image_info.resize(slot + 128);
}
texture_info[slot] = tex_info;
need_texture_info = true;
image_info[slot] = tex_info;
need_image_info = true;
}
}
catch (sycl::exception const &e) {
set_error("GPU texture allocation failed: runtime exception \"" + string(e.what()) + "\"");
set_error("GPU image allocation failed: runtime exception \"" + string(e.what()) + "\"");
}
}
void OneapiDevice::tex_copy_to(device_texture &mem)
void OneapiDevice::image_copy_to(device_image &mem)
{
if (!mem.device_pointer) {
tex_alloc(mem);
image_alloc(mem);
}
else {
if (mem.data_height > 0) {
/* 2D/3D texture -- Tile optimized */
/* 2D/3D image -- Tile optimized */
sycl::ext::oneapi::experimental::image_descriptor desc = image_desc(mem);
sycl::queue *queue = reinterpret_cast<sycl::queue *>(device_queue_);
try {
/* Copy data from host to the texture properly based on the texture description */
/* Copy data from host to the image properly based on the image description */
thread_scoped_lock lock(device_mem_map_mutex);
const Mem &cmem = device_mem_map[&mem];
sycl::ext::oneapi::experimental::image_mem_handle image_handle{
@ -891,7 +892,7 @@ void OneapiDevice::tex_copy_to(device_texture &mem)
# endif
}
catch (sycl::exception const &e) {
set_error("oneAPI texture copy error: got runtime exception \"" + string(e.what()) + "\"");
set_error("oneAPI image copy error: got runtime exception \"" + string(e.what()) + "\"");
}
}
else {
@ -900,7 +901,7 @@ void OneapiDevice::tex_copy_to(device_texture &mem)
}
}
void OneapiDevice::tex_free(device_texture &mem)
void OneapiDevice::image_free(device_image &mem)
{
if (mem.device_pointer) {
thread_scoped_lock lock(device_mem_map_mutex);
@ -910,24 +911,24 @@ void OneapiDevice::tex_free(device_texture &mem)
sycl::queue *queue = reinterpret_cast<sycl::queue *>(device_queue_);
if (cmem.texobject) {
/* Free bindless texture itself. */
/* Free bindless image itself. */
sycl::ext::oneapi::experimental::sampled_image_handle image(cmem.texobject);
sycl::ext::oneapi::experimental::destroy_image_handle(image, *queue);
}
if (cmem.array) {
/* Free texture memory. */
/* Free image memory. */
sycl::ext::oneapi::experimental::image_mem_handle imgHandle{
(sycl::ext::oneapi::experimental::image_mem_handle::raw_handle_type)cmem.array};
try {
/* We have allocated only standard textures, so we also deallocate only them. */
/* We have allocated only standard image, so we also deallocate only them. */
sycl::ext::oneapi::experimental::free_image_mem(
imgHandle, sycl::ext::oneapi::experimental::image_type::standard, *queue);
}
catch (sycl::exception const &e) {
set_error("oneAPI texture deallocation error: got runtime exception \"" +
string(e.what()) + "\"");
set_error("oneAPI image deallocation error: got runtime exception \"" + string(e.what()) +
"\"");
}
stats.mem_free(mem.memory_size());

8
intern/cycles/device/oneapi/device_impl.h
View file

@ -96,10 +96,10 @@ class OneapiDevice : public GPUDevice {
void global_copy_to(device_memory &mem);
void global_free(device_memory &mem);
/* Texture memory. */
void tex_alloc(device_texture &mem);
void tex_copy_to(device_texture &mem);
void tex_free(device_texture &mem);
/* Image memory. */
void image_alloc(device_image &mem);
void image_copy_to(device_image &mem);
void image_free(device_image &mem);
/* Host side memory, override for more efficient copies. */
void *host_alloc(const MemoryType type, const size_t size) override;

6
intern/cycles/device/oneapi/queue.cpp
View file

@ -55,7 +55,7 @@ int OneapiDeviceQueue::num_sort_partitions(int max_num_paths, uint /*max_scene_s
void OneapiDeviceQueue::init_execution()
{
oneapi_device_->load_texture_info();
oneapi_device_->load_image_info();
SyclQueue *device_queue = oneapi_device_->sycl_queue();
void *kg_dptr = oneapi_device_->kernel_globals_device_pointer();
@ -76,8 +76,8 @@ bool OneapiDeviceQueue::enqueue(DeviceKernel kernel,
return false;
}
/* Update texture info in case memory moved to host. */
if (oneapi_device_->load_texture_info()) {
/* Update image info in case memory moved to host. */
if (oneapi_device_->load_image_info()) {
if (!synchronize()) {
return false;
}

2
intern/cycles/device/optix/device_impl.cpp
View file

@ -105,7 +105,7 @@ OptiXDevice::~OptiXDevice()
free_bvh_memory_delayed();
sbt_data.free();
texture_info.free();
image_info.free();
launch_params.free();
/* Unload modules. */

4
intern/cycles/kernel/CMakeLists.txt
View file

@ -235,9 +235,9 @@ set(SRC_KERNEL_UTIL_HEADERS
util/colorspace.h
util/differential.h
util/ies.h
util/image_3d.h
util/lookup_table.h
util/nanovdb.h
util/texture_3d.h
util/profiler.h
)
@ -292,13 +292,13 @@ set(SRC_UTIL_HEADERS
../util/rect.h
../util/static_assert.h
../util/transform.h
../util/texture.h
../util/types.h
../util/types_base.h
../util/types_float2.h
../util/types_float3.h
../util/types_float4.h
../util/types_float8.h
../util/types_image.h
../util/types_int2.h
../util/types_int3.h
../util/types_int4.h

2
intern/cycles/kernel/data_arrays.h
View file

@ -79,7 +79,7 @@ KERNEL_DATA_ARRAY(float, lookup_table)
KERNEL_DATA_ARRAY(float, sample_pattern_lut)
/* image textures */
KERNEL_DATA_ARRAY(TextureInfo, texture_info)
KERNEL_DATA_ARRAY(KernelImageInfo, image_info)
/* ies lights */
KERNEL_DATA_ARRAY(float, ies)

4
intern/cycles/kernel/device/cpu/globals.h
View file

@ -13,8 +13,8 @@
# include "kernel/osl/globals.h"
#endif
#include "util/guiding.h" // IWYU pragma: keep
#include "util/texture.h" // IWYU pragma: keep
#include "util/guiding.h" // IWYU pragma: keep
#include "util/types_image.h" // IWYU pragma: keep
#include "util/unique_ptr.h"
CCL_NAMESPACE_BEGIN

34
intern/cycles/kernel/device/cpu/image.h
View file

@ -8,6 +8,7 @@
#include "kernel/device/cpu/globals.h"
#include "util/half.h"
#include "util/types_image.h"
CCL_NAMESPACE_BEGIN
@ -31,7 +32,7 @@ ccl_device_inline float frac(const float x, int *ix)
return x - (float)i;
}
template<typename TexT, typename OutT = float4> struct TextureInterpolator {
template<typename TexT, typename OutT = float4> struct ImageInterpolator {
static ccl_always_inline OutT zero()
{
@ -129,7 +130,7 @@ template<typename TexT, typename OutT = float4> struct TextureInterpolator {
/* ******** 2D interpolation ******** */
static ccl_always_inline OutT interp_closest(const TextureInfo &info, const float x, float y)
static ccl_always_inline OutT interp_closest(const KernelImageInfo &info, const float x, float y)
{
const int width = info.width;
const int height = info.height;
@ -164,7 +165,7 @@ template<typename TexT, typename OutT = float4> struct TextureInterpolator {
return read(data, ix, iy, width, height);
}
static ccl_always_inline OutT interp_linear(const TextureInfo &info, const float x, float y)
static ccl_always_inline OutT interp_linear(const KernelImageInfo &info, const float x, float y)
{
const int width = info.width;
const int height = info.height;
@ -218,7 +219,7 @@ template<typename TexT, typename OutT = float4> struct TextureInterpolator {
ty * tx * read(data, nix, niy, width, height);
}
static ccl_always_inline OutT interp_cubic(const TextureInfo &info, const float x, float y)
static ccl_always_inline OutT interp_cubic(const KernelImageInfo &info, const float x, float y)
{
const int width = info.width;
const int height = info.height;
@ -307,7 +308,7 @@ template<typename TexT, typename OutT = float4> struct TextureInterpolator {
#undef DATA
}
static ccl_always_inline OutT interp(const TextureInfo &info, const float x, float y)
static ccl_always_inline OutT interp(const KernelImageInfo &info, const float x, float y)
{
switch (info.interpolation) {
case INTERPOLATION_CLOSEST:
@ -322,9 +323,9 @@ template<typename TexT, typename OutT = float4> struct TextureInterpolator {
#undef SET_CUBIC_SPLINE_WEIGHTS
ccl_device float4 kernel_tex_image_interp(KernelGlobals kg, const int id, const float x, float y)
ccl_device float4 kernel_image_interp(KernelGlobals kg, const int id, const float x, float y)
{
const TextureInfo &info = kernel_data_fetch(texture_info, id);
const KernelImageInfo &info = kernel_data_fetch(image_info, id);
if (UNLIKELY(!info.data)) {
return zero_float4();
@ -332,33 +333,32 @@ ccl_device float4 kernel_tex_image_interp(KernelGlobals kg, const int id, const
switch (info.data_type) {
case IMAGE_DATA_TYPE_HALF: {
const float f = TextureInterpolator<half, float>::interp(info, x, y);
const float f = ImageInterpolator<half, float>::interp(info, x, y);
return make_float4(f, f, f, 1.0f);
}
case IMAGE_DATA_TYPE_BYTE: {
const float f = TextureInterpolator<uchar, float>::interp(info, x, y);
const float f = ImageInterpolator<uchar, float>::interp(info, x, y);
return make_float4(f, f, f, 1.0f);
}
case IMAGE_DATA_TYPE_USHORT: {
const float f = TextureInterpolator<uint16_t, float>::interp(info, x, y);
const float f = ImageInterpolator<uint16_t, float>::interp(info, x, y);
return make_float4(f, f, f, 1.0f);
}
case IMAGE_DATA_TYPE_FLOAT: {
const float f = TextureInterpolator<float, float>::interp(info, x, y);
const float f = ImageInterpolator<float, float>::interp(info, x, y);
return make_float4(f, f, f, 1.0f);
}
case IMAGE_DATA_TYPE_HALF4:
return TextureInterpolator<half4>::interp(info, x, y);
return ImageInterpolator<half4>::interp(info, x, y);
case IMAGE_DATA_TYPE_BYTE4:
return TextureInterpolator<uchar4>::interp(info, x, y);
return ImageInterpolator<uchar4>::interp(info, x, y);
case IMAGE_DATA_TYPE_USHORT4:
return TextureInterpolator<ushort4>::interp(info, x, y);
return ImageInterpolator<ushort4>::interp(info, x, y);
case IMAGE_DATA_TYPE_FLOAT4:
return TextureInterpolator<float4>::interp(info, x, y);
return ImageInterpolator<float4>::interp(info, x, y);
default:
assert(0);
return make_float4(
TEX_IMAGE_MISSING_R, TEX_IMAGE_MISSING_G, TEX_IMAGE_MISSING_B, TEX_IMAGE_MISSING_A);
return IMAGE_MISSING_RGBA;
}
}

8
intern/cycles/kernel/device/cuda/compat.h
View file

@ -75,12 +75,12 @@ typedef unsigned long long uint64_t;
/* GPU texture objects */
typedef unsigned long long CUtexObject;
typedef CUtexObject ccl_gpu_tex_object_2D;
typedef CUtexObject ccl_gpu_image_object_2D;
template<typename T>
ccl_device_forceinline T ccl_gpu_tex_object_read_2D(const ccl_gpu_tex_object_2D texobj,
const float x,
const float y)
ccl_device_forceinline T ccl_gpu_image_object_read_2D(const ccl_gpu_image_object_2D texobj,
const float x,
const float y)
{
return tex2D<T>(texobj, x, y);
}

2
intern/cycles/kernel/device/cuda/globals.h
View file

@ -12,7 +12,7 @@
#include "kernel/util/profiler.h"
#include "util/color.h"
#include "util/texture.h"
#include "util/types_image.h"
CCL_NAMESPACE_BEGIN

38
intern/cycles/kernel/device/gpu/image.h
View file

@ -55,11 +55,11 @@ ccl_device float cubic_h1(const float a)
/* Fast bicubic texture lookup using 4 bilinear lookups, adapted from CUDA samples. */
template<typename T>
ccl_device_noinline T kernel_tex_image_interp_bicubic(const ccl_global TextureInfo &info,
float x,
float y)
ccl_device_noinline T kernel_image_interp_bicubic(const ccl_global KernelImageInfo &info,
float x,
float y)
{
ccl_gpu_tex_object_2D tex = (ccl_gpu_tex_object_2D)info.data;
ccl_gpu_image_object_2D tex = (ccl_gpu_image_object_2D)info.data;
x = (x * info.width) - 0.5f;
y = (y * info.height) - 0.5f;
@ -77,27 +77,27 @@ ccl_device_noinline T kernel_tex_image_interp_bicubic(const ccl_global TextureIn
float y0 = (py + cubic_h0(fy) + 0.5f) / info.height;
float y1 = (py + cubic_h1(fy) + 0.5f) / info.height;
return cubic_g0(fy) * (g0x * ccl_gpu_tex_object_read_2D<T>(tex, x0, y0) +
g1x * ccl_gpu_tex_object_read_2D<T>(tex, x1, y0)) +
cubic_g1(fy) * (g0x * ccl_gpu_tex_object_read_2D<T>(tex, x0, y1) +
g1x * ccl_gpu_tex_object_read_2D<T>(tex, x1, y1));
return cubic_g0(fy) * (g0x * ccl_gpu_image_object_read_2D<T>(tex, x0, y0) +
g1x * ccl_gpu_image_object_read_2D<T>(tex, x1, y0)) +
cubic_g1(fy) * (g0x * ccl_gpu_image_object_read_2D<T>(tex, x0, y1) +
g1x * ccl_gpu_image_object_read_2D<T>(tex, x1, y1));
}
ccl_device float4 kernel_tex_image_interp(KernelGlobals kg, const int id, const float x, float y)
ccl_device float4 kernel_image_interp(KernelGlobals kg, const int id, const float x, float y)
{
const ccl_global TextureInfo &info = kernel_data_fetch(texture_info, id);
const ccl_global KernelImageInfo &info = kernel_data_fetch(image_info, id);
/* float4, byte4, ushort4 and half4 */
const int texture_type = info.data_type;
if (texture_type == IMAGE_DATA_TYPE_FLOAT4 || texture_type == IMAGE_DATA_TYPE_BYTE4 ||
texture_type == IMAGE_DATA_TYPE_HALF4 || texture_type == IMAGE_DATA_TYPE_USHORT4)
const int image_type = info.data_type;
if (image_type == IMAGE_DATA_TYPE_FLOAT4 || image_type == IMAGE_DATA_TYPE_BYTE4 ||
image_type == IMAGE_DATA_TYPE_HALF4 || image_type == IMAGE_DATA_TYPE_USHORT4)
{
if (info.interpolation == INTERPOLATION_CUBIC || info.interpolation == INTERPOLATION_SMART) {
return kernel_tex_image_interp_bicubic<float4>(info, x, y);
return kernel_image_interp_bicubic<float4>(info, x, y);
}
else {
ccl_gpu_tex_object_2D tex = (ccl_gpu_tex_object_2D)info.data;
return ccl_gpu_tex_object_read_2D<float4>(tex, x, y);
ccl_gpu_image_object_2D tex = (ccl_gpu_image_object_2D)info.data;
return ccl_gpu_image_object_read_2D<float4>(tex, x, y);
}
}
/* float, byte and half */
@ -105,11 +105,11 @@ ccl_device float4 kernel_tex_image_interp(KernelGlobals kg, const int id, const
float f;
if (info.interpolation == INTERPOLATION_CUBIC || info.interpolation == INTERPOLATION_SMART) {
f = kernel_tex_image_interp_bicubic<float>(info, x, y);
f = kernel_image_interp_bicubic<float>(info, x, y);
}
else {
ccl_gpu_tex_object_2D tex = (ccl_gpu_tex_object_2D)info.data;
f = ccl_gpu_tex_object_read_2D<float>(tex, x, y);
ccl_gpu_image_object_2D tex = (ccl_gpu_image_object_2D)info.data;
f = ccl_gpu_image_object_read_2D<float>(tex, x, y);
}
return make_float4(f, f, f, 1.0f);

8
intern/cycles/kernel/device/hip/compat.h
View file

@ -77,12 +77,12 @@ typedef unsigned long long uint64_t;
#define ccl_gpu_ballot(predicate) __ballot(predicate)
/* GPU texture objects */
typedef hipTextureObject_t ccl_gpu_tex_object_2D;
typedef hipTextureObject_t ccl_gpu_image_object_2D;
template<typename T>
ccl_device_forceinline T ccl_gpu_tex_object_read_2D(const ccl_gpu_tex_object_2D texobj,
const float x,
const float y)
ccl_device_forceinline T ccl_gpu_image_object_read_2D(const ccl_gpu_image_object_2D texobj,
const float x,
const float y)
{
return tex2D<T>(texobj, x, y);
}

2
intern/cycles/kernel/device/hip/globals.h
View file

@ -12,7 +12,7 @@
#include "kernel/util/profiler.h"
#include "util/color.h"
#include "util/texture.h"
#include "util/types_image.h"
CCL_NAMESPACE_BEGIN

4
intern/cycles/kernel/device/hiprt/globals.h
View file

@ -11,8 +11,8 @@
#include "kernel/integrator/state.h"
#include "kernel/util/profiler.h" // IWYU pragma: export
#include "util/color.h" // IWYU pragma: export
#include "util/texture.h" // IWYU pragma: export
#include "util/color.h" // IWYU pragma: export
#include "util/types_image.h" // IWYU pragma: export
/* The size of global stack available to each thread (memory reserved for each thread in
* global_stack_buffer). */

8
intern/cycles/kernel/device/metal/context_begin.h
View file

@ -24,11 +24,11 @@ class MetalKernelContext {
{}
/* texture fetch adapter functions */
using ccl_gpu_tex_object_2D = uint64_t;
using ccl_gpu_image_object_2D = uint64_t;
template<typename T>
inline __attribute__((__always_inline__))
T ccl_gpu_tex_object_read_2D(ccl_gpu_tex_object_2D tex, const float x, float y) const {
T ccl_gpu_image_object_read_2D(ccl_gpu_image_object_2D tex, const float x, float y) const {
kernel_assert(0);
return 0;
}
@ -36,14 +36,14 @@ class MetalKernelContext {
// texture2d
template<>
inline __attribute__((__always_inline__))
float4 ccl_gpu_tex_object_read_2D(ccl_gpu_tex_object_2D tex, const float x, float y) const {
float4 ccl_gpu_image_object_read_2D(ccl_gpu_image_object_2D tex, const float x, float y) const {
const uint tid(tex);
const uint sid(tex >> 32);
return ((ccl_global Texture2DParamsMetal*)metal_ancillaries->textures)[tid].tex.sample(metal_samplers[sid], float2(x, y));
}
template<>
inline __attribute__((__always_inline__))
float ccl_gpu_tex_object_read_2D(ccl_gpu_tex_object_2D tex, const float x, float y) const {
float ccl_gpu_image_object_read_2D(ccl_gpu_image_object_2D tex, const float x, float y) const {
const uint tid(tex);
const uint sid(tex >> 32);
return ((ccl_global Texture2DParamsMetal*)metal_ancillaries->textures)[tid].tex.sample(metal_samplers[sid], float2(x, y)).x;

2
intern/cycles/kernel/device/metal/globals.h
View file

@ -10,7 +10,7 @@
#include "kernel/util/profiler.h"
#include "util/color.h"
#include "util/texture.h"
#include "util/types_image.h"
CCL_NAMESPACE_BEGIN

37
intern/cycles/kernel/device/oneapi/compat.h
View file

@ -243,13 +243,13 @@ ccl_device_forceinline int __float_as_int(const float x)
static_assert(
sizeof(sycl::ext::oneapi::experimental::sampled_image_handle::raw_image_handle_type) ==
sizeof(uint64_t));
typedef uint64_t ccl_gpu_tex_object_2D;
typedef uint64_t ccl_gpu_tex_object_3D;
typedef uint64_t ccl_gpu_image_object_2D;
typedef uint64_t ccl_gpu_image_object_3D;
template<typename T>
ccl_device_forceinline T ccl_gpu_tex_object_read_2D(const ccl_gpu_tex_object_2D texobj,
const float x,
const float y)
ccl_device_forceinline T ccl_gpu_image_object_read_2D(const ccl_gpu_image_object_2D texobj,
const float x,
const float y)
{
/* Generic implementation not possible due to limitation with SYCL bindless sampled images
* not being able to read in a format, which is different from the supported data type of
@ -260,9 +260,8 @@ ccl_device_forceinline T ccl_gpu_tex_object_read_2D(const ccl_gpu_tex_object_2D
}
template<>
ccl_device_forceinline float ccl_gpu_tex_object_read_2D<float>(const ccl_gpu_tex_object_2D texobj,
const float x,
const float y)
ccl_device_forceinline float ccl_gpu_image_object_read_2D<float>(
const ccl_gpu_image_object_2D texobj, const float x, const float y)
{
sycl::ext::oneapi::experimental::sampled_image_handle image(
(sycl::ext::oneapi::experimental::sampled_image_handle::raw_image_handle_type)texobj);
@ -270,8 +269,8 @@ ccl_device_forceinline float ccl_gpu_tex_object_read_2D<float>(const ccl_gpu_tex
}
template<>
ccl_device_forceinline float4 ccl_gpu_tex_object_read_2D<float4>(
const ccl_gpu_tex_object_2D texobj, const float x, const float y)
ccl_device_forceinline float4 ccl_gpu_image_object_read_2D<float4>(
const ccl_gpu_image_object_2D texobj, const float x, const float y)
{
sycl::ext::oneapi::experimental::sampled_image_handle image(
(sycl::ext::oneapi::experimental::sampled_image_handle::raw_image_handle_type)texobj);
@ -280,10 +279,10 @@ ccl_device_forceinline float4 ccl_gpu_tex_object_read_2D<float4>(
}
template<typename T>
ccl_device_forceinline T ccl_gpu_tex_object_read_3D(const ccl_gpu_tex_object_3D texobj,
const float x,
const float y,
const float z)
ccl_device_forceinline T ccl_gpu_image_object_read_3D(const ccl_gpu_image_object_3D texobj,
const float x,
const float y,
const float z)
{
/* A generic implementation is not possible due to limitations with SYCL bindless sampled images
* not being able to read in a format that is different from the supported data type of
@ -295,10 +294,8 @@ ccl_device_forceinline T ccl_gpu_tex_object_read_3D(const ccl_gpu_tex_object_3D
}
template<>
ccl_device_forceinline float ccl_gpu_tex_object_read_3D<float>(const ccl_gpu_tex_object_3D texobj,
const float x,
const float y,
const float z)
ccl_device_forceinline float ccl_gpu_image_object_read_3D<float>(
const ccl_gpu_image_object_3D texobj, const float x, const float y, const float z)
{
sycl::ext::oneapi::experimental::sampled_image_handle image(
(sycl::ext::oneapi::experimental::sampled_image_handle::raw_image_handle_type)texobj);
@ -306,8 +303,8 @@ ccl_device_forceinline float ccl_gpu_tex_object_read_3D<float>(const ccl_gpu_tex
}
template<>
ccl_device_forceinline float4 ccl_gpu_tex_object_read_3D<float4>(
const ccl_gpu_tex_object_3D texobj, const float x, const float y, const float z)
ccl_device_forceinline float4 ccl_gpu_image_object_read_3D<float4>(
const ccl_gpu_image_object_3D texobj, const float x, const float y, const float z)
{
sycl::ext::oneapi::experimental::sampled_image_handle image(
(sycl::ext::oneapi::experimental::sampled_image_handle::raw_image_handle_type)texobj);

2
intern/cycles/kernel/device/oneapi/globals.h
View file

@ -12,7 +12,7 @@
#include "kernel/util/profiler.h"
#include "util/color.h"
#include "util/texture.h"
#include "util/types_image.h"
CCL_NAMESPACE_BEGIN

8
intern/cycles/kernel/device/optix/compat.h
View file

@ -62,12 +62,12 @@ typedef unsigned long long uint64_t;
/* GPU texture objects */
typedef unsigned long long CUtexObject;
typedef CUtexObject ccl_gpu_tex_object_2D;
typedef CUtexObject ccl_gpu_image_object_2D;
template<typename T>
ccl_device_forceinline T ccl_gpu_tex_object_read_2D(const ccl_gpu_tex_object_2D texobj,
const float x,
const float y)
ccl_device_forceinline T ccl_gpu_image_object_read_2D(const ccl_gpu_image_object_2D texobj,
const float x,
const float y)
{
return tex2D<T>(texobj, x, y);
}

2
intern/cycles/kernel/device/optix/globals.h
View file

@ -12,7 +12,7 @@
#include "kernel/util/profiler.h"
#include "util/color.h"
#include "util/texture.h"
#include "util/types_image.h"
CCL_NAMESPACE_BEGIN

8
intern/cycles/kernel/geom/volume.h
View file

@ -18,9 +18,7 @@
#include "kernel/geom/attribute.h"
#include "kernel/geom/object.h"
#include "kernel/sample/lcg.h"
#include "kernel/util/texture_3d.h"
#include "kernel/util/image_3d.h"
CCL_NAMESPACE_BEGIN
@ -88,13 +86,13 @@ ccl_device float4 volume_attribute_float4(KernelGlobals kg,
}
if (desc.element == ATTR_ELEMENT_VOXEL) {
/* todo: optimize this so we don't have to transform both here and in
* kernel_tex_image_interp_3d when possible. Also could optimize for the
* kernel_image_interp_3d when possible. Also could optimize for the
* common case where transform is translation/scale only. */
float3 P = sd->P;
object_inverse_position_transform(kg, sd, &P);
const InterpolationType interp = (sd->flag & SD_VOLUME_CUBIC) ? INTERPOLATION_CUBIC :
INTERPOLATION_NONE;
return kernel_tex_image_interp_3d(kg, sd, desc.offset, P, interp, stochastic);
return kernel_image_interp_3d(kg, sd, desc.offset, P, interp, stochastic);
}
return zero_float4();
}

10
intern/cycles/kernel/osl/services.cpp
View file

@ -6,6 +6,7 @@
* here, so for now we just put here. In the future it might be better
* to have dedicated file for such tweaks.
*/
#include "util/types_image.h"
#if (defined(__GNUC__) && !defined(__clang__)) && defined(NDEBUG)
# pragma GCC diagnostic ignored "-Wmaybe-uninitialized"
# pragma GCC diagnostic ignored "-Wuninitialized"
@ -40,7 +41,7 @@
#include "kernel/svm/bevel.h"
#include "kernel/util/ies.h"
#include "kernel/util/texture_3d.h"
#include "kernel/util/image_3d.h"
CCL_NAMESPACE_BEGIN
@ -1168,11 +1169,10 @@ bool OSLRenderServices::texture(OSLUStringHash filename,
float4 rgba;
if (id == -1) {
rgba = make_float4(
TEX_IMAGE_MISSING_R, TEX_IMAGE_MISSING_G, TEX_IMAGE_MISSING_B, TEX_IMAGE_MISSING_A);
rgba = IMAGE_MISSING_RGBA;
}
else {
rgba = kernel_tex_image_interp(kernel_globals, id, s, 1.0f - t);
rgba = kernel_image_interp(kernel_globals, id, s, 1.0f - t);
}
result[0] = rgba[0];
@ -1286,7 +1286,7 @@ bool OSLRenderServices::texture3d(OSLUStringHash filename,
/* Packed texture. */
const int slot = handle->svm_slots[0].y;
const float3 P_float3 = make_float3(P.x, P.y, P.z);
float4 rgba = kernel_tex_image_interp_3d(
float4 rgba = kernel_image_interp_3d(
kernel_globals, globals->sd, slot, P_float3, INTERPOLATION_NONE, false);
result[0] = rgba[0];

6
intern/cycles/kernel/osl/services_gpu.h
View file

@ -17,7 +17,7 @@
#include "kernel/util/differential.h"
#include "kernel/util/ies.h"
#include "kernel/util/texture_3d.h"
#include "kernel/util/image_3d.h"
#include "util/hash.h"
#include "util/transform.h"
@ -1024,7 +1024,7 @@ ccl_device_extern bool rs_texture(ccl_private ShaderGlobals *sg,
switch (type) {
case OSL_TEXTURE_HANDLE_TYPE_SVM: {
const float4 rgba = kernel_tex_image_interp(nullptr, slot, s, 1.0f - t);
const float4 rgba = kernel_image_interp(nullptr, slot, s, 1.0f - t);
if (nchannels > 0) {
result[0] = rgba.x;
}
@ -1072,7 +1072,7 @@ ccl_device_extern bool rs_texture3d(ccl_private ShaderGlobals *sg,
switch (type) {
case OSL_TEXTURE_HANDLE_TYPE_SVM: {
const float4 rgba = kernel_tex_image_interp_3d(
const float4 rgba = kernel_image_interp_3d(
nullptr, sg->sd, slot, *P, INTERPOLATION_NONE, false);
if (nchannels > 0) {
result[0] = rgba.x;

8
intern/cycles/kernel/svm/image.h
View file

@ -14,6 +14,7 @@
#include "kernel/svm/util.h"
#include "util/color.h"
#include "util/types_image.h"
CCL_NAMESPACE_BEGIN
@ -21,11 +22,10 @@ ccl_device float4
svm_image_texture(KernelGlobals kg, const int id, const float x, float y, const uint flags)
{
if (id == -1) {
return make_float4(
TEX_IMAGE_MISSING_R, TEX_IMAGE_MISSING_G, TEX_IMAGE_MISSING_B, TEX_IMAGE_MISSING_A);
return IMAGE_MISSING_RGBA;
}
float4 r = kernel_tex_image_interp(kg, id, x, y);
float4 r = kernel_image_interp(kg, id, x, y);
const float alpha = r.w;
if ((flags & NODE_IMAGE_ALPHA_UNASSOCIATE) && alpha != 1.0f && alpha != 0.0f) {
@ -74,7 +74,7 @@ ccl_device_noinline int svm_node_tex_image(KernelGlobals kg,
}
/* TODO(lukas): Consider moving tile information out of the SVM node.
* TextureInfo seems a reasonable candidate. */
* KernelImageInfo seems a reasonable candidate. */
int id = -1;
const int num_nodes = (int)node.y;
if (num_nodes > 0) {

2
intern/cycles/kernel/svm/sky.h
View file

@ -172,7 +172,7 @@ ccl_device float3 sky_radiance_nishita(KernelGlobals kg,
const float x = fractf((-direction.y - M_PI_2_F + sun_rotation) * M_1_2PI_F);
/* Undo the non-linear transformation from the sky LUT */
const float y = copysignf(sqrtf(fabsf(dir_elevation) * M_2_PI_F), dir_elevation) * 0.5f + 0.5f;
xyz += make_float3(kernel_tex_image_interp(kg, texture_id, x, y));
xyz += make_float3(kernel_image_interp(kg, texture_id, x, y));
/* Convert to RGB */
return xyz_to_rgb_clamped(kg, xyz);

40
intern/cycles/kernel/util/texture_3d.h → intern/cycles/kernel/util/image_3d.h
View file

@ -7,7 +7,7 @@
#include "kernel/globals.h"
#include "kernel/sample/lcg.h"
#include "util/texture.h"
#include "util/types_image.h"
#if !defined(__KERNEL_METAL__) && !defined(__KERNEL_ONEAPI__)
# ifdef WITH_NANOVDB
@ -88,7 +88,7 @@ ccl_device_inline float3 interp_stochastic(const float3 P,
/** \} */
template<typename OutT, typename Acc>
ccl_device OutT kernel_tex_image_interp_trilinear_nanovdb(ccl_private Acc &acc, const float3 P)
ccl_device OutT kernel_image_interp_trilinear_nanovdb(ccl_private Acc &acc, const float3 P)
{
const float3 floor_P = floor(P);
const float3 t = P - floor_P;
@ -116,7 +116,7 @@ ccl_device OutT kernel_tex_image_interp_trilinear_nanovdb(ccl_private Acc &acc,
}
template<typename OutT, typename Acc>
ccl_device OutT kernel_tex_image_interp_tricubic_nanovdb(ccl_private Acc &acc, const float3 P)
ccl_device OutT kernel_image_interp_tricubic_nanovdb(ccl_private Acc &acc, const float3 P)
{
# if defined(__KERNEL_HIP__)
/* Explicitly unroll for HIP compiler to unroll the loop. Without this the render result is wrong
@ -180,7 +180,7 @@ __attribute__((noinline))
# else
ccl_device_noinline
# endif
OutT kernel_tex_image_interp_nanovdb(const ccl_global TextureInfo &info,
OutT kernel_image_interp_nanovdb(const ccl_global KernelImageInfo &info,
float3 P,
const InterpolationType interp)
{
@ -193,22 +193,22 @@ OutT kernel_tex_image_interp_nanovdb(const ccl_global TextureInfo &info,
nanovdb::CachedReadAccessor<T> acc(grid->tree().root());
if (interp == INTERPOLATION_LINEAR) {
return kernel_tex_image_interp_trilinear_nanovdb<OutT>(acc, P);
return kernel_image_interp_trilinear_nanovdb<OutT>(acc, P);
}
return kernel_tex_image_interp_tricubic_nanovdb<OutT>(acc, P);
return kernel_image_interp_tricubic_nanovdb<OutT>(acc, P);
}
#endif /* WITH_NANOVDB */
ccl_device float4 kernel_tex_image_interp_3d(KernelGlobals kg,
ccl_private ShaderData *sd,
const int id,
float3 P,
InterpolationType interp,
const bool stochastic)
ccl_device float4 kernel_image_interp_3d(KernelGlobals kg,
ccl_private ShaderData *sd,
const int id,
float3 P,
InterpolationType interp,
const bool stochastic)
{
#ifdef WITH_NANOVDB
const ccl_global TextureInfo &info = kernel_data_fetch(texture_info, id);
const ccl_global KernelImageInfo &info = kernel_data_fetch(image_info, id);
if (info.use_transform_3d) {
P = transform_point(&info.transform_3d, P);
@ -225,23 +225,22 @@ ccl_device float4 kernel_tex_image_interp_3d(KernelGlobals kg,
const ImageDataType data_type = (ImageDataType)info.data_type;
if (data_type == IMAGE_DATA_TYPE_NANOVDB_FLOAT) {
const float f = kernel_tex_image_interp_nanovdb<float, float>(info, P, interpolation);
const float f = kernel_image_interp_nanovdb<float, float>(info, P, interpolation);
return make_float4(f, f, f, 1.0f);
}
if (data_type == IMAGE_DATA_TYPE_NANOVDB_FLOAT3) {
const float3 f = kernel_tex_image_interp_nanovdb<float3, packed_float3>(
info, P, interpolation);
const float3 f = kernel_image_interp_nanovdb<float3, packed_float3>(info, P, interpolation);
return make_float4(f, 1.0f);
}
if (data_type == IMAGE_DATA_TYPE_NANOVDB_FLOAT4) {
return kernel_tex_image_interp_nanovdb<float4, float4>(info, P, interpolation);
return kernel_image_interp_nanovdb<float4, float4>(info, P, interpolation);
}
if (data_type == IMAGE_DATA_TYPE_NANOVDB_FPN) {
const float f = kernel_tex_image_interp_nanovdb<float, nanovdb::FpN>(info, P, interpolation);
const float f = kernel_image_interp_nanovdb<float, nanovdb::FpN>(info, P, interpolation);
return make_float4(f, f, f, 1.0f);
}
if (data_type == IMAGE_DATA_TYPE_NANOVDB_FP16) {
const float f = kernel_tex_image_interp_nanovdb<float, nanovdb::Fp16>(info, P, interpolation);
const float f = kernel_image_interp_nanovdb<float, nanovdb::Fp16>(info, P, interpolation);
return make_float4(f, f, f, 1.0f);
}
if (data_type == IMAGE_DATA_TYPE_NANOVDB_EMPTY) {
@ -256,8 +255,7 @@ ccl_device float4 kernel_tex_image_interp_3d(KernelGlobals kg,
(void)stochastic;
#endif
return make_float4(
TEX_IMAGE_MISSING_R, TEX_IMAGE_MISSING_G, TEX_IMAGE_MISSING_B, TEX_IMAGE_MISSING_A);
return IMAGE_MISSING_RGBA;
}
#ifndef __KERNEL_GPU__

46
intern/cycles/scene/image.cpp
View file

@ -16,7 +16,7 @@
#include "util/log.h"
#include "util/progress.h"
#include "util/task.h"
#include "util/texture.h"
#include "util/types_image.h"
#ifdef WITH_OSL
# include <OSL/oslexec.h>
@ -180,7 +180,7 @@ vector<int4> ImageHandle::get_svm_slots() const
return svm_slots;
}
device_texture *ImageHandle::image_memory() const
device_image *ImageHandle::image_memory() const
{
if (slots.empty()) {
return nullptr;
@ -565,7 +565,7 @@ void ImageManager::device_load_image(Device *device,
img->mem.reset();
}
img->mem = make_unique<device_texture>(
img->mem = make_unique<device_image>(
device, img->mem_name.c_str(), slot, type, img->params.interpolation, img->params.extension);
img->mem->info.use_transform_3d = img->metadata.use_transform_3d;
img->mem->info.transform_3d = img->metadata.transform_3d;
@ -577,10 +577,10 @@ void ImageManager::device_load_image(Device *device,
const thread_scoped_lock device_lock(device_mutex);
float *pixels = (float *)img->mem->alloc(1, 1);
pixels[0] = TEX_IMAGE_MISSING_R;
pixels[1] = TEX_IMAGE_MISSING_G;
pixels[2] = TEX_IMAGE_MISSING_B;
pixels[3] = TEX_IMAGE_MISSING_A;
pixels[0] = IMAGE_MISSING_RGBA.x;
pixels[1] = IMAGE_MISSING_RGBA.y;
pixels[2] = IMAGE_MISSING_RGBA.z;
pixels[3] = IMAGE_MISSING_RGBA.w;
}
}
else if (type == IMAGE_DATA_TYPE_FLOAT) {
@ -589,7 +589,7 @@ void ImageManager::device_load_image(Device *device,
const thread_scoped_lock device_lock(device_mutex);
float *pixels = (float *)img->mem->alloc(1, 1);
pixels[0] = TEX_IMAGE_MISSING_R;
pixels[0] = IMAGE_MISSING_RGBA.x;
}
}
else if (type == IMAGE_DATA_TYPE_BYTE4) {
@ -598,10 +598,10 @@ void ImageManager::device_load_image(Device *device,
const thread_scoped_lock device_lock(device_mutex);
uchar *pixels = (uchar *)img->mem->alloc(1, 1);
pixels[0] = (TEX_IMAGE_MISSING_R * 255);
pixels[1] = (TEX_IMAGE_MISSING_G * 255);
pixels[2] = (TEX_IMAGE_MISSING_B * 255);
pixels[3] = (TEX_IMAGE_MISSING_A * 255);
pixels[0] = (IMAGE_MISSING_RGBA.x * 255);
pixels[1] = (IMAGE_MISSING_RGBA.y * 255);
pixels[2] = (IMAGE_MISSING_RGBA.z * 255);
pixels[3] = (IMAGE_MISSING_RGBA.w * 255);
}
}
else if (type == IMAGE_DATA_TYPE_BYTE) {
@ -610,7 +610,7 @@ void ImageManager::device_load_image(Device *device,
const thread_scoped_lock device_lock(device_mutex);
uchar *pixels = (uchar *)img->mem->alloc(1, 1);
pixels[0] = (TEX_IMAGE_MISSING_R * 255);
pixels[0] = (IMAGE_MISSING_RGBA.x * 255);
}
}
else if (type == IMAGE_DATA_TYPE_HALF4) {
@ -619,10 +619,10 @@ void ImageManager::device_load_image(Device *device,
const thread_scoped_lock device_lock(device_mutex);
half *pixels = (half *)img->mem->alloc(1, 1);
pixels[0] = TEX_IMAGE_MISSING_R;
pixels[1] = TEX_IMAGE_MISSING_G;
pixels[2] = TEX_IMAGE_MISSING_B;
pixels[3] = TEX_IMAGE_MISSING_A;
pixels[0] = IMAGE_MISSING_RGBA.x;
pixels[1] = IMAGE_MISSING_RGBA.y;
pixels[2] = IMAGE_MISSING_RGBA.z;
pixels[3] = IMAGE_MISSING_RGBA.w;
}
}
else if (type == IMAGE_DATA_TYPE_USHORT) {
@ -631,7 +631,7 @@ void ImageManager::device_load_image(Device *device,
const thread_scoped_lock device_lock(device_mutex);
uint16_t *pixels = (uint16_t *)img->mem->alloc(1, 1);
pixels[0] = (TEX_IMAGE_MISSING_R * 65535);
pixels[0] = (IMAGE_MISSING_RGBA.x * 65535);
}
}
else if (type == IMAGE_DATA_TYPE_USHORT4) {
@ -640,10 +640,10 @@ void ImageManager::device_load_image(Device *device,
const thread_scoped_lock device_lock(device_mutex);
uint16_t *pixels = (uint16_t *)img->mem->alloc(1, 1);
pixels[0] = (TEX_IMAGE_MISSING_R * 65535);
pixels[1] = (TEX_IMAGE_MISSING_G * 65535);
pixels[2] = (TEX_IMAGE_MISSING_B * 65535);
pixels[3] = (TEX_IMAGE_MISSING_A * 65535);
pixels[0] = (IMAGE_MISSING_RGBA.x * 65535);
pixels[1] = (IMAGE_MISSING_RGBA.y * 65535);
pixels[2] = (IMAGE_MISSING_RGBA.z * 65535);
pixels[3] = (IMAGE_MISSING_RGBA.w * 65535);
}
}
else if (type == IMAGE_DATA_TYPE_HALF) {
@ -652,7 +652,7 @@ void ImageManager::device_load_image(Device *device,
const thread_scoped_lock device_lock(device_mutex);
half *pixels = (half *)img->mem->alloc(1, 1);
pixels[0] = TEX_IMAGE_MISSING_R;
pixels[0] = IMAGE_MISSING_RGBA.x;
}
}
#ifdef WITH_NANOVDB

4
intern/cycles/scene/image.h
View file

@ -104,7 +104,7 @@ class ImageHandle {
ImageMetaData metadata();
int svm_slot(const int slot_index = 0) const;
vector<int4> get_svm_slots() const;
device_texture *image_memory() const;
device_image *image_memory() const;
VDBImageLoader *vdb_loader() const;
@ -162,7 +162,7 @@ class ImageManager {
bool builtin;
string mem_name;
unique_ptr<device_texture> mem;
unique_ptr<device_image> mem;
int users;
thread_mutex mutex;

2
intern/cycles/scene/image_vdb.cpp
View file

@ -7,7 +7,7 @@
#include "util/log.h"
#include "util/nanovdb.h"
#include "util/openvdb.h"
#include "util/texture.h"
#include "util/types_image.h"
#ifdef WITH_OPENVDB
# include <openvdb/tools/Dense.h>

13
intern/cycles/scene/volume.cpp
View file

@ -22,7 +22,6 @@
#include "util/nanovdb.h"
#include "util/path.h"
#include "util/progress.h"
#include "util/texture.h"
#include "util/types.h"
#include "bvh/octree.h"
@ -658,17 +657,17 @@ void GeometryManager::create_volume_mesh(const Scene *scene, Volume *volume, Pro
continue;
}
/* Create NanoVDB grid handle from texture memory. */
device_texture *texture = handle.image_memory();
if (texture == nullptr || texture->host_pointer == nullptr ||
texture->info.data_type == IMAGE_DATA_TYPE_NANOVDB_EMPTY ||
!is_nanovdb_type(texture->info.data_type))
/* Create NanoVDB grid handle from image memory. */
device_image *image = handle.image_memory();
if (image == nullptr || image->host_pointer == nullptr ||
image->info.data_type == IMAGE_DATA_TYPE_NANOVDB_EMPTY ||
!is_nanovdb_type(image->info.data_type))
{
continue;
}
nanovdb::GridHandle grid(
nanovdb::HostBuffer::createFull(texture->memory_size(), texture->host_pointer));
nanovdb::HostBuffer::createFull(image->memory_size(), image->host_pointer));
/* Add padding based on the maximum velocity vector. */
if (attr.std == ATTR_STD_VOLUME_VELOCITY && scene->need_motion() != Scene::MOTION_NONE) {

2
intern/cycles/util/CMakeLists.txt
View file

@ -104,7 +104,6 @@ set(SRC_HEADERS
system.h
task.h
tbb.h
texture.h
thread.h
time.h
transform.h
@ -114,6 +113,7 @@ set(SRC_HEADERS
types_float3.h
types_float4.h
types_float8.h
types_image.h
types_int2.h
types_int3.h
types_int4.h

2
intern/cycles/util/image_metadata.cpp
View file

@ -11,7 +11,7 @@
#include "util/image_metadata.h"
#include "util/log.h"
#include "util/param.h"
#include "util/texture.h"
#include "util/types_image.h"
CCL_NAMESPACE_BEGIN

2
intern/cycles/util/image_metadata.h
View file

@ -10,7 +10,7 @@
#include "util/colorspace.h"
#include "util/string.h"
#include "util/texture.h"
#include "util/types_image.h"
CCL_NAMESPACE_BEGIN

20
intern/cycles/util/texture.h → intern/cycles/util/types_image.h
View file

@ -9,16 +9,10 @@
CCL_NAMESPACE_BEGIN
/* Color to use when textures are not found. */
enum {
TEX_IMAGE_MISSING_R = 1,
TEX_IMAGE_MISSING_G = 0,
TEX_IMAGE_MISSING_B = 1,
TEX_IMAGE_MISSING_A = 1
};
/* Color to use when images are not found. */
#define IMAGE_MISSING_RGBA make_float4(1, 0, 1, 1)
/* Interpolation types for textures
* CUDA also use texture space to store other objects. */
/* Interpolation types for images. */
enum InterpolationType {
INTERPOLATION_NONE = ~0,
INTERPOLATION_LINEAR = 0,
@ -29,6 +23,7 @@ enum InterpolationType {
INTERPOLATION_NUM_TYPES,
};
/* Image data types supported by the kernel. */
enum ImageDataType {
IMAGE_DATA_TYPE_FLOAT4 = 0,
IMAGE_DATA_TYPE_BYTE4 = 1,
@ -65,7 +60,7 @@ enum ImageAlphaType {
IMAGE_ALPHA_NUM_TYPES,
};
/* Extension types for textures.
/* Extension types for image.
*
* Defines how the image is extrapolated past its original bounds. */
enum ExtensionType {
@ -81,8 +76,9 @@ enum ExtensionType {
EXTENSION_NUM_TYPES,
};
struct TextureInfo {
/* Pointer, offset or texture depending on device. */
/* Kernel data structure to describe image. */
struct KernelImageInfo {
/* Pointer, offset or image/texture object depending on device. */
uint64_t data = 0;
/* Data Type */
uint data_type = IMAGE_DATA_NUM_TYPES;