diff --git a/Makefile b/Makefile index b75af92ac..568846063 100644 --- a/Makefile +++ b/Makefile @@ -139,10 +139,6 @@ ifneq (${USE_SIMD},) MY_CMAKE_FLAGS += -DUSE_SIMD:STRING="${USE_SIMD}" endif -ifneq (${USE_BATCHED},) -MY_CMAKE_FLAGS += -DUSE_BATCHED:STRING="${USE_BATCHED}" -endif - ifneq (${VEC_REPORT},) MY_CMAKE_FLAGS += -DVEC_REPORT:BOOL="${VEC_REPORT}" endif @@ -379,7 +375,7 @@ help: @echo " avx, avx2, avx512f)" @echo " OSL_USE_OPTIX=1 Build the OptiX test renderer" @echo " USE_BATCHED=targets Build batched SIMD execution of shaders for (comma-separated choices:" - @echo " 0, b8_AVX, b8_AVX2, b8_AVX2_noFMA," + @echo " 0, b4_SSE2, b8_AVX, b8_AVX2, b8_AVX2_noFMA," @echo " b8_AVX512, b8_AVX512_noFMA," @echo " b16_AVX512, b16_AVX512_noFMA)" @echo " VEC_REPORT=0 Generate compiler vectorization reports" diff --git a/src/cmake/compiler.cmake b/src/cmake/compiler.cmake index fbd4aa10f..f13c6a4ff 100644 --- a/src/cmake/compiler.cmake +++ b/src/cmake/compiler.cmake @@ -338,7 +338,7 @@ endif () # # The USE_BATCHED option may be set to indicate that support for batched # SIMD shader execution be compiled along with targe specific libraries -set (USE_BATCHED "" CACHE STRING "Build batched SIMD shader execution for (0, b4_SSE2, b8_AVX, b8_AVX2, b8_AVX2_noFMA, b8_AVX512, b8_AVX512_noFMA, b16_AVX512, b16_AVX512_noFMA)") +set_cache (USE_BATCHED "" "Build batched SIMD shader execution for (0, b4_SSE2, b8_AVX, b8_AVX2, b8_AVX2_noFMA, b8_AVX512, b8_AVX512_noFMA, b16_AVX512, b16_AVX512_noFMA)") option (VEC_REPORT "Enable compiler's reporting system for vectorization" OFF) set (BATCHED_SUPPORT_DEFINES "") set (BATCHED_TARGET_LIBS "") diff --git a/src/include/OSL/batched_texture.h b/src/include/OSL/batched_texture.h index ca3bc893f..5f10cc5ff 100644 --- a/src/include/OSL/batched_texture.h +++ b/src/include/OSL/batched_texture.h @@ -18,9 +18,11 @@ using OIIO::Tex::Wrap; struct UniformTextureOptions { // Options that must be the same for all points we're texturing at once - int firstchannel = 0; ///< First channel of the lookup - int subimage = 0; ///< Subimage or face ID - ustring subimagename; ///< Subimage name + int firstchannel = 0; ///< First channel of the lookup + int subimage = 0; ///< Subimage or face ID + ustring subimagename; ///< Subimage name +#if defined(OIIO_TEXTUREOPTBATCH_VERSION) && OIIO_TEXTUREOPTBATCH_VERSION >= 2 + // Future expansion of an ideal v2 of OIIO's TextureOptBatch. But not yet. Tex::Wrap swrap = Tex::Wrap::Default; ///< Wrap mode in the s direction Tex::Wrap twrap = Tex::Wrap::Default; ///< Wrap mode in the t direction Tex::Wrap rwrap @@ -28,8 +30,19 @@ struct UniformTextureOptions { Tex::MipMode mipmode = Tex::MipMode::Default; ///< Mip mode Tex::InterpMode interpmode = Tex::InterpMode::SmartBicubic; ///< Interpolation mode - int anisotropic = 32; ///< Maximum anisotropic ratio - int conservative_filter = 1; ///< True: over-blur rather than alias + int anisotropic = 32; ///< Maximum anisotropic ratio + int conservative_filter = 1; ///< True: over-blur rather than alias +#else + // Original (v1) sizing and layout of the TextureOptBatch struct. + int swrap = int(Tex::Wrap::Default); ///< Wrap mode in the s direction + int twrap = int(Tex::Wrap::Default); ///< Wrap mode in the t direction + int rwrap = int(Tex::Wrap::Default); ///< Wrap mode in r (volumetric) + int mipmode = int(Tex::MipMode::Default); ///< Mip mode + int interpmode = int( + Tex::InterpMode::SmartBicubic); ///< Interpolation mode + int anisotropic = 32; ///< Maximum anisotropic ratio + int conservative_filter = 1; ///< True: over-blur rather than alias +#endif float fill = 0.0f; ///< Fill value for missing channels const float* missingcolor = nullptr; ///< Color for missing texture }; diff --git a/src/liboslexec/batched_llvm_gen.cpp b/src/liboslexec/batched_llvm_gen.cpp index 0b8aa41ee..d78554309 100644 --- a/src/liboslexec/batched_llvm_gen.cpp +++ b/src/liboslexec/batched_llvm_gen.cpp @@ -4274,8 +4274,15 @@ llvm_batched_texture_options(BatchedBackendLLVM& rop, int opnum, llvm::Value* wide_const_fzero_value = rop.ll.wide_constant(0.0f); llvm::Value* wide_const_fone_value = rop.ll.wide_constant(1.0f); llvm::Value* const_zero_value = rop.ll.constant(0); - llvm::Value* wrap_default_value = rop.ll.constant( +#if defined(OIIO_TEXTUREOPTBATCH_VERSION) && OIIO_TEXTUREOPTBATCH_VERSION >= 2 + // Possible future TextureOptBatch v2 -- not active yet + llvm::Value* wrap_default_value = rop.ll.constant8( + static_cast(Tex::Wrap::Default)); +#else + // OIIO <= 3.0 + llvm::Value* wrap_default_value = rop.ll.constant( static_cast(Tex::Wrap::Default)); +#endif llvm::Value* sblur = wide_const_fzero_value; llvm::Value* tblur = wide_const_fzero_value; @@ -4293,10 +4300,19 @@ llvm_batched_texture_options(BatchedBackendLLVM& rop, int opnum, llvm::Value* swrap = wrap_default_value; llvm::Value* twrap = wrap_default_value; llvm::Value* rwrap = wrap_default_value; - llvm::Value* mipmode = rop.ll.constant( +#if defined(OIIO_TEXTUREOPTBATCH_VERSION) && OIIO_TEXTUREOPTBATCH_VERSION >= 2 + // Possible future TextureOptBatch v2 -- not active yet + llvm::Value* mipmode = rop.ll.constant8( + static_cast(Tex::MipMode::Default)); + llvm::Value* interpmode = rop.ll.constant8( + static_cast(Tex::InterpMode::SmartBicubic)); +#else + // OIIO <= 3.0 + llvm::Value* mipmode = rop.ll.constant( static_cast(Tex::MipMode::Default)); llvm::Value* interpmode = rop.ll.constant( static_cast(Tex::InterpMode::SmartBicubic)); +#endif llvm::Value* anisotropic = rop.ll.constant(32); llvm::Value* conservative_filter = rop.ll.constant(1); llvm::Value* fill = rop.ll.constant(0.0f); @@ -4403,24 +4419,46 @@ llvm_batched_texture_options(BatchedBackendLLVM& rop, int opnum, continue; \ } -#define PARAM_UNIFORM_STRING_CODE(paramname, decoder, llvm_decoder, fieldname) \ - if (name == Strings::paramname && valtype == TypeDesc::STRING) { \ - if (valIsVarying) { \ - is_##fieldname##_uniform = false; \ - continue; \ - } \ - llvm::Value* val = nullptr; \ - if (Val.is_constant()) { \ - int mode = decoder(Val.get_string()); \ - val = rop.ll.constant(mode); \ - } else { \ - val = rop.llvm_load_value(Val); \ - llvm::Value* scalar_value_uh \ - = rop.ll.call_function("osl_gen_ustringhash_pod", val); \ - val = rop.ll.call_function(#llvm_decoder, scalar_value_uh); \ - } \ - fieldname = val; \ - continue; \ +#define PARAM_UNIFORM_STRING_INT_CODE(paramname, decoder, llvm_decoder, \ + fieldname) \ + if (name == Strings::paramname && valtype == TypeDesc::STRING) { \ + if (valIsVarying) { \ + is_##fieldname##_uniform = false; \ + continue; \ + } \ + llvm::Value* val = nullptr; \ + if (Val.is_constant()) { \ + int mode = int(decoder(Val.get_string())); \ + val = rop.ll.constant(mode); \ + } else { \ + val = rop.llvm_load_value(Val); \ + llvm::Value* scalar_value_uh \ + = rop.ll.call_function("osl_gen_ustringhash_pod", val); \ + val = rop.ll.call_function(#llvm_decoder, scalar_value_uh); \ + } \ + fieldname = val; \ + continue; \ + } + +#define PARAM_UNIFORM_STRING_UINT8_CODE(paramname, decoder, llvm_decoder, \ + fieldname) \ + if (name == Strings::paramname && valtype == TypeDesc::STRING) { \ + if (valIsVarying) { \ + is_##fieldname##_uniform = false; \ + continue; \ + } \ + llvm::Value* val = nullptr; \ + if (Val.is_constant()) { \ + int mode = int(decoder(Val.get_string())); \ + val = rop.ll.constant8(uint8_t(mode)); \ + } else { \ + val = rop.llvm_load_value(Val); \ + llvm::Value* scalar_value_uh \ + = rop.ll.call_function("osl_gen_ustringhash_pod", val); \ + val = rop.ll.call_function(#llvm_decoder, scalar_value_uh); \ + } \ + fieldname = val; \ + continue; \ } if (name == Strings::wrap && valtype == TypeDesc::STRING) { @@ -4434,7 +4472,7 @@ llvm_batched_texture_options(BatchedBackendLLVM& rop, int opnum, } llvm::Value* val = nullptr; if (Val.is_constant()) { - int mode = TextureOpt::decode_wrapmode(Val.get_string()); + int mode = int(TextureOpt::decode_wrapmode(Val.get_string())); val = rop.ll.constant(mode); } else { val = rop.llvm_load_value(Val); @@ -4450,14 +4488,33 @@ llvm_batched_texture_options(BatchedBackendLLVM& rop, int opnum, } continue; } - PARAM_UNIFORM_STRING_CODE(swrap, OIIO::TextureOpt::decode_wrapmode, - osl_texture_decode_wrapmode, swrap) - PARAM_UNIFORM_STRING_CODE(twrap, OIIO::TextureOpt::decode_wrapmode, - osl_texture_decode_wrapmode, twrap) +#if defined(OIIO_TEXTUREOPTBATCH_VERSION) && OIIO_TEXTUREOPTBATCH_VERSION >= 2 + // Possible future TextureOptBatch v2 -- not active yet + PARAM_UNIFORM_STRING_UINT8_CODE(swrap, OIIO::Tex::decode_wrapmode, + osl_texture_decode_wrapmode, swrap) + PARAM_UNIFORM_STRING_UINT8_CODE(twrap, OIIO::Tex::decode_wrapmode, + osl_texture_decode_wrapmode, twrap) + if (tex3d) { + PARAM_UNIFORM_STRING_UINT8_CODE(rwrap, OIIO::Tex::decode_wrapmode, + osl_texture_decode_wrapmode, rwrap) + } + PARAM_UNIFORM_STRING_UINT8_CODE(interp, tex_interp_to_code, + osl_texture_decode_interpmode, + interpmode) +#else + // OIIO <= 3.0 + PARAM_UNIFORM_STRING_INT_CODE(swrap, OIIO::TextureOpt::decode_wrapmode, + osl_texture_decode_wrapmode, swrap) + PARAM_UNIFORM_STRING_INT_CODE(twrap, OIIO::TextureOpt::decode_wrapmode, + osl_texture_decode_wrapmode, twrap) if (tex3d) { - PARAM_UNIFORM_STRING_CODE(rwrap, OIIO::TextureOpt::decode_wrapmode, - osl_texture_decode_wrapmode, rwrap) + PARAM_UNIFORM_STRING_INT_CODE(rwrap, + OIIO::TextureOpt::decode_wrapmode, + osl_texture_decode_wrapmode, rwrap) } + PARAM_UNIFORM_STRING_INT_CODE(interp, tex_interp_to_code, + osl_texture_decode_interpmode, interpmode) +#endif PARAM_UNIFORM_FLOAT(fill) PARAM_UNIFORM_INT(firstchannel) @@ -4479,10 +4536,6 @@ llvm_batched_texture_options(BatchedBackendLLVM& rop, int opnum, continue; } - PARAM_UNIFORM_STRING_CODE(interp, tex_interp_to_code, - osl_texture_decode_interpmode, interpmode) - - if (name == Strings::alpha && valtype == TypeDesc::FLOAT) { OSL_ASSERT( valIsVarying @@ -4569,7 +4622,7 @@ llvm_batched_texture_options(BatchedBackendLLVM& rop, int opnum, #undef PARAM_WIDE_FLOAT_S_T_R #undef PARAM_UNIFORM_FLOAT #undef PARAM_UNIFORM_INT -#undef PARAM_UNIFORM_STRING_CODE +#undef PARAM_UNIFORM_STRING_INT_CODE } // The LLVMMemberIndex will be the same for any width of BatchedTextureOptions, diff --git a/src/liboslexec/batched_llvm_instance.cpp b/src/liboslexec/batched_llvm_instance.cpp index 332e814f6..4e5a46c33 100644 --- a/src/liboslexec/batched_llvm_instance.cpp +++ b/src/liboslexec/batched_llvm_instance.cpp @@ -745,14 +745,24 @@ BatchedBackendLLVM::llvm_type_batched_texture_options() sg_types.push_back(ll.type_wide_float()); // rnd // Uniform values of the batch - sg_types.push_back(ll.type_int()); // firstchannel - sg_types.push_back(ll.type_int()); // subimage - sg_types.push_back(vp); // subimagename - sg_types.push_back(ll.type_int()); // swrap - sg_types.push_back(ll.type_int()); // twrap - sg_types.push_back(ll.type_int()); // rwrap - sg_types.push_back(ll.type_int()); // mipmode - sg_types.push_back(ll.type_int()); // interpmode + sg_types.push_back(ll.type_int()); // firstchannel + sg_types.push_back(ll.type_int()); // subimage + sg_types.push_back(vp); // subimagename +#if defined(OIIO_TEXTUREOPTBATCH_VERSION) && OIIO_TEXTUREOPTBATCH_VERSION >= 2 + // Possible future TextureOptBatch v2 -- not active yet + sg_types.push_back(ll.type_int8()); // swrap + sg_types.push_back(ll.type_int8()); // twrap + sg_types.push_back(ll.type_int8()); // rwrap + sg_types.push_back(ll.type_int8()); // mipmode + sg_types.push_back(ll.type_int8()); // interpmode +#else + // OIIO <= 3.0 + sg_types.push_back(ll.type_int()); // swrap + sg_types.push_back(ll.type_int()); // twrap + sg_types.push_back(ll.type_int()); // rwrap + sg_types.push_back(ll.type_int()); // mipmode + sg_types.push_back(ll.type_int()); // interpmode +#endif sg_types.push_back(ll.type_int()); // anisotropic sg_types.push_back(ll.type_int()); // conservative_filter sg_types.push_back(ll.type_float()); // fill diff --git a/src/liboslexec/constfold.cpp b/src/liboslexec/constfold.cpp index 460c77aa6..2e7664f27 100644 --- a/src/liboslexec/constfold.cpp +++ b/src/liboslexec/constfold.cpp @@ -2501,10 +2501,10 @@ DECLFOLDER(constfold_texture) // Keep from repeating the same tedious code for {s,t,r, }{width,blur,wrap} #define CHECK(field, ctype, osltype) \ if (name == Strings::field && !field##_set) { \ - if (valuetype == osltype && *(ctype*)value == opt.field) \ + if (valuetype == osltype && *(ctype*)value == (ctype)opt.field) \ elide = true; \ else if (osltype == TypeDesc::FLOAT && valuetype == TypeDesc::INT \ - && *(int*)value == opt.field) \ + && *(int*)value == (int)opt.field) \ elide = true; \ else \ field##_set = true; \ @@ -2520,8 +2520,8 @@ DECLFOLDER(constfold_texture) { \ if (valuetype == osltype) { \ ctype* v = (ctype*)value; \ - if (*v == opt.s##field && *v == opt.t##field \ - && *v == opt.r##field) \ + if (*v == (ctype)opt.s##field && *v == (ctype)opt.t##field \ + && *v == (ctype)opt.r##field) \ elide = true; \ else { \ s##field##_set = true; \ @@ -2530,8 +2530,8 @@ DECLFOLDER(constfold_texture) } \ } else if (osltype == TypeDesc::FLOAT && valuetype == TypeDesc::INT) { \ int* v = (int*)value; \ - if (*v == opt.s##field && *v == opt.t##field \ - && *v == opt.r##field) \ + if (*v == (ctype)opt.s##field && *v == (ctype)opt.t##field \ + && *v == (ctype)opt.r##field) \ elide = true; \ else { \ s##field##_set = true; \ @@ -2573,7 +2573,8 @@ DECLFOLDER(constfold_texture) else if (name == Strings::interp && !interp_set) { if (value && valuetype == TypeDesc::STRING - && tex_interp_to_code(*(ustring*)value) == opt.interpmode) + && tex_interp_to_code(*(ustring*)value) + == (int)opt.interpmode) elide = true; else interp_set = true; diff --git a/src/liboslexec/llvm_gen.cpp b/src/liboslexec/llvm_gen.cpp index a01c0f008..5b7b1d824 100644 --- a/src/liboslexec/llvm_gen.cpp +++ b/src/liboslexec/llvm_gen.cpp @@ -2574,7 +2574,8 @@ llvm_gen_texture_options(BackendLLVM& rop, int opnum, int first_optional_arg, bool sblur_set = false, tblur_set = false, rblur_set = false; bool swrap_set = false, twrap_set = false, rwrap_set = false; bool firstchannel_set = false, fill_set = false, interp_set = false; - bool time_set = false, subimage_set = false; + // bool time_set = false; + bool subimage_set = false; Opcode& op(rop.inst()->ops()[opnum]); for (int a = first_optional_arg; a < op.nargs(); ++a) { @@ -2645,8 +2646,8 @@ llvm_gen_texture_options(BackendLLVM& rop, int opnum, int first_optional_arg, #define PARAM_STRING_CODE(paramname, decoder, fieldname) \ if (name == Strings::paramname && valtype == TypeDesc::STRING) { \ if (Val.is_constant()) { \ - int code = decoder(Val.get_string()); \ - if (!paramname##_set && code == optdefaults.fieldname) \ + int code = (int)decoder(Val.get_string()); \ + if (!paramname##_set && code == (int)optdefaults.fieldname) \ continue; \ if (code >= 0) { \ llvm::Value* val = rop.ll.constant(code); \ @@ -2672,7 +2673,7 @@ llvm_gen_texture_options(BackendLLVM& rop, int opnum, int first_optional_arg, if (name == Strings::wrap && valtype == TypeDesc::STRING) { if (Val.is_constant()) { - int mode = TextureOpt::decode_wrapmode(Val.get_string()); + int mode = (int)TextureOpt::decode_wrapmode(Val.get_string()); llvm::Value* val = rop.ll.constant(mode); rop.ll.call_function("osl_texture_set_stwrap_code", opt, val); if (tex3d) @@ -2692,7 +2693,6 @@ llvm_gen_texture_options(BackendLLVM& rop, int opnum, int first_optional_arg, PARAM_STRING_CODE(rwrap, TextureOpt::decode_wrapmode, rwrap) PARAM_FLOAT(fill) - PARAM_FLOAT(time) PARAM_INT(firstchannel) PARAM_INT(subimage) @@ -2751,6 +2751,16 @@ llvm_gen_texture_options(BackendLLVM& rop, int opnum, int first_optional_arg, rop.ll.constant(nchans), val); continue; } + + // PARAM_FLOAT(time) + if (name == Strings::time + && (valtype == TypeDesc::FLOAT || valtype == TypeDesc::INT)) { + // NOTE: currently no supported 3d texture format makes use of + // time. So there is no time in the TextureOpt struct, but will + // silently accept and ignore the time option. + continue; + } + rop.shadingcontext()->errorfmt( "Unknown texture{} optional argument: \"{}\", <{}> ({}:{})", tex3d ? "3d" : "", name, valtype, op.sourcefile(), op.sourceline()); diff --git a/src/liboslexec/optexture.cpp b/src/liboslexec/optexture.cpp index 11dd54c67..42b91421c 100644 --- a/src/liboslexec/optexture.cpp +++ b/src/liboslexec/optexture.cpp @@ -28,8 +28,10 @@ namespace pvt { OSL_SHADEOP OSL_HOSTDEVICE void osl_init_texture_options(OpaqueExecContextPtr oec, void* opt) { +#if defined(OIIO_TEXTUREOPT_VERSION) && OIIO_TEXTUREOPT_VERSION >= 2 + new (opt) TextureOpt; +#else // TODO: Simplify when TextureOpt() has __device__ marker. - // new (opt) TextureOpt; TextureOpt* o = reinterpret_cast(opt); o->firstchannel = 0; o->subimage = 0; @@ -46,20 +48,21 @@ osl_init_texture_options(OpaqueExecContextPtr oec, void* opt) o->twidth = 1.0f; o->fill = 0.0f; o->missingcolor = nullptr; - o->time = 0.0f; + o->time = 0.0f; // Deprecated o->rnd = -1.0f; - o->samples = 1; + o->samples = 1; // Deprecated o->rwrap = TextureOpt::WrapDefault; o->rblur = 0.0f; o->rwidth = 1.0f; -#ifdef OIIO_TEXTURESYSTEM_SUPPORTS_COLORSPACE +# ifdef OIIO_TEXTURESYSTEM_SUPPORTS_COLORSPACE o->colortransformid = 0; int* envlayout = (int*)&o->colortransformid + 1; -#else +# else int* envlayout = (int*)&o->rwidth + 1; -#endif +# endif // envlayout is private so we access it from the last public member for now. *envlayout = 0; +#endif } @@ -89,7 +92,7 @@ decode_wrapmode(ustringhash_pod name_) OSL_SHADEOP OSL_HOSTDEVICE int osl_texture_decode_wrapmode(ustringhash_pod name_) { - return decode_wrapmode(name_); + return (int)decode_wrapmode(name_); } OSL_SHADEOP OSL_HOSTDEVICE void @@ -202,7 +205,8 @@ osl_texture_set_fill(void* opt, float x) OSL_SHADEOP OSL_HOSTDEVICE void osl_texture_set_time(void* opt, float x) { - ((TextureOpt*)opt)->time = x; + // Not used by the texture system + // ((TextureOpt*)opt)->time = x; } OSL_SHADEOP OSL_HOSTDEVICE int diff --git a/src/liboslexec/oslexec_pvt.h b/src/liboslexec/oslexec_pvt.h index d9b8f1222..d1a0f35a7 100644 --- a/src/liboslexec/oslexec_pvt.h +++ b/src/liboslexec/oslexec_pvt.h @@ -23,9 +23,12 @@ # include #endif +#include + #include #include #include +#include #include #include @@ -2491,13 +2494,13 @@ tex_interp_to_code(ustringhash modename) { int mode = -1; if (modename == Hashes::smartcubic) - mode = TextureOpt::InterpSmartBicubic; + mode = (int)TextureOpt::InterpSmartBicubic; else if (modename == Hashes::linear) - mode = TextureOpt::InterpBilinear; + mode = (int)TextureOpt::InterpBilinear; else if (modename == Hashes::cubic) - mode = TextureOpt::InterpBicubic; + mode = (int)TextureOpt::InterpBicubic; else if (modename == Hashes::closest) - mode = TextureOpt::InterpClosest; + mode = (int)TextureOpt::InterpClosest; return mode; } @@ -2507,13 +2510,13 @@ tex_interp_to_code(ustring modename) { int mode = -1; if (modename == Strings::smartcubic) - mode = TextureOpt::InterpSmartBicubic; + mode = (int)TextureOpt::InterpSmartBicubic; else if (modename == Strings::linear) - mode = TextureOpt::InterpBilinear; + mode = (int)TextureOpt::InterpBilinear; else if (modename == Strings::cubic) - mode = TextureOpt::InterpBicubic; + mode = (int)TextureOpt::InterpBicubic; else if (modename == Strings::closest) - mode = TextureOpt::InterpClosest; + mode = (int)TextureOpt::InterpClosest; return mode; } #endif