AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-561273 - hip samples on TheRock build using HIP LANGUAGE and hip-lang package. (#1794)

Browse Source
This commit is contained in:
Jaydeep
2026-01-29 13:45:58 +05:30
committed by GitHub
parent fa6f071751
commit 190d9a8e27
30 changed files with 740 additions and 430 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip-tests/samples/2_Cookbook/11_texture_driver/CMakeLists.txt
+30 -18
View File
@@ -18,9 +18,7 @@
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
# THE SOFTWARE.
project(texture2dDrv)
cmake_minimum_required(VERSION 3.10)
cmake_minimum_required(VERSION 3.21)
if(UNIX)
if(NOT DEFINED ROCM_PATH)
@@ -34,36 +32,50 @@ if(UNIX)
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH})
endif()
# Find hip
find_package(hip)
project(texture2dDrv LANGUAGES CXX HIP)
# Set compiler and linker
set(CMAKE_CXX_COMPILER ${HIP_HIPCC_EXECUTABLE})
set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE})
set(CMAKE_BUILD_TYPE Release)
# Detect GPU architectures and setup OFFLOAD_ARCH_FLAGS
include(${CMAKE_CURRENT_SOURCE_DIR}/../../common/DetectGPUArchs.cmake)
# Create the excutable
if(TARGET build_cookbook)
set(EXCLUDE_OPTION EXCLUDE_FROM_ALL)
set(EXCLUDE_OPTION EXCLUDE_FROM_ALL)
else()
set(EXCLUDE_OPTION )
set(EXCLUDE_OPTION)
endif()
# Mark source files as HIP code
set_source_files_properties(texture2dDrv.cpp PROPERTIES LANGUAGE HIP)
set_source_files_properties(tex2dKernel.cpp PROPERTIES LANGUAGE HIP)
add_executable(texture2dDrv ${EXCLUDE_OPTION} texture2dDrv.cpp)
# Generate code object
add_custom_target(
codeobj3
COMMAND ${HIP_HIPCC_EXECUTABLE} --cuda-device-only ${CMAKE_CURRENT_SOURCE_DIR}/tex2dKernel.cpp -o tex2dKernel.code
COMMENT "codeobj3 generated"
# Generate code object using CMake's HIP language support
# Note: Build for all detected architectures to support multiple GPUs (like hipcc does)
add_custom_command(
OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/tex2dKernel.code
COMMAND ${CMAKE_HIP_COMPILER} -x hip --cuda-device-only ${OFFLOAD_ARCH_FLAGS} ${CMAKE_CURRENT_SOURCE_DIR}/tex2dKernel.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/tex2dKernel.code
DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/tex2dKernel.cpp
COMMENT "Generating code object tex2dKernel.code for ${CMAKE_HIP_ARCHITECTURES}"
VERBATIM
)
add_custom_target(codeobj3 ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/tex2dKernel.code)
add_dependencies(texture2dDrv codeobj3)
target_include_directories(texture2dDrv PRIVATE ../../common)
# Link with HIP
target_link_libraries(texture2dDrv hip::host)
# Set RPATH so executable can find HIP libraries at runtime
if(UNIX)
set_target_properties(texture2dDrv PROPERTIES
BUILD_RPATH "${ROCM_PATH}/lib"
INSTALL_RPATH "${ROCM_PATH}/lib"
)
endif()
if(TARGET build_cookbook)
add_dependencies(build_cookbook texture2dDrv)
add_dependencies(build_cookbook texture2dDrv)
endif()
projects/hip-tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp
+5 -4
View File
@@ -24,6 +24,7 @@ THE SOFTWARE.
#include <iostream>
#include <fstream>
#include <vector>
#include <cstring>
#include "hip_helper.h"
#define fileName "tex2dKernel.code"
@@ -119,7 +120,7 @@ template <typename T> bool runTest(hipModule_t& module, const char* refName, con
unsigned int height = 256;
unsigned int size = width * height * sizeof(T);
T* hData = (T*)malloc(size);
memset(hData, 0, size);
std::memset(hData, 0, size);
for (int i = 0; i < height; i++) {
for (int j = 0; j < width; j++) {
initVal(hData[i * width + j]);
@@ -136,12 +137,12 @@ template <typename T> bool runTest(hipModule_t& module, const char* refName, con
hipMemcpyHostToDevice));
hipResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
std::memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = hipResourceTypeArray;
resDesc.res.array.array = array;
hipTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
std::memset(&texDesc, 0, sizeof(texDesc));
texDesc.addressMode[0] = hipAddressModeClamp;
texDesc.addressMode[1] = hipAddressModeClamp;
texDesc.filterMode = hipFilterModePoint;
@@ -180,7 +181,7 @@ template <typename T> bool runTest(hipModule_t& module, const char* refName, con
checkHipErrors(hipDeviceSynchronize());
T* hOutputData = (T*)malloc(size);
memset(hOutputData, 0, size);
std::memset(hOutputData, 0, size);
checkHipErrors(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost));
for (int i = 0; i < height; i++) {