diff --git a/bin/hipify-perl b/bin/hipify-perl index 3e39a5a11d..e375d22928 100755 --- a/bin/hipify-perl +++ b/bin/hipify-perl @@ -321,6 +321,12 @@ while (@ARGV) { $ft{'err'} += s/\bcudaDevAttrMaxSharedMemoryPerMultiprocessor\b/hipDeviceAttributeMaxSharedMemoryPerMultiprocessor/g; $ft{'err'} += s/\bcudaDevAttrMemoryClockRate\b/hipDeviceAttributeMemoryClockRate/g; $ft{'err'} += s/\bcudaDevAttrGlobalMemoryBusWidth\b/hipDeviceAttributeMemoryBusWidth/g; + $ft{'err'} += s/\bcudaDevAttrMaxTexture1DWidth\b/hipDeviceAttributeMaxTexture1DWidth/g; + $ft{'err'} += s/\bcudaDevAttrMaxTexture2DWidth\b/hipDeviceAttributeMaxTexture2DWidth/g; + $ft{'err'} += s/\bcudaDevAttrMaxTexture2DHeight\b/hipDeviceAttributeMaxTexture2DHeight/g; + $ft{'err'} += s/\bcudaDevAttrMaxTexture3DWidth\b/hipDeviceAttributeMaxTexture3DWidth/g; + $ft{'err'} += s/\bcudaDevAttrMaxTexture3DHeight\b/hipDeviceAttributeMaxTexture3DHeight/g; + $ft{'err'} += s/\bcudaDevAttrMaxTexture3DDepth\b/hipDeviceAttributeMaxTexture3DDepth/g; #-------- $ft{'dev'} += s/\bcudaDeviceAttr\b/hipDeviceAttribute_t/g; $ft{'dev'} += s/\bcudaDeviceGetAttribute\b/hipDeviceGetAttribute/g; diff --git a/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md b/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md index 663e6ad082..afd05b7d87 100644 --- a/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md +++ b/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md @@ -519,12 +519,12 @@ | 18 |*`cudaDevAttrIntegrated`* |*`hipDeviceAttributeIntegrated`* | | 19 |*`cudaDevAttrCanMapHostMemory`* | | | 20 |*`cudaDevAttrComputeMode`* |*`hipDeviceAttributeComputeMode`* | -| 21 |*`cudaDevAttrMaxTexture1DWidth`* | | -| 22 |*`cudaDevAttrMaxTexture2DWidth`* | | -| 23 |*`cudaDevAttrMaxTexture2DHeight`* | | -| 24 |*`cudaDevAttrMaxTexture3DWidth`* | | -| 25 |*`cudaDevAttrMaxTexture3DHeight`* | | -| 26 |*`cudaDevAttrMaxTexture3DDepth`* | | +| 21 |*`cudaDevAttrMaxTexture1DWidth`* |*`hipDeviceAttributeMaxTexture1DWidth`* | +| 22 |*`cudaDevAttrMaxTexture2DWidth`* |*`hipDeviceAttributeMaxTexture2DWidth`* | +| 23 |*`cudaDevAttrMaxTexture2DHeight`* |*`hipDeviceAttributeMaxTexture2DHeight`* | +| 24 |*`cudaDevAttrMaxTexture3DWidth`* |*`hipDeviceAttributeMaxTexture3DWidth`* | +| 25 |*`cudaDevAttrMaxTexture3DHeight`* |*`hipDeviceAttributeMaxTexture3DHeight`* | +| 26 |*`cudaDevAttrMaxTexture3DDepth`* |*`hipDeviceAttributeMaxTexture3DDepth`* | | 27 |*`cudaDevAttrMaxTexture2DLayeredWidth`* | | | 28 |*`cudaDevAttrMaxTexture2DLayeredHeight`* | | | 29 |*`cudaDevAttrMaxTexture2DLayeredLayers`* | | diff --git a/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp b/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp index 51dc9665c3..cc4502a866 100644 --- a/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp +++ b/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp @@ -328,17 +328,17 @@ const std::map CUDA_DRIVER_TYPE_NAME_MAP{ // cudaDevAttrComputeMode {"CU_DEVICE_ATTRIBUTE_COMPUTE_MODE", {"hipDeviceAttributeComputeMode", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 20 // cudaDevAttrMaxTexture1DWidth - {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_WIDTH", {"hipDeviceAttributeMaxTexture1DWidth", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 21 + {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_WIDTH", {"hipDeviceAttributeMaxTexture1DWidth", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 21 // cudaDevAttrMaxTexture2DWidth - {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_WIDTH", {"hipDeviceAttributeMaxTexture2DWidth", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 22 + {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_WIDTH", {"hipDeviceAttributeMaxTexture2DWidth", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 22 // cudaDevAttrMaxTexture2DHeight - {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_HEIGHT", {"hipDeviceAttributeMaxTexture2DHeight", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 23 + {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_HEIGHT", {"hipDeviceAttributeMaxTexture2DHeight", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 23 // cudaDevAttrMaxTexture3DWidth - {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_WIDTH", {"hipDeviceAttributeMaxTexture3DWidth", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 24 + {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_WIDTH", {"hipDeviceAttributeMaxTexture3DWidth", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 24 // cudaDevAttrMaxTexture3DHeight - {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_HEIGHT", {"hipDeviceAttributeMaxTexture3DHeight", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 25 + {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_HEIGHT", {"hipDeviceAttributeMaxTexture3DHeight", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 25 // cudaDevAttrMaxTexture3DDepth - {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_DEPTH", {"hipDeviceAttributeMaxTexture3DDepth", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 26 + {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_DEPTH", {"hipDeviceAttributeMaxTexture3DDepth", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 26 // cudaDevAttrMaxTexture2DLayeredWidth {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_WIDTH", {"hipDeviceAttributeMaxTexture2DLayeredWidth", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 27 // cudaDevAttrMaxTexture2DLayeredHeight diff --git a/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp b/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp index 6e86cf8fda..662655edb9 100644 --- a/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp +++ b/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp @@ -255,17 +255,17 @@ const std::map CUDA_RUNTIME_TYPE_NAME_MAP { // CU_DEVICE_ATTRIBUTE_COMPUTE_MODE {"cudaDevAttrComputeMode", {"hipDeviceAttributeComputeMode", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 20 // CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_WIDTH - {"cudaDevAttrMaxTexture1DWidth", {"hipDeviceAttributeMaxTexture1DWidth", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 21 + {"cudaDevAttrMaxTexture1DWidth", {"hipDeviceAttributeMaxTexture1DWidth", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 21 // CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_WIDTH - {"cudaDevAttrMaxTexture2DWidth", {"hipDeviceAttributeMaxTexture2DWidth", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 22 + {"cudaDevAttrMaxTexture2DWidth", {"hipDeviceAttributeMaxTexture2DWidth", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 22 // CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_HEIGHT - {"cudaDevAttrMaxTexture2DHeight", {"hipDeviceAttributeMaxTexture2DHeight", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 23 + {"cudaDevAttrMaxTexture2DHeight", {"hipDeviceAttributeMaxTexture2DHeight", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 23 // CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_WIDTH - {"cudaDevAttrMaxTexture3DWidth", {"hipDeviceAttributeMaxTexture3DWidth", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 24 + {"cudaDevAttrMaxTexture3DWidth", {"hipDeviceAttributeMaxTexture3DWidth", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 24 // CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_HEIGHT - {"cudaDevAttrMaxTexture3DHeight", {"hipDeviceAttributeMaxTexture3DHeight", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 25 + {"cudaDevAttrMaxTexture3DHeight", {"hipDeviceAttributeMaxTexture3DHeight", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 25 // CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_DEPTH - {"cudaDevAttrMaxTexture3DDepth", {"hipDeviceAttributeMaxTexture3DDepth", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 26 + {"cudaDevAttrMaxTexture3DDepth", {"hipDeviceAttributeMaxTexture3DDepth", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 26 // CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_WIDTH {"cudaDevAttrMaxTexture2DLayeredWidth", {"hipDeviceAttributeMaxTexture2DLayeredWidth", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 27 // CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_HEIGHT diff --git a/include/hip/hcc_detail/texture_functions.h b/include/hip/hcc_detail/texture_functions.h index 1f22e01177..bb0a9e7223 100644 --- a/include/hip/hcc_detail/texture_functions.h +++ b/include/hip/hcc_detail/texture_functions.h @@ -22,7 +22,6 @@ THE SOFTWARE. #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_TEXTURE_FUNCTIONS_H #define HIP_INCLUDE_HIP_HCC_DETAIL_TEXTURE_FUNCTIONS_H - #include #include @@ -118,8 +117,6 @@ union TData { #define TEXTURE_RETURN_UINT return texel.u.x; -#define TEXTURE_RETURN_FLOAT return texel.f.x; - #define TEXTURE_RETURN_SIGNED return texel.i.x; #define TEXTURE_RETURN_UNSIGNED return texel.u.x; @@ -136,8 +133,6 @@ union TData { #define TEXTURE_RETURN_UINT_X return make_uint1(texel.u.x); -#define TEXTURE_RETURN_FLOAT_X return make_float1(texel.f.x); - #define TEXTURE_RETURN_CHAR_XY return make_char2(texel.i.x, texel.i.y); #define TEXTURE_RETURN_UCHAR_XY return make_uchar2(texel.u.x, texel.u.y); @@ -150,8 +145,6 @@ union TData { #define TEXTURE_RETURN_UINT_XY return make_uint2(texel.u.x, texel.u.y); -#define TEXTURE_RETURN_FLOAT_XY return make_float2(texel.f.x, texel.f.y); - #define TEXTURE_RETURN_CHAR_XYZW return make_char4(texel.i.x, texel.i.y, texel.i.z, texel.i.w); #define TEXTURE_RETURN_UCHAR_XYZW return make_uchar4(texel.u.x, texel.u.y, texel.u.z, texel.u.w); @@ -164,10 +157,30 @@ union TData { #define TEXTURE_RETURN_UINT_XYZW return make_uint4(texel.u.x, texel.u.y, texel.u.z, texel.u.w); -#define TEXTURE_RETURN_FLOAT_XYZW return make_float4(texel.f.x, texel.f.y, texel.f.z, texel.f.w); +#define HIP_AD_FORMAT_NOT_INITIALIZED 0 + +#define TEXTURE_RETURN_FLOAT return (texFormatToSize[texRef.format] == 1)? texel.f.x : (float)texel.u.x/texFormatToSize[texRef.format]; + +#define TEXTURE_RETURN_FLOAT_X return (texFormatToSize[texRef.format] == 1)? make_float1(texel.f.x) : make_float1((float)texel.u.x/texFormatToSize[texRef.format]); + +#define TEXTURE_RETURN_FLOAT_XY return (texFormatToSize[texRef.format] == 1)? make_float2(texel.f.x, texel.f.y) : make_float2((float)texel.u.x/texFormatToSize[texRef.format], (float)texel.u.y/texFormatToSize[texRef.format]); + +#define TEXTURE_RETURN_FLOAT_XYZW return (texFormatToSize[texRef.format] == 1)? make_float4(texel.f.x, texel.f.y, texel.f.z, texel.f.w) : make_float4((float)texel.u.x/texFormatToSize[texRef.format], (float)texel.u.y/texFormatToSize[texRef.format], (float)texel.u.z/texFormatToSize[texRef.format], (float)texel.u.w/texFormatToSize[texRef.format]) ; extern "C" { + __device__ __constant__ static int texFormatToSize[] = { + [HIP_AD_FORMAT_NOT_INITIALIZED] = 1 , + [HIP_AD_FORMAT_UNSIGNED_INT8] = UCHAR_MAX , + [HIP_AD_FORMAT_UNSIGNED_INT16]= USHRT_MAX, + [HIP_AD_FORMAT_UNSIGNED_INT32]= 1 , + [HIP_AD_FORMAT_SIGNED_INT8] = SCHAR_MAX, + [HIP_AD_FORMAT_SIGNED_INT16] = SHRT_MAX, + [HIP_AD_FORMAT_SIGNED_INT32] = 1 , + [HIP_AD_FORMAT_HALF] = 1 , + [HIP_AD_FORMAT_FLOAT] = 1 +}; + __device__ __hip_float4_vector_value_type __ockl_image_sample_1D( unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s, @@ -3863,6 +3876,7 @@ __TEXTURE_FUNCTIONS_DECL__ float tex1D(texture texRef, TEXTURE_RETURN_FLOAT; } ////// + template __TEXTURE_FUNCTIONS_DECL__ float tex1D(texture texRef, float x) { TEXTURE_REF_PARAMETERS_INIT; diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index e3c10766e9..e6cdfdd304 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -116,7 +116,10 @@ typedef struct hipDeviceProp_t { int gcnArch; ///< AMD GCN Arch Value. Eg: 803, 701 int integrated; ///< APU vs dGPU int cooperativeLaunch; ///< HIP device supports cooperative launch - int cooperativeMultiDeviceLaunch; ///< HIP device supports cooperative launch on multiple devices + int cooperativeMultiDeviceLaunch; ///< HIP device supports cooperative launch on multiple devices + int maxTexture1D; ///< Maximum number of elements in 1D images + int maxTexture2D[2]; ///< Maximum dimensions (width, height) of 2D images, in image elements + int maxTexture3D[3]; ///< Maximum dimensions (width, height, depth) of 3D images, in image elements } hipDeviceProp_t; @@ -295,6 +298,14 @@ typedef enum hipDeviceAttribute_t { hipDeviceAttributeIntegrated, ///< iGPU hipDeviceAttributeCooperativeLaunch, ///< Support cooperative launch hipDeviceAttributeCooperativeMultiDeviceLaunch, ///< Support cooperative launch on multiple devices + + hipDeviceAttributeMaxTexture1DWidth, ///< Maximum number of elements in 1D images + hipDeviceAttributeMaxTexture2DWidth, ///< Maximum dimension width of 2D images in image elements + hipDeviceAttributeMaxTexture2DHeight, ///< Maximum dimension height of 2D images in image elements + hipDeviceAttributeMaxTexture3DWidth, ///< Maximum dimension width of 3D images in image elements + hipDeviceAttributeMaxTexture3DHeight, ///< Maximum dimensions height of 3D images in image elements + hipDeviceAttributeMaxTexture3DDepth ///< Maximum dimensions depth of 3D images in image elements + } hipDeviceAttribute_t; enum hipComputeMode { diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h index aca04067e3..2c294c91a3 100644 --- a/include/hip/nvcc_detail/hip_runtime_api.h +++ b/include/hip/nvcc_detail/hip_runtime_api.h @@ -772,6 +772,13 @@ inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t* p_prop, int dev p_prop->maxSharedMemoryPerMultiProcessor = cdprop.sharedMemPerMultiprocessor; p_prop->isMultiGpuBoard = cdprop.isMultiGpuBoard; + p_prop->maxTexture1D = cdprop.maxTexture1D; + p_prop->maxTexture2D[0] = cdprop.maxTexture2D[0]; + p_prop->maxTexture2D[1] = cdprop.maxTexture2D[1]; + p_prop->maxTexture3D[0] = cdprop.maxTexture3D[0]; + p_prop->maxTexture3D[1] = cdprop.maxTexture3D[1]; + p_prop->maxTexture3D[2] = cdprop.maxTexture3D[2]; + return hipCUDAErrorTohipError(cerror); } @@ -858,6 +865,24 @@ inline static hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t att case hipDeviceAttributeIntegrated: cdattr = cudaDevAttrIntegrated; break; + case hipDeviceAttributeMaxTexture1DWidth: + cdattr = cudaDevAttrMaxTexture1DWidth; + break; + case hipDeviceAttributeMaxTexture2DWidth: + cdattr = cudaDevAttrMaxTexture2DWidth; + break; + case hipDeviceAttributeMaxTexture2DHeight: + cdattr = cudaDevAttrMaxTexture2DHeight; + break; + case hipDeviceAttributeMaxTexture3DWidth: + cdattr = cudaDevAttrMaxTexture3DWidth; + break; + case hipDeviceAttributeMaxTexture3DHeight: + cdattr = cudaDevAttrMaxTexture3DHeight; + break; + case hipDeviceAttributeMaxTexture3DDepth: + cdattr = cudaDevAttrMaxTexture3DDepth; + break; default: cerror = cudaErrorInvalidValue; break; diff --git a/samples/1_Utils/hipCommander/c.cmd b/samples/1_Utils/hipCommander/c.cmd index db11071203..4cb980eccb 100644 --- a/samples/1_Utils/hipCommander/c.cmd +++ b/samples/1_Utils/hipCommander/c.cmd @@ -1,3 +1,3 @@ -loop,1000; H2D; NullKernel; D2H; endloop; +loop(1000); H2D; NullKernel; D2H; endloop; streamsync; -printTiming, 1000 +printTiming(1000) diff --git a/samples/1_Utils/hipCommander/classic.cmd b/samples/1_Utils/hipCommander/classic.cmd index c149eec5f7..7b1e4273c9 100644 --- a/samples/1_Utils/hipCommander/classic.cmd +++ b/samples/1_Utils/hipCommander/classic.cmd @@ -1 +1 @@ -H2D; NullKernel, D2H, streamsync +H2D; NullKernel; D2H; streamsync diff --git a/samples/1_Utils/hipCommander/hipCommander.cpp b/samples/1_Utils/hipCommander/hipCommander.cpp index 457cdfa7d3..21b5505623 100644 --- a/samples/1_Utils/hipCommander/hipCommander.cpp +++ b/samples/1_Utils/hipCommander/hipCommander.cpp @@ -7,11 +7,6 @@ #include #include -#ifdef __HIP_PLATFORM_HCC__ -#include -#include -#include -#endif #include @@ -283,100 +278,6 @@ class Command { #define FILENAME "nullkernel.hsaco" #define KERNEL_NAME "NullKernel" - -#ifdef __HIP_PLATFORM_HCC__ -//================================================================================================= -// Use Aql to launch the NULL kernel. -class AqlKernelCommand : public Command { - public: - AqlKernelCommand(CommandStream* cmdStream, const std::vector args) - : Command(cmdStream, args) { - hc::accelerator_view* av; - HIPCHECK(hipHccGetAcceleratorView(cmdStream->currentStream(), &av)); - - hc::accelerator acc = av->get_accelerator(); - - hsa_region_t systemRegion = *(hsa_region_t*)acc.get_hsa_am_system_region(); - - _hsaAgent = *(hsa_agent_t*)acc.get_hsa_agent(); - - std::ifstream file(FILENAME, std::ios::binary | std::ios::ate); - std::streamsize fsize = file.tellg(); - file.seekg(0, std::ios::beg); - - std::vector buffer(fsize); - if (file.read(buffer.data(), fsize)) { - uint64_t elfSize = ElfSize(&buffer[0]); - - assert(fsize == elfSize); - - // TODO - replace module load code with explicit module load and unload. - - hipModule_t module; - HIPCHECK(hipModuleLoadData(&module, &buffer[0])); - HIPCHECK(hipModuleGetFunction(&_function, module, KERNEL_NAME)); - - } else { - failed("could not open code object '%s'\n", FILENAME); - } - }; - - ~AqlKernelCommand(){}; - - void run() override { -#define LEN 64 - uint32_t len = LEN; - uint32_t one = 1; - - float* Ad = NULL; - - size_t argSize = 36; - char argBuffer[argSize]; - *(uint32_t*)(&argBuffer[0]) = len; - *(uint32_t*)(&argBuffer[4]) = one; - *(uint32_t*)(&argBuffer[8]) = one; - *(uint32_t*)(&argBuffer[12]) = len; - *(uint32_t*)(&argBuffer[16]) = one; - *(uint32_t*)(&argBuffer[20]) = one; - *(float**)(&argBuffer[24]) = Ad; // Ad pointer argument - - - void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0], - HIP_LAUNCH_PARAM_BUFFER_SIZE, &argSize, HIP_LAUNCH_PARAM_END}; - - hipModuleLaunchKernel(_function, len, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config); - }; - - - public: - hsa_queue_t _hsaQueue; - hsa_agent_t _hsaAgent; - - hipFunction_t _function; - - private: - static uint64_t ElfSize(const void* emi) { - const Elf64_Ehdr* ehdr = (const Elf64_Ehdr*)emi; - const Elf64_Shdr* shdr = (const Elf64_Shdr*)((char*)emi + ehdr->e_shoff); - - uint64_t max_offset = ehdr->e_shoff; - uint64_t total_size = max_offset + ehdr->e_shentsize * ehdr->e_shnum; - - for (uint16_t i = 0; i < ehdr->e_shnum; ++i) { - uint64_t cur_offset = static_cast(shdr[i].sh_offset); - if (max_offset < cur_offset) { - max_offset = cur_offset; - total_size = max_offset; - if (SHT_NOBITS != shdr[i].sh_type) { - total_size += static_cast(shdr[i].sh_size); - } - } - } - return total_size; - } -}; -#endif - //================================================================================================= // HCC optimizes away fully NULL kernel calls, so run one that is nearly null: class ModuleKernelCommand : public Command { @@ -391,26 +292,18 @@ class ModuleKernelCommand : public Command { void run() override { #define LEN 64 - uint32_t len = LEN; - uint32_t one = 1; + float *X = NULL; + HIPCHECK(hipMalloc((void**)&X, sizeof(float))); + struct { + float *Ad; + }args; + args.Ad = X; + size_t argSize = sizeof(args); - float* Ad = NULL; - - size_t argSize = 36; - char argBuffer[argSize]; - *(uint32_t*)(&argBuffer[0]) = len; - *(uint32_t*)(&argBuffer[4]) = one; - *(uint32_t*)(&argBuffer[8]) = one; - *(uint32_t*)(&argBuffer[12]) = len; - *(uint32_t*)(&argBuffer[16]) = one; - *(uint32_t*)(&argBuffer[20]) = one; - *(float**)(&argBuffer[24]) = Ad; // Ad pointer argument - - - void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0], + void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &argSize, HIP_LAUNCH_PARAM_END}; - hipModuleLaunchKernel(_function, len, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config); + hipModuleLaunchKernel(_function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config); }; @@ -447,33 +340,6 @@ class KernelCommand : public Command { hipStream_t _stream; }; - -#ifdef __HIP_PLATFORM_HCC__ -//================================================================================================= -class PfeCommand : public Command { - public: - PfeCommand(CommandStream* cmdStream, const std::vector args, - hipStream_t stream = 0) - : Command(cmdStream, args) { - HIPCHECK(hipHccGetAcceleratorView(stream, &_av)); - } - - ~PfeCommand() {} - - - void run() override { - static const int gridX = 64; - static const int groupX = 64; - auto cf = hc::parallel_for_each(*_av, hc::extent<1>(gridX).tile(groupX), - [=](hc::index<1>& idx) __HC__ {}); - } - - private: - hc::accelerator_view* _av; -}; -#endif - - //================================================================================================= class CopyCommand : public Command { enum MemType { PinnedHost, UnpinnedHost, Device }; @@ -708,10 +574,8 @@ CommandStream::CommandStream(std::string commandStreamString, int iterations) std::vector tokens; tokenize(commandStreamString, ';', tokens); - - std::for_each(tokens.begin(), tokens.end(), [&](const std::string s) { this->parse(s); }); - setStream(0); + std::for_each(tokens.begin(), tokens.end(), [&](const std::string s) { this->parse(s); }); } @@ -831,14 +695,6 @@ void CommandStream::parse(std::string fullCmd) { } else if (c == "vectoraddkernel") { cmd = new KernelCommand(cmdStream, args, KernelCommand::VectorAdd); -#ifdef __HIP_PLATFORM_HCC__ - } else if (c == "nullpfe") { - cmd = new PfeCommand(cmdStream, args); - - } else if (c == "aqlkernel") { - cmd = new AqlKernelCommand(cmdStream, args); -#endif - } else if (c == "devicesync") { cmd = new DeviceSyncCommand(cmdStream, args); diff --git a/samples/1_Utils/hipCommander/l2.hcm b/samples/1_Utils/hipCommander/l2.hcm index b541bd6a66..6b14f7b829 100644 --- a/samples/1_Utils/hipCommander/l2.hcm +++ b/samples/1_Utils/hipCommander/l2.hcm @@ -1,3 +1,3 @@ -setstream,1; +setstream(1); NullKernel; streamsync; -loop,10000; H2D; NullKernel; streamsync; endloop,1; +loop(10000); H2D; NullKernel; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/loop.hcm b/samples/1_Utils/hipCommander/loop.hcm index db11071203..4cb980eccb 100644 --- a/samples/1_Utils/hipCommander/loop.hcm +++ b/samples/1_Utils/hipCommander/loop.hcm @@ -1,3 +1,3 @@ -loop,1000; H2D; NullKernel; D2H; endloop; +loop(1000); H2D; NullKernel; D2H; endloop; streamsync; -printTiming, 1000 +printTiming(1000) diff --git a/samples/1_Utils/hipCommander/loop2.hcm b/samples/1_Utils/hipCommander/loop2.hcm index b8a14aa156..ae753d0722 100644 --- a/samples/1_Utils/hipCommander/loop2.hcm +++ b/samples/1_Utils/hipCommander/loop2.hcm @@ -1,2 +1,2 @@ -setstream,1; -loop,1000; NullKernel; syncstream; endloop,1, +setstream(1); +loop(1000); NullKernel; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/nullkernel.hsaco b/samples/1_Utils/hipCommander/nullkernel.hsaco index 585b55cce5..da6a3e6823 100755 Binary files a/samples/1_Utils/hipCommander/nullkernel.hsaco and b/samples/1_Utils/hipCommander/nullkernel.hsaco differ diff --git a/samples/1_Utils/hipCommander/perf/latency_hostsync.hcm b/samples/1_Utils/hipCommander/perf/latency_hostsync.hcm index 511ab355d5..f042b446e3 100644 --- a/samples/1_Utils/hipCommander/perf/latency_hostsync.hcm +++ b/samples/1_Utils/hipCommander/perf/latency_hostsync.hcm @@ -1,8 +1,8 @@ -setstream,1; +setstream(1); NullKernel; streamsync; -loop,100000; NullKernel; streamsync; endloop,1; +loop(100000); NullKernel; streamsync; endloop(1); -loop,100000; H2D; streamsync; NullKernel; streamsync; endloop,1; +loop(100000); H2D; streamsync; NullKernel; streamsync; endloop(1); -loop,100000; H2D; streamsync; NullKernel; streamsync; D2H; streamsync; endloop,1; +loop(100000); H2D; streamsync; NullKernel; streamsync; D2H; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/latency_nosync.hcm b/samples/1_Utils/hipCommander/perf/latency_nosync.hcm index c89d738be9..682d9d8b30 100644 --- a/samples/1_Utils/hipCommander/perf/latency_nosync.hcm +++ b/samples/1_Utils/hipCommander/perf/latency_nosync.hcm @@ -1,5 +1,5 @@ -setstream,1; +setstream(1); NullKernel; streamsync; -loop,100000; NullKernel; streamsync; endloop,1; -loop,100000; H2D; NullKernel; streamsync; endloop,1; -loop,100000; H2D; NullKernel; D2H; streamsync; endloop,1; +loop(100000); NullKernel; streamsync; endloop(1); +loop(100000); H2D; NullKernel; streamsync; endloop(1); +loop(100000); H2D; NullKernel; D2H; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/latency_nullstream.hcm b/samples/1_Utils/hipCommander/perf/latency_nullstream.hcm index 69345b23b0..87968a4df9 100644 --- a/samples/1_Utils/hipCommander/perf/latency_nullstream.hcm +++ b/samples/1_Utils/hipCommander/perf/latency_nullstream.hcm @@ -1,7 +1,7 @@ -setstream,0; +setstream(0); NullKernel; streamsync; -loop,100000; NullKernel; streamsync; endloop,1; +loop(100000); NullKernel; streamsync; endloop(1); -loop,100000; H2D; NullKernel; streamsync; endloop,1; +loop(100000); H2D; NullKernel; streamsync; endloop(1); -loop,100000; H2D; NullKernel; D2H; streamsync; endloop,1; +loop(100000); H2D; NullKernel; D2H; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/modulelaunch_latency.hcm b/samples/1_Utils/hipCommander/perf/modulelaunch_latency.hcm index d1d4091fad..576208135c 100644 --- a/samples/1_Utils/hipCommander/perf/modulelaunch_latency.hcm +++ b/samples/1_Utils/hipCommander/perf/modulelaunch_latency.hcm @@ -1,5 +1,5 @@ setstream(1); NullKernel; streamsync; loop(100); ModuleKernel; streamsync; endloop(1); -loop(100); AqlKernel; streamsync; endloop(1); +loop(100); streamsync; endloop(1); loop(3000); NullKernel; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/setstream.hcm b/samples/1_Utils/hipCommander/setstream.hcm index a7bdd093b8..22f1931ac4 100644 --- a/samples/1_Utils/hipCommander/setstream.hcm +++ b/samples/1_Utils/hipCommander/setstream.hcm @@ -1,3 +1,3 @@ -setstream,1; -setstream,2; H2D; NullKernel; D2H; +setstream(1); +setstream(2); H2D; NullKernel; D2H; streamsync diff --git a/src/hip_device.cpp b/src/hip_device.cpp index 746eaa2ddf..5c19ad42c6 100644 --- a/src/hip_device.cpp +++ b/src/hip_device.cpp @@ -276,6 +276,24 @@ hipError_t ihipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device case hipDeviceAttributeIntegrated: *pi = prop->integrated; break; + case hipDeviceAttributeMaxTexture1DWidth: + *pi = prop->maxTexture1D; + break; + case hipDeviceAttributeMaxTexture2DWidth: + *pi = prop->maxTexture2D[0]; + break; + case hipDeviceAttributeMaxTexture2DHeight: + *pi = prop->maxTexture2D[1]; + break; + case hipDeviceAttributeMaxTexture3DWidth: + *pi = prop->maxTexture3D[0]; + break; + case hipDeviceAttributeMaxTexture3DHeight: + *pi = prop->maxTexture3D[1]; + break; + case hipDeviceAttributeMaxTexture3DDepth: + *pi = prop->maxTexture3D[2]; + break; default: e = hipErrorInvalidValue; break; diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index c0a3827376..1b11b3c67e 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -43,7 +43,7 @@ THE SOFTWARE. #include #include #include "hsa/hsa_ext_amd.h" - +#include "hsa/hsa_ext_image.h" #include "hip/hip_runtime.h" #include "hip_hcc_internal.h" #include "trace_helper.h" @@ -913,6 +913,18 @@ hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop) { if(agent_profile == HSA_PROFILE_FULL) { prop->integrated = 1; } + + err = hsa_agent_get_info(_hsaAgent, (hsa_agent_info_t)HSA_EXT_AGENT_INFO_IMAGE_1D_MAX_ELEMENTS, + &prop->maxTexture1D); + DeviceErrorCheck(err); + + err = hsa_agent_get_info(_hsaAgent, (hsa_agent_info_t)HSA_EXT_AGENT_INFO_IMAGE_2D_MAX_ELEMENTS, + prop->maxTexture2D); + DeviceErrorCheck(err); + + err = hsa_agent_get_info(_hsaAgent, (hsa_agent_info_t)HSA_EXT_AGENT_INFO_IMAGE_3D_MAX_ELEMENTS, + prop->maxTexture3D); + DeviceErrorCheck(err); return e; } @@ -1399,7 +1411,7 @@ void ihipInit() { // Make sure the hip visible devices are within the deviceCnt range for (int i = 0; i < g_hip_visible_devices.size(); i++) { - if (g_hip_visible_devices[i] >= deviceCnt) { + if ((g_hip_visible_devices[i] >= deviceCnt) ||(g_hip_visible_devices[i] < 0)){ // Make sure any DeviceID after invalid DeviceID will be erased. g_hip_visible_devices.resize(i); break; @@ -1415,17 +1427,21 @@ void ihipInit() { g_deviceArray = new ihipDevice_t*[deviceCnt]; g_deviceCnt = 0; - for (int i = 0; i < accs.size(); i++) { - // check if the device id is included in the HIP_VISIBLE_DEVICES env variable - if (!accs[i].get_is_emulated()) { - if (std::find(g_hip_visible_devices.begin(), g_hip_visible_devices.end(), (i - 1)) == - g_hip_visible_devices.end() && - g_visible_device) { - // If device is not in visible devices list, ignore - continue; + + if(g_visible_device) { + for (int i = 0; i < g_hip_visible_devices.size(); i++) { + int devIndex = g_hip_visible_devices[i]; + if (!accs[devIndex+1].get_is_emulated()) { + g_deviceArray[g_deviceCnt] = new ihipDevice_t(g_deviceCnt, deviceCnt, accs[devIndex+1]); + g_deviceCnt++; + } + } + }else { + for (int i = 0; i < accs.size(); i++) { + if (!accs[i].get_is_emulated()) { + g_deviceArray[g_deviceCnt] = new ihipDevice_t(g_deviceCnt, deviceCnt, accs[i]); + g_deviceCnt++; } - g_deviceArray[g_deviceCnt] = new ihipDevice_t(g_deviceCnt, deviceCnt, accs[i]); - g_deviceCnt++; } } diff --git a/src/hip_module.cpp b/src/hip_module.cpp index 9680299540..594555e2b8 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -993,7 +993,7 @@ hipError_t ihipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* bloc tmp_min = min(maxWavesSGPRSLimited, tmp_min); activeWavefronts = min(maxWavesVGPRSLimited, tmp_min); - if (maxActivWaves < activeWavefronts) { + if (maxActivWaves <= activeWavefronts) { maxActivWaves = activeWavefronts; maxWavefronts = wavefrontsPerWG; } diff --git a/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp b/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp index ec82861013..1e1076b806 100644 --- a/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp +++ b/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp @@ -109,5 +109,12 @@ int main(int argc, char* argv[]) { CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxSharedMemoryPerMultiprocessor, props.maxSharedMemoryPerMultiProcessor)); CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeIntegrated, props.integrated)); + + CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxTexture1DWidth, props.maxTexture1D)); + CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxTexture2DWidth, props.maxTexture2D[0])); + CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxTexture2DHeight, props.maxTexture2D[1])); + CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxTexture3DWidth, props.maxTexture3D[0])); + CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxTexture3DHeight, props.maxTexture3D[1])); + CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxTexture3DDepth, props.maxTexture3D[2])); passed(); };