Cycles: Adapt shared kernel/device/gpu layer for MSL
This patch adapts the shared kernel entrypoints so that they can be compiled as MSL (Metal Shading Language). Where possible, the adaptations avoid changes in common code. In MSL, kernel function inputs are explicitly bound to resources. In the case of argument buffers, we declare a struct containing the kernel arguments, accessible via device pointer. This differs from CUDA and HIP where kernel function arguments are declared as traditional C-style function parameters. This patch adapts the entrypoints declared in kernel.h so that they can be translated via a new `ccl_gpu_kernel_signature` macro into the required parameter struct + kernel entrypoint pairing for MSL. MSL buffer attribution must be applied to function parameters or non-static class data members. To allow universal access to the integrator state, kernel data, and texture fetch adapters, we wrap all of the shared kernel code in a `MetalKernelContext` class. This is achieved by bracketing the appropriate kernel headers with "context_begin.h" and "context_end.h" on Metal. When calling deeper into the kernel code, we must reference the context class (e.g. `context.integrator_init_from_camera`). This extra prefixing is performed by a set of defines in "context_end.h". These will require explicit maintenance if entrypoints change. We invite discussion on more maintainable ways to enforce correctness. Lambda expressions are not supported on MSL, so a new `ccl_gpu_kernel_lambda` macro generates an inline function object and optionally capturing any required state. This yields the same behaviour. This approach is applied to all parallel_... implementations which are templated by operation. The lambda expressions in the film_convert... kernels don't adapt cleanly to use function objects. However, these entrypoints can be macro-generated more concisely to avoid lambda expressions entirely, instead relying on constant folding to handle the pixel/channel conversions. A separate implementation of `gpu_parallel_active_index_array` is provided for Metal to workaround some subtle differences in SIMD width, and also to encapsulate some required thread parameters which must be declared as explicit entrypoint function parameters. Ref T92212 Reviewed By: brecht Maniphest Tasks: T92212 Differential Revision: https://developer.blender.org/D13109
This commit is contained in:
Michael Jones
Michael Jones
parent
4648c4990c
commit
3a4c8f406a
|
|
@ -39,6 +39,10 @@ set(SRC_KERNEL_DEVICE_HIP
|
|||
device/hip/kernel.cpp
|
||||
)
|
||||
|
||||
set(SRC_KERNEL_DEVICE_METAL
|
||||
device/metal/kernel.metal
|
||||
)
|
||||
|
||||
set(SRC_KERNEL_DEVICE_OPTIX
|
||||
device/optix/kernel.cu
|
||||
device/optix/kernel_shader_raytrace.cu
|
||||
|
|
@ -79,6 +83,13 @@ set(SRC_KERNEL_DEVICE_OPTIX_HEADERS
|
|||
device/optix/globals.h
|
||||
)
|
||||
|
||||
set(SRC_KERNEL_DEVICE_METAL_HEADERS
|
||||
device/metal/compat.h
|
||||
device/metal/context_begin.h
|
||||
device/metal/context_end.h
|
||||
device/metal/globals.h
|
||||
)
|
||||
|
||||
set(SRC_KERNEL_CLOSURE_HEADERS
|
||||
closure/alloc.h
|
||||
closure/bsdf.h
|
||||
|
|
@ -368,6 +379,7 @@ if(WITH_CYCLES_CUDA_BINARIES)
|
|||
${SRC_KERNEL_HEADERS}
|
||||
${SRC_KERNEL_DEVICE_GPU_HEADERS}
|
||||
${SRC_KERNEL_DEVICE_CUDA_HEADERS}
|
||||
${SRC_KERNEL_DEVICE_METAL_HEADERS}
|
||||
${SRC_UTIL_HEADERS}
|
||||
)
|
||||
set(cuda_cubins)
|
||||
|
|
@ -723,12 +735,14 @@ cycles_add_library(cycles_kernel "${LIB}"
|
|||
${SRC_KERNEL_DEVICE_CUDA}
|
||||
${SRC_KERNEL_DEVICE_HIP}
|
||||
${SRC_KERNEL_DEVICE_OPTIX}
|
||||
${SRC_KERNEL_DEVICE_METAL}
|
||||
${SRC_KERNEL_HEADERS}
|
||||
${SRC_KERNEL_DEVICE_CPU_HEADERS}
|
||||
${SRC_KERNEL_DEVICE_GPU_HEADERS}
|
||||
${SRC_KERNEL_DEVICE_CUDA_HEADERS}
|
||||
${SRC_KERNEL_DEVICE_HIP_HEADERS}
|
||||
${SRC_KERNEL_DEVICE_OPTIX_HEADERS}
|
||||
${SRC_KERNEL_DEVICE_METAL_HEADERS}
|
||||
)
|
||||
|
||||
source_group("bake" FILES ${SRC_KERNEL_BAKE_HEADERS})
|
||||
|
|
@ -740,6 +754,7 @@ source_group("device\\cuda" FILES ${SRC_KERNEL_DEVICE_CUDA} ${SRC_KERNEL_DEVICE_
|
|||
source_group("device\\gpu" FILES ${SRC_KERNEL_DEVICE_GPU_HEADERS})
|
||||
source_group("device\\hip" FILES ${SRC_KERNEL_DEVICE_HIP} ${SRC_KERNEL_DEVICE_HIP_HEADERS})
|
||||
source_group("device\\optix" FILES ${SRC_KERNEL_DEVICE_OPTIX} ${SRC_KERNEL_DEVICE_OPTIX_HEADERS})
|
||||
source_group("device\\metal" FILES ${SRC_KERNEL_DEVICE_METAL} ${SRC_KERNEL_DEVICE_METAL_HEADERS})
|
||||
source_group("film" FILES ${SRC_KERNEL_FILM_HEADERS})
|
||||
source_group("geom" FILES ${SRC_KERNEL_GEOM_HEADERS})
|
||||
source_group("integrator" FILES ${SRC_KERNEL_INTEGRATOR_HEADERS})
|
||||
|
|
@ -772,6 +787,8 @@ delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_DEVICE_HIP}" ${CYCLES_
|
|||
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_DEVICE_HIP_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/device/hip)
|
||||
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_DEVICE_OPTIX}" ${CYCLES_INSTALL_PATH}/source/kernel/device/optix)
|
||||
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_DEVICE_OPTIX_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/device/optix)
|
||||
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_DEVICE_METAL}" ${CYCLES_INSTALL_PATH}/source/kernel/device/metal)
|
||||
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_DEVICE_METAL_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/device/metal)
|
||||
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_FILM_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/film)
|
||||
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_GEOM_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/geom)
|
||||
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_INTEGRATOR_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/integrator)
|
||||
|
|
|
|||
|
|
@ -75,6 +75,7 @@ typedef unsigned long long uint64_t;
|
|||
#define ccl_gpu_block_idx_x (blockIdx.x)
|
||||
#define ccl_gpu_grid_dim_x (gridDim.x)
|
||||
#define ccl_gpu_warp_size (warpSize)
|
||||
#define ccl_gpu_thread_mask(thread_warp) uint(0xFFFFFFFF >> (ccl_gpu_warp_size - thread_warp))
|
||||
|
||||
#define ccl_gpu_global_id_x() (ccl_gpu_block_idx_x * ccl_gpu_block_dim_x + ccl_gpu_thread_idx_x)
|
||||
#define ccl_gpu_global_size_x() (ccl_gpu_grid_dim_x * ccl_gpu_block_dim_x)
|
||||
|
|
|
|||
|
|
@ -93,11 +93,35 @@
|
|||
/* Compute number of threads per block and minimum blocks per multiprocessor
|
||||
* given the maximum number of registers per thread. */
|
||||
|
||||
#define ccl_gpu_kernel(block_num_threads, thread_num_registers) \
|
||||
#define ccl_gpu_kernel_threads(block_num_threads) \
|
||||
extern "C" __global__ void __launch_bounds__(block_num_threads)
|
||||
|
||||
#define ccl_gpu_kernel_threads_registers(block_num_threads, thread_num_registers) \
|
||||
extern "C" __global__ void __launch_bounds__(block_num_threads, \
|
||||
GPU_MULTIPRESSOR_MAX_REGISTERS / \
|
||||
(block_num_threads * thread_num_registers))
|
||||
|
||||
/* allow ccl_gpu_kernel to accept 1 or 2 parameters */
|
||||
#define SELECT_MACRO(_1, _2, NAME, ...) NAME
|
||||
#define ccl_gpu_kernel(...) \
|
||||
SELECT_MACRO(__VA_ARGS__, ccl_gpu_kernel_threads_registers, ccl_gpu_kernel_threads)(__VA_ARGS__)
|
||||
|
||||
#define ccl_gpu_kernel_signature(name, ...) kernel_gpu_##name(__VA_ARGS__)
|
||||
|
||||
#define ccl_gpu_kernel_call(x) x
|
||||
|
||||
/* define a function object where "func" is the lambda body, and additional parameters are used to
|
||||
* specify captured state */
|
||||
#define ccl_gpu_kernel_lambda(func, ...) \
|
||||
struct KernelLambda { \
|
||||
__VA_ARGS__; \
|
||||
__device__ int operator()(const int state) \
|
||||
{ \
|
||||
return (func); \
|
||||
} \
|
||||
} ccl_gpu_kernel_lambda_pass; \
|
||||
ccl_gpu_kernel_lambda_pass
|
||||
|
||||
/* sanity checks */
|
||||
|
||||
#if GPU_KERNEL_BLOCK_NUM_THREADS > GPU_BLOCK_MAX_THREADS
|
||||
|
|
|
|||
|
|
@ -65,7 +65,9 @@ ccl_device float cubic_h1(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 TextureInfo &info, float x, float y)
|
||||
ccl_device_noinline T kernel_tex_image_interp_bicubic(ccl_global const TextureInfo &info,
|
||||
float x,
|
||||
float y)
|
||||
{
|
||||
ccl_gpu_tex_object tex = (ccl_gpu_tex_object)info.data;
|
||||
|
||||
|
|
@ -94,7 +96,7 @@ ccl_device_noinline T kernel_tex_image_interp_bicubic(const TextureInfo &info, f
|
|||
/* Fast tricubic texture lookup using 8 trilinear lookups. */
|
||||
template<typename T>
|
||||
ccl_device_noinline T
|
||||
kernel_tex_image_interp_tricubic(const TextureInfo &info, float x, float y, float z)
|
||||
kernel_tex_image_interp_tricubic(ccl_global const TextureInfo &info, float x, float y, float z)
|
||||
{
|
||||
ccl_gpu_tex_object tex = (ccl_gpu_tex_object)info.data;
|
||||
|
||||
|
|
@ -169,7 +171,7 @@ ccl_device T kernel_tex_image_interp_tricubic_nanovdb(S &s, float x, float y, fl
|
|||
|
||||
template<typename T>
|
||||
ccl_device_noinline T kernel_tex_image_interp_nanovdb(
|
||||
const TextureInfo &info, float x, float y, float z, uint interpolation)
|
||||
ccl_global const TextureInfo &info, float x, float y, float z, uint interpolation)
|
||||
{
|
||||
using namespace nanovdb;
|
||||
|
||||
|
|
@ -191,7 +193,7 @@ ccl_device_noinline T kernel_tex_image_interp_nanovdb(
|
|||
|
||||
ccl_device float4 kernel_tex_image_interp(KernelGlobals kg, int id, float x, float y)
|
||||
{
|
||||
const TextureInfo &info = kernel_tex_fetch(__texture_info, id);
|
||||
ccl_global const TextureInfo &info = kernel_tex_fetch(__texture_info, id);
|
||||
|
||||
/* float4, byte4, ushort4 and half4 */
|
||||
const int texture_type = info.data_type;
|
||||
|
|
@ -226,7 +228,7 @@ ccl_device float4 kernel_tex_image_interp_3d(KernelGlobals kg,
|
|||
float3 P,
|
||||
InterpolationType interp)
|
||||
{
|
||||
const TextureInfo &info = kernel_tex_fetch(__texture_info, id);
|
||||
ccl_global const TextureInfo &info = kernel_tex_fetch(__texture_info, id);
|
||||
|
||||
if (info.use_transform_3d) {
|
||||
P = transform_point(&info.transform_3d, P);
|
||||
|
|
|
|||
File diff suppressed because it is too large
Load Diff
|
|
@ -31,10 +31,43 @@ CCL_NAMESPACE_BEGIN
|
|||
# define GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE 512
|
||||
#endif
|
||||
|
||||
#ifdef __KERNEL_METAL__
|
||||
struct ActiveIndexContext {
|
||||
ActiveIndexContext(int _thread_index,
|
||||
int _global_index,
|
||||
int _threadgroup_size,
|
||||
int _simdgroup_size,
|
||||
int _simd_lane_index,
|
||||
int _simd_group_index,
|
||||
int _num_simd_groups,
|
||||
threadgroup int *_simdgroup_offset)
|
||||
: thread_index(_thread_index),
|
||||
global_index(_global_index),
|
||||
blocksize(_threadgroup_size),
|
||||
ccl_gpu_warp_size(_simdgroup_size),
|
||||
thread_warp(_simd_lane_index),
|
||||
warp_index(_simd_group_index),
|
||||
num_warps(_num_simd_groups),
|
||||
warp_offset(_simdgroup_offset)
|
||||
{
|
||||
}
|
||||
|
||||
const int thread_index, global_index, blocksize, ccl_gpu_warp_size, thread_warp, warp_index,
|
||||
num_warps;
|
||||
threadgroup int *warp_offset;
|
||||
|
||||
template<uint blocksizeDummy, typename IsActiveOp>
|
||||
void active_index_array(const uint num_states,
|
||||
ccl_global int *indices,
|
||||
ccl_global int *num_indices,
|
||||
IsActiveOp is_active_op)
|
||||
{
|
||||
const uint state_index = global_index;
|
||||
#else
|
||||
template<uint blocksize, typename IsActiveOp>
|
||||
__device__ void gpu_parallel_active_index_array(const uint num_states,
|
||||
int *indices,
|
||||
int *num_indices,
|
||||
ccl_global int *indices,
|
||||
ccl_global int *num_indices,
|
||||
IsActiveOp is_active_op)
|
||||
{
|
||||
extern ccl_gpu_shared int warp_offset[];
|
||||
|
|
@ -45,43 +78,62 @@ __device__ void gpu_parallel_active_index_array(const uint num_states,
|
|||
const uint warp_index = thread_index / ccl_gpu_warp_size;
|
||||
const uint num_warps = blocksize / ccl_gpu_warp_size;
|
||||
|
||||
/* Test if state corresponding to this thread is active. */
|
||||
const uint state_index = ccl_gpu_block_idx_x * blocksize + thread_index;
|
||||
const uint is_active = (state_index < num_states) ? is_active_op(state_index) : 0;
|
||||
#endif
|
||||
|
||||
/* For each thread within a warp compute how many other active states precede it. */
|
||||
const uint thread_mask = 0xFFFFFFFF >> (ccl_gpu_warp_size - thread_warp);
|
||||
const uint thread_offset = ccl_gpu_popc(ccl_gpu_ballot(is_active) & thread_mask);
|
||||
/* Test if state corresponding to this thread is active. */
|
||||
const uint is_active = (state_index < num_states) ? is_active_op(state_index) : 0;
|
||||
|
||||
/* Last thread in warp stores number of active states for each warp. */
|
||||
if (thread_warp == ccl_gpu_warp_size - 1) {
|
||||
warp_offset[warp_index] = thread_offset + is_active;
|
||||
}
|
||||
/* For each thread within a warp compute how many other active states precede it. */
|
||||
const uint thread_offset = ccl_gpu_popc(ccl_gpu_ballot(is_active) &
|
||||
ccl_gpu_thread_mask(thread_warp));
|
||||
|
||||
ccl_gpu_syncthreads();
|
||||
|
||||
/* Last thread in block converts per-warp sizes to offsets, increments global size of
|
||||
* index array and gets offset to write to. */
|
||||
if (thread_index == blocksize - 1) {
|
||||
/* TODO: parallelize this. */
|
||||
int offset = 0;
|
||||
for (int i = 0; i < num_warps; i++) {
|
||||
int num_active = warp_offset[i];
|
||||
warp_offset[i] = offset;
|
||||
offset += num_active;
|
||||
/* Last thread in warp stores number of active states for each warp. */
|
||||
if (thread_warp == ccl_gpu_warp_size - 1) {
|
||||
warp_offset[warp_index] = thread_offset + is_active;
|
||||
}
|
||||
|
||||
const uint block_num_active = warp_offset[warp_index] + thread_offset + is_active;
|
||||
warp_offset[num_warps] = atomic_fetch_and_add_uint32(num_indices, block_num_active);
|
||||
ccl_gpu_syncthreads();
|
||||
|
||||
/* Last thread in block converts per-warp sizes to offsets, increments global size of
|
||||
* index array and gets offset to write to. */
|
||||
if (thread_index == blocksize - 1) {
|
||||
/* TODO: parallelize this. */
|
||||
int offset = 0;
|
||||
for (int i = 0; i < num_warps; i++) {
|
||||
int num_active = warp_offset[i];
|
||||
warp_offset[i] = offset;
|
||||
offset += num_active;
|
||||
}
|
||||
|
||||
const uint block_num_active = warp_offset[warp_index] + thread_offset + is_active;
|
||||
warp_offset[num_warps] = atomic_fetch_and_add_uint32(num_indices, block_num_active);
|
||||
}
|
||||
|
||||
ccl_gpu_syncthreads();
|
||||
|
||||
/* Write to index array. */
|
||||
if (is_active) {
|
||||
const uint block_offset = warp_offset[num_warps];
|
||||
indices[block_offset + warp_offset[warp_index] + thread_offset] = state_index;
|
||||
}
|
||||
}
|
||||
|
||||
ccl_gpu_syncthreads();
|
||||
#ifdef __KERNEL_METAL__
|
||||
}; /* end class ActiveIndexContext */
|
||||
|
||||
/* Write to index array. */
|
||||
if (is_active) {
|
||||
const uint block_offset = warp_offset[num_warps];
|
||||
indices[block_offset + warp_offset[warp_index] + thread_offset] = state_index;
|
||||
}
|
||||
}
|
||||
/* inject the required thread params into a struct, and redirect to its templated member function
|
||||
*/
|
||||
# define gpu_parallel_active_index_array \
|
||||
ActiveIndexContext(metal_local_id, \
|
||||
metal_global_id, \
|
||||
metal_local_size, \
|
||||
simdgroup_size, \
|
||||
simd_lane_index, \
|
||||
simd_group_index, \
|
||||
num_simd_groups, \
|
||||
simdgroup_offset) \
|
||||
.active_index_array
|
||||
#endif
|
||||
|
||||
CCL_NAMESPACE_END
|
||||
|
|
|
|||
|
|
@ -33,10 +33,12 @@ CCL_NAMESPACE_BEGIN
|
|||
# define GPU_PARALLEL_PREFIX_SUM_DEFAULT_BLOCK_SIZE 512
|
||||
#endif
|
||||
|
||||
template<uint blocksize>
|
||||
__device__ void gpu_parallel_prefix_sum(int *counter, int *prefix_sum, const int num_values)
|
||||
__device__ void gpu_parallel_prefix_sum(const int global_id,
|
||||
ccl_global int *counter,
|
||||
ccl_global int *prefix_sum,
|
||||
const int num_values)
|
||||
{
|
||||
if (!(ccl_gpu_block_idx_x == 0 && ccl_gpu_thread_idx_x == 0)) {
|
||||
if (global_id != 0) {
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -33,16 +33,16 @@ CCL_NAMESPACE_BEGIN
|
|||
#endif
|
||||
#define GPU_PARALLEL_SORTED_INDEX_INACTIVE_KEY (~0)
|
||||
|
||||
template<uint blocksize, typename GetKeyOp>
|
||||
__device__ void gpu_parallel_sorted_index_array(const uint num_states,
|
||||
template<typename GetKeyOp>
|
||||
__device__ void gpu_parallel_sorted_index_array(const uint state_index,
|
||||
const uint num_states,
|
||||
const int num_states_limit,
|
||||
int *indices,
|
||||
int *num_indices,
|
||||
int *key_counter,
|
||||
int *key_prefix_sum,
|
||||
ccl_global int *indices,
|
||||
ccl_global int *num_indices,
|
||||
ccl_global int *key_counter,
|
||||
ccl_global int *key_prefix_sum,
|
||||
GetKeyOp get_key_op)
|
||||
{
|
||||
const uint state_index = ccl_gpu_block_idx_x * blocksize + ccl_gpu_thread_idx_x;
|
||||
const int key = (state_index < num_states) ? get_key_op(state_index) :
|
||||
GPU_PARALLEL_SORTED_INDEX_INACTIVE_KEY;
|
||||
|
||||
|
|
|
|||
|
|
@ -74,6 +74,7 @@ typedef unsigned long long uint64_t;
|
|||
#define ccl_gpu_block_idx_x (blockIdx.x)
|
||||
#define ccl_gpu_grid_dim_x (gridDim.x)
|
||||
#define ccl_gpu_warp_size (warpSize)
|
||||
#define ccl_gpu_thread_mask(thread_warp) uint(0xFFFFFFFF >> (ccl_gpu_warp_size - thread_warp))
|
||||
|
||||
#define ccl_gpu_global_id_x() (ccl_gpu_block_idx_x * ccl_gpu_block_dim_x + ccl_gpu_thread_idx_x)
|
||||
#define ccl_gpu_global_size_x() (ccl_gpu_grid_dim_x * ccl_gpu_block_dim_x)
|
||||
|
|
|
|||
|
|
@ -36,11 +36,35 @@
|
|||
/* Compute number of threads per block and minimum blocks per multiprocessor
|
||||
* given the maximum number of registers per thread. */
|
||||
|
||||
#define ccl_gpu_kernel(block_num_threads, thread_num_registers) \
|
||||
#define ccl_gpu_kernel_threads(block_num_threads) \
|
||||
extern "C" __global__ void __launch_bounds__(block_num_threads)
|
||||
|
||||
#define ccl_gpu_kernel_threads_registers(block_num_threads, thread_num_registers) \
|
||||
extern "C" __global__ void __launch_bounds__(block_num_threads, \
|
||||
GPU_MULTIPRESSOR_MAX_REGISTERS / \
|
||||
(block_num_threads * thread_num_registers))
|
||||
|
||||
/* allow ccl_gpu_kernel to accept 1 or 2 parameters */
|
||||
#define SELECT_MACRO(_1, _2, NAME, ...) NAME
|
||||
#define ccl_gpu_kernel(...) \
|
||||
SELECT_MACRO(__VA_ARGS__, ccl_gpu_kernel_threads_registers, ccl_gpu_kernel_threads)(__VA_ARGS__)
|
||||
|
||||
#define ccl_gpu_kernel_signature(name, ...) kernel_gpu_##name(__VA_ARGS__)
|
||||
|
||||
#define ccl_gpu_kernel_call(x) x
|
||||
|
||||
/* define a function object where "func" is the lambda body, and additional parameters are used to
|
||||
* specify captured state */
|
||||
#define ccl_gpu_kernel_lambda(func, ...) \
|
||||
struct KernelLambda { \
|
||||
__VA_ARGS__; \
|
||||
__device__ int operator()(const int state) \
|
||||
{ \
|
||||
return (func); \
|
||||
} \
|
||||
} ccl_gpu_kernel_lambda_pass; \
|
||||
ccl_gpu_kernel_lambda_pass
|
||||
|
||||
/* sanity checks */
|
||||
|
||||
#if GPU_KERNEL_BLOCK_NUM_THREADS > GPU_BLOCK_MAX_THREADS
|
||||
|
|
|
|||
|
|
@ -58,6 +58,95 @@ using namespace metal;
|
|||
|
||||
#define kernel_assert(cond)
|
||||
|
||||
#define ccl_gpu_global_id_x() metal_global_id
|
||||
#define ccl_gpu_warp_size simdgroup_size
|
||||
#define ccl_gpu_thread_idx_x simd_group_index
|
||||
#define ccl_gpu_thread_mask(thread_warp) uint64_t((1ull << thread_warp) - 1)
|
||||
|
||||
#define ccl_gpu_ballot(predicate) ((uint64_t)((simd_vote::vote_t)simd_ballot(predicate)))
|
||||
#define ccl_gpu_popc(x) popcount(x)
|
||||
|
||||
// clang-format off
|
||||
|
||||
/* kernel.h adapters */
|
||||
|
||||
#define ccl_gpu_kernel(...)
|
||||
|
||||
/* convert a comma-separated list into a semicolon-separated list (so that we can generate a struct based on kernel entrypoint parameters) */
|
||||
#define FN0()
|
||||
#define FN1(p1) p1;
|
||||
#define FN2(p1, p2) p1; p2;
|
||||
#define FN3(p1, p2, p3) p1; p2; p3;
|
||||
#define FN4(p1, p2, p3, p4) p1; p2; p3; p4;
|
||||
#define FN5(p1, p2, p3, p4, p5) p1; p2; p3; p4; p5;
|
||||
#define FN6(p1, p2, p3, p4, p5, p6) p1; p2; p3; p4; p5; p6;
|
||||
#define FN7(p1, p2, p3, p4, p5, p6, p7) p1; p2; p3; p4; p5; p6; p7;
|
||||
#define FN8(p1, p2, p3, p4, p5, p6, p7, p8) p1; p2; p3; p4; p5; p6; p7; p8;
|
||||
#define FN9(p1, p2, p3, p4, p5, p6, p7, p8, p9) p1; p2; p3; p4; p5; p6; p7; p8; p9;
|
||||
#define FN10(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10;
|
||||
#define FN11(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11;
|
||||
#define FN12(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12;
|
||||
#define FN13(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13;
|
||||
#define FN14(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14;
|
||||
#define FN15(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14; p15;
|
||||
#define FN16(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, p16) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14; p15; p16;
|
||||
#define GET_LAST_ARG(p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, p16, ...) p16
|
||||
#define PARAMS_MAKER(...) GET_LAST_ARG(__VA_ARGS__, FN16, FN15, FN14, FN13, FN12, FN11, FN10, FN9, FN8, FN7, FN6, FN5, FN4, FN3, FN2, FN1, FN0)
|
||||
|
||||
/* generate a struct containing the entrypoint parameters and a "run" method which can access them implicitly via this-> */
|
||||
#define ccl_gpu_kernel_signature(name, ...) \
|
||||
struct kernel_gpu_##name \
|
||||
{ \
|
||||
PARAMS_MAKER(__VA_ARGS__)(__VA_ARGS__) \
|
||||
void run(thread MetalKernelContext& context, \
|
||||
threadgroup int *simdgroup_offset, \
|
||||
const uint metal_global_id, \
|
||||
const ushort metal_local_id, \
|
||||
const ushort metal_local_size, \
|
||||
uint simdgroup_size, \
|
||||
uint simd_lane_index, \
|
||||
uint simd_group_index, \
|
||||
uint num_simd_groups) ccl_global const; \
|
||||
}; \
|
||||
kernel void kernel_metal_##name(device const kernel_gpu_##name *params_struct, \
|
||||
constant KernelParamsMetal &ccl_restrict _launch_params_metal, \
|
||||
constant MetalAncillaries *_metal_ancillaries, \
|
||||
threadgroup int *simdgroup_offset[[ threadgroup(0) ]], \
|
||||
const uint metal_global_id [[thread_position_in_grid]], \
|
||||
const ushort metal_local_id [[thread_position_in_threadgroup]], \
|
||||
const ushort metal_local_size [[threads_per_threadgroup]], \
|
||||
uint simdgroup_size [[threads_per_simdgroup]], \
|
||||
uint simd_lane_index [[thread_index_in_simdgroup]], \
|
||||
uint simd_group_index [[simdgroup_index_in_threadgroup]], \
|
||||
uint num_simd_groups [[simdgroups_per_threadgroup]]) { \
|
||||
MetalKernelContext context(_launch_params_metal, _metal_ancillaries); \
|
||||
INIT_DEBUG_BUFFER \
|
||||
params_struct->run(context, simdgroup_offset, metal_global_id, metal_local_id, metal_local_size, simdgroup_size, simd_lane_index, simd_group_index, num_simd_groups); \
|
||||
} \
|
||||
void kernel_gpu_##name::run(thread MetalKernelContext& context, \
|
||||
threadgroup int *simdgroup_offset, \
|
||||
const uint metal_global_id, \
|
||||
const ushort metal_local_id, \
|
||||
const ushort metal_local_size, \
|
||||
uint simdgroup_size, \
|
||||
uint simd_lane_index, \
|
||||
uint simd_group_index, \
|
||||
uint num_simd_groups) ccl_global const
|
||||
|
||||
#define ccl_gpu_kernel_call(x) context.x
|
||||
|
||||
/* define a function object where "func" is the lambda body, and additional parameters are used to specify captured state */
|
||||
#define ccl_gpu_kernel_lambda(func, ...) \
|
||||
struct KernelLambda \
|
||||
{ \
|
||||
KernelLambda(ccl_private MetalKernelContext &_context) : context(_context) {} \
|
||||
ccl_private MetalKernelContext &context; \
|
||||
__VA_ARGS__; \
|
||||
int operator()(const int state) const { return (func); } \
|
||||
}ccl_gpu_kernel_lambda_pass(context); ccl_gpu_kernel_lambda_pass
|
||||
|
||||
// clang-format on
|
||||
|
||||
/* make_type definitions with Metal style element initializers */
|
||||
#ifdef make_float2
|
||||
# undef make_float2
|
||||
|
|
@ -124,3 +213,38 @@ using namespace metal;
|
|||
#define logf(x) trigmode::log(float(x))
|
||||
|
||||
#define NULL 0
|
||||
|
||||
/* texture bindings and sampler setup */
|
||||
|
||||
struct Texture2DParamsMetal {
|
||||
texture2d<float, access::sample> tex;
|
||||
};
|
||||
struct Texture3DParamsMetal {
|
||||
texture3d<float, access::sample> tex;
|
||||
};
|
||||
|
||||
struct MetalAncillaries {
|
||||
device Texture2DParamsMetal *textures_2d;
|
||||
device Texture3DParamsMetal *textures_3d;
|
||||
};
|
||||
|
||||
enum SamplerType {
|
||||
SamplerFilterNearest_AddressRepeat,
|
||||
SamplerFilterNearest_AddressClampEdge,
|
||||
SamplerFilterNearest_AddressClampZero,
|
||||
|
||||
SamplerFilterLinear_AddressRepeat,
|
||||
SamplerFilterLinear_AddressClampEdge,
|
||||
SamplerFilterLinear_AddressClampZero,
|
||||
|
||||
SamplerCount
|
||||
};
|
||||
|
||||
constant constexpr array<sampler, SamplerCount> metal_samplers = {
|
||||
sampler(address::repeat, filter::nearest),
|
||||
sampler(address::clamp_to_edge, filter::nearest),
|
||||
sampler(address::clamp_to_zero, filter::nearest),
|
||||
sampler(address::repeat, filter::linear),
|
||||
sampler(address::clamp_to_edge, filter::linear),
|
||||
sampler(address::clamp_to_zero, filter::linear),
|
||||
};
|
||||
|
|
@ -0,0 +1,79 @@
|
|||
/*
|
||||
* Copyright 2021 Blender Foundation
|
||||
*
|
||||
* Licensed under the Apache License, Version 2.0 (the "License");
|
||||
* you may not use this file except in compliance with the License.
|
||||
* You may obtain a copy of the License at
|
||||
*
|
||||
* http://www.apache.org/licenses/LICENSE-2.0
|
||||
*
|
||||
* Unless required by applicable law or agreed to in writing, software
|
||||
* distributed under the License is distributed on an "AS IS" BASIS,
|
||||
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||
* See the License for the specific language governing permissions and
|
||||
* limitations under the License.
|
||||
*/
|
||||
|
||||
// clang-format off
|
||||
|
||||
/* Open the Metal kernel context class
|
||||
* Necessary to access resource bindings */
|
||||
class MetalKernelContext {
|
||||
public:
|
||||
constant KernelParamsMetal &launch_params_metal;
|
||||
constant MetalAncillaries *metal_ancillaries;
|
||||
|
||||
MetalKernelContext(constant KernelParamsMetal &_launch_params_metal, constant MetalAncillaries * _metal_ancillaries)
|
||||
: launch_params_metal(_launch_params_metal), metal_ancillaries(_metal_ancillaries)
|
||||
{}
|
||||
|
||||
/* texture fetch adapter functions */
|
||||
typedef uint64_t ccl_gpu_tex_object;
|
||||
|
||||
template<typename T>
|
||||
inline __attribute__((__always_inline__))
|
||||
T ccl_gpu_tex_object_read_2D(ccl_gpu_tex_object tex, float x, float y) const {
|
||||
kernel_assert(0);
|
||||
return 0;
|
||||
}
|
||||
template<typename T>
|
||||
inline __attribute__((__always_inline__))
|
||||
T ccl_gpu_tex_object_read_3D(ccl_gpu_tex_object tex, float x, float y, float z) const {
|
||||
kernel_assert(0);
|
||||
return 0;
|
||||
}
|
||||
|
||||
// texture2d
|
||||
template<>
|
||||
inline __attribute__((__always_inline__))
|
||||
float4 ccl_gpu_tex_object_read_2D(ccl_gpu_tex_object tex, float x, float y) const {
|
||||
const uint tid(tex);
|
||||
const uint sid(tex >> 32);
|
||||
return metal_ancillaries->textures_2d[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 tex, float x, float y) const {
|
||||
const uint tid(tex);
|
||||
const uint sid(tex >> 32);
|
||||
return metal_ancillaries->textures_2d[tid].tex.sample(metal_samplers[sid], float2(x, y)).x;
|
||||
}
|
||||
|
||||
// texture3d
|
||||
template<>
|
||||
inline __attribute__((__always_inline__))
|
||||
float4 ccl_gpu_tex_object_read_3D(ccl_gpu_tex_object tex, float x, float y, float z) const {
|
||||
const uint tid(tex);
|
||||
const uint sid(tex >> 32);
|
||||
return metal_ancillaries->textures_3d[tid].tex.sample(metal_samplers[sid], float3(x, y, z));
|
||||
}
|
||||
template<>
|
||||
inline __attribute__((__always_inline__))
|
||||
float ccl_gpu_tex_object_read_3D(ccl_gpu_tex_object tex, float x, float y, float z) const {
|
||||
const uint tid(tex);
|
||||
const uint sid(tex >> 32);
|
||||
return metal_ancillaries->textures_3d[tid].tex.sample(metal_samplers[sid], float3(x, y, z)).x;
|
||||
}
|
||||
# include "kernel/device/gpu/image.h"
|
||||
|
||||
// clang-format on
|
||||
|
|
@ -0,0 +1,23 @@
|
|||
/*
|
||||
* Copyright 2021 Blender Foundation
|
||||
*
|
||||
* Licensed under the Apache License, Version 2.0 (the "License");
|
||||
* you may not use this file except in compliance with the License.
|
||||
* You may obtain a copy of the License at
|
||||
*
|
||||
* http://www.apache.org/licenses/LICENSE-2.0
|
||||
*
|
||||
* Unless required by applicable law or agreed to in writing, software
|
||||
* distributed under the License is distributed on an "AS IS" BASIS,
|
||||
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||
* See the License for the specific language governing permissions and
|
||||
* limitations under the License.
|
||||
*/
|
||||
}
|
||||
; /* end of MetalKernelContext class definition */
|
||||
|
||||
/* Silently redirect into the MetalKernelContext instance */
|
||||
/* NOTE: These macros will need maintaining as entrypoints change */
|
||||
|
||||
#undef kernel_integrator_state
|
||||
#define kernel_integrator_state context.launch_params_metal.__integrator_state
|
||||
|
|
@ -0,0 +1,51 @@
|
|||
/*
|
||||
* Copyright 2021 Blender Foundation
|
||||
*
|
||||
* Licensed under the Apache License, Version 2.0 (the "License");
|
||||
* you may not use this file except in compliance with the License.
|
||||
* You may obtain a copy of the License at
|
||||
*
|
||||
* http://www.apache.org/licenses/LICENSE-2.0
|
||||
*
|
||||
* Unless required by applicable law or agreed to in writing, software
|
||||
* distributed under the License is distributed on an "AS IS" BASIS,
|
||||
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||
* See the License for the specific language governing permissions and
|
||||
* limitations under the License.
|
||||
*/
|
||||
|
||||
/* Constant Globals */
|
||||
|
||||
#include "kernel/types.h"
|
||||
#include "kernel/util/profiling.h"
|
||||
|
||||
#include "kernel/integrator/state.h"
|
||||
|
||||
CCL_NAMESPACE_BEGIN
|
||||
|
||||
typedef struct KernelParamsMetal {
|
||||
|
||||
#define KERNEL_TEX(type, name) ccl_constant type *name;
|
||||
#include "kernel/textures.h"
|
||||
#undef KERNEL_TEX
|
||||
|
||||
const IntegratorStateGPU __integrator_state;
|
||||
const KernelData data;
|
||||
|
||||
} KernelParamsMetal;
|
||||
|
||||
typedef struct KernelGlobalsGPU {
|
||||
int unused[1];
|
||||
} KernelGlobalsGPU;
|
||||
|
||||
typedef ccl_global const KernelGlobalsGPU *ccl_restrict KernelGlobals;
|
||||
|
||||
#define kernel_data launch_params_metal.data
|
||||
#define kernel_integrator_state launch_params_metal.__integrator_state
|
||||
|
||||
/* data lookup defines */
|
||||
|
||||
#define kernel_tex_fetch(tex, index) launch_params_metal.tex[index]
|
||||
#define kernel_tex_array(tex) launch_params_metal.tex
|
||||
|
||||
CCL_NAMESPACE_END
|
||||
|
|
@ -0,0 +1,25 @@
|
|||
/*
|
||||
* Copyright 2021 Blender Foundation
|
||||
*
|
||||
* Licensed under the Apache License, Version 2.0 (the "License");
|
||||
* you may not use this file except in compliance with the License.
|
||||
* You may obtain a copy of the License at
|
||||
*
|
||||
* http://www.apache.org/licenses/LICENSE-2.0
|
||||
*
|
||||
* Unless required by applicable law or agreed to in writing, software
|
||||
* distributed under the License is distributed on an "AS IS" BASIS,
|
||||
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||
* See the License for the specific language governing permissions and
|
||||
* limitations under the License.
|
||||
*/
|
||||
|
||||
/* Metal kernel entry points */
|
||||
|
||||
// clang-format off
|
||||
|
||||
#include "kernel/device/metal/compat.h"
|
||||
#include "kernel/device/metal/globals.h"
|
||||
#include "kernel/device/gpu/kernel.h"
|
||||
|
||||
// clang-format on
|
||||
|
|
@ -76,6 +76,7 @@ typedef unsigned long long uint64_t;
|
|||
#define ccl_gpu_block_idx_x (blockIdx.x)
|
||||
#define ccl_gpu_grid_dim_x (gridDim.x)
|
||||
#define ccl_gpu_warp_size (warpSize)
|
||||
#define ccl_gpu_thread_mask(thread_warp) uint(0xFFFFFFFF >> (ccl_gpu_warp_size - thread_warp))
|
||||
|
||||
#define ccl_gpu_global_id_x() (ccl_gpu_block_idx_x * ccl_gpu_block_dim_x + ccl_gpu_thread_idx_x)
|
||||
#define ccl_gpu_global_size_x() (ccl_gpu_grid_dim_x * ccl_gpu_block_dim_x)
|
||||
|
|
|
|||
Loading…
Reference in New Issue