AI/rocm-systems
2
0
انشعاب 0
کد مسائل تقاضاهای واکشی Actions Packages پروژه‌ها انتشارها دانشنامه فعالیت

SWDEV-299127 - Merge 'develop' into 'amd-staging'

فهرست منبع 
Change-Id: Id10ef743148a5b638a2cc246b0a55e1a7cfada3f


[ROCm/hip commit: 474fd7f695]
This commit is contained in:
Jenkins
2021-10-03 19:11:20 -04:00
والد 8eaf01d188 266355ffb3
کامیت 5cf5dca9e3
46فایلهای تغییر یافته به همراه1956 افزوده شده و 298 حذف شده
دانلود پرونده وصله دانلود فایل تغییرات diff Expand all files Collapse all files
projects/hip/bin/hipcc
+4 -1
مشاهده پرونده
@@ -229,6 +229,7 @@ my $needCXXFLAGS = 0; # need to add CXX flags to compile step
my $needCFLAGS = 0; # need to add C flags to compile step
my $needLDFLAGS = 1; # need to add LDFLAGS to compile step.
my $fileTypeFlag = 0; # to see if -x flag is mentioned
my $hasOMPTargets = 0; # If OMP targets is mentioned
my $hasC = 0; # options contain a c-style file
my $hasCXX = 0; # options contain a cpp-style file (NVCC must force recognition as GPU file)
my $hasCU = 0; # options contain a cu-style file (HCC must force recognition as GPU file)
@@ -517,6 +518,8 @@ foreach $arg (@ARGV)
$hasC = 0;
$hasCXX = 0;
$hasHIP = 1;
} elsif ($arg =~ '-fopenmp-targets=') {
$hasOMPTargets = 1;
} elsif ($arg =~ m/^-/) {
# options start with -
if ($arg eq '-fgpu-rdc') {
@@ -555,7 +558,7 @@ foreach $arg (@ARGV)
$toolArgs .= " -x c";
} elsif (($arg =~ /\.cpp$/) or ($arg =~ /\.cxx$/) or ($arg =~ /\.cc$/) or ($arg =~ /\.C$/)) {
$needCXXFLAGS = 1;
if ($HIP_COMPILE_CXX_AS_HIP eq '0' or $HIP_PLATFORM ne "amd") {
if ($HIP_COMPILE_CXX_AS_HIP eq '0' or $HIP_PLATFORM ne "amd" or $hasOMPTargets eq 1) {
$hasCXX = 1;
} elsif ($HIP_PLATFORM eq "amd") {
$hasHIP = 1;
projects/hip/include/hip/hip_runtime_api.h
+879 -204
مشاهده پرونده
تفاوت فایلی نمایش داده نمی شود زیرا این فایل بسیار بزرگ است Diff را بارگزاری کن
projects/hip/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp
+11 -11
مشاهده پرونده
@@ -38,9 +38,9 @@ THE SOFTWARE.
__global__ void EmptyKernel() { }
void print_timing(std::string test, const std::array<float, TOTAL_RUN_COUNT> &results, int batch = 1) {
float total_us = 0.0f, mean_us = 0.0f, stddev_us = 0.0f;
// skip warm-up runs
auto start_iter = std::next(results.begin(), WARMUP_RUN_COUNT);
auto end_iter = results.end();
@@ -48,7 +48,7 @@ void print_timing(std::string test, const std::array<float, TOTAL_RUN_COUNT> &re
// mean
std::for_each(start_iter, end_iter, [&](const float &run_ms) {
total_us += (run_ms * 1000) / batch;
});
});
mean_us = total_us / TIMING_RUN_COUNT;
// stddev
@@ -63,18 +63,18 @@ void print_timing(std::string test, const std::array<float, TOTAL_RUN_COUNT> &re
printf("\n %s: %.1f us, std: %.1f us\n", test.c_str(), mean_us, stddev_us);
}
int main() {
int main() {
hipStream_t stream0 = 0;
hipDevice_t device;
hipDeviceGet(&device, 0);
hipCtx_t context;
hipCtxCreate(&context, 0, device);
hipCtx_t context;
hipCtxCreate(&context, 0, device);
hipModule_t module;
hipFunction_t function;
hipModuleLoad(&module, FILE_NAME);
hipModuleGetFunction(&function, module, KERNEL_NAME);
void* params = nullptr;
std::array<float, TOTAL_RUN_COUNT> results;
hipEvent_t start, stop;
hipEventCreate(&start);
@@ -83,7 +83,7 @@ int main() {
/************************************************************************************/
/* HIP kernel launch enqueue rate: */
/* Measure time taken to enqueue a kernel on the GPU */
/************************************************************************************/
/************************************************************************************/
// Timing hipModuleLaunchKernel
for (auto i = 0; i < TOTAL_RUN_COUNT; ++i) {
@@ -104,8 +104,8 @@ int main() {
print_timing("hipLaunchKernelGGL enqueue rate", results);
/***********************************************************************************/
/* Single dispatch execution latency using HIP events: */
/* Measures latency to start & finish executing a kernel with GPU-scope visibility */
/* Single dispatch execution latency using HIP events: */
/* Measures latency to start & finish executing a kernel with GPU-scope visibility */
/***********************************************************************************/
//Timing around the dispatch
@@ -120,7 +120,7 @@ int main() {
/*********************************************************************************/
/* Batch dispatch execution latency using HIP events: */
/* Measures latency to start & finish executing each dispatch in a batch */
/* Measures latency to start & finish executing each dispatch in a batch */
/*********************************************************************************/
for (auto i = 0; i < TOTAL_RUN_COUNT; ++i) {
projects/hip/samples/2_Cookbook/13_occupancy/occupancy.cpp
+22 -22
مشاهده پرونده
@@ -29,23 +29,23 @@ THE SOFTWARE.
// Device (Kernel) function
__global__ void multiply(float* C, float* A, float* B, int N){
int tx = blockDim.x*blockIdx.x+threadIdx.x;
if (tx < N){
C[tx] = A[tx] * B[tx];
}
}
// CPU implementation
void multiplyCPU(float* C, float* A, float* B, int N){
for(unsigned int i=0; i<N; i++){
C[i] = A[i] * B[i];
for(unsigned int i=0; i<N; i++){
C[i] = A[i] * B[i];
}
}
void launchKernel(float* C, float* A, float* B, bool manual){
hipDeviceProp_t devProp;
HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
@@ -59,9 +59,9 @@ void launchKernel(float* C, float* A, float* B, bool manual){
int mingridSize = 0;
int gridSize = 0;
int blockSize = 0;
if (manual){
blockSize = threadsperblock;
blockSize = threadsperblock;
gridSize = blocks;
std::cout << std::endl << "Manual Configuration with block size " << blockSize << std::endl;
}
@@ -69,15 +69,15 @@ void launchKernel(float* C, float* A, float* B, bool manual){
HIP_CHECK(hipOccupancyMaxPotentialBlockSize(&mingridSize, &blockSize, multiply, 0, 0));
std::cout << std::endl << "Automatic Configuation based on hipOccupancyMaxPotentialBlockSize " << std::endl;
std::cout << "Suggested blocksize is " << blockSize << ", Minimum gridsize is " << mingridSize << std::endl;
gridSize = (NUM/blockSize)+1;
gridSize = (NUM/blockSize)+1;
}
// Record the start event
HIP_CHECK(hipEventRecord(start, NULL));
HIP_CHECK(hipEventRecord(start, NULL));
// Launching the Kernel from Host
hipLaunchKernelGGL(multiply, dim3(gridSize), dim3(blockSize), 0, 0, C, A, B, NUM);
// Record the stop event
HIP_CHECK(hipEventRecord(stop, NULL));
HIP_CHECK(hipEventSynchronize(stop));
@@ -88,7 +88,7 @@ void launchKernel(float* C, float* A, float* B, bool manual){
//Calculate Occupancy
int numBlock = 0;
HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, multiply, blockSize, 0));
if(devProp.maxThreadsPerMultiProcessor){
std::cout << "Theoretical Occupancy is " << (double)numBlock* blockSize/devProp.maxThreadsPerMultiProcessor * 100 << "%" << std::endl;
}
@@ -106,26 +106,26 @@ int main() {
C0 = (float *)malloc(NUM * sizeof(float));
C1 = (float *)malloc(NUM * sizeof(float));
cpuC = (float *)malloc(NUM * sizeof(float));
for(i=0; i< NUM; i++){
A[i] = i;
B[i] = i;
}
// allocate the memory on the device side
// allocate the memory on the device side
HIP_CHECK(hipMalloc((void**)&Ad, NUM * sizeof(float)));
HIP_CHECK(hipMalloc((void**)&Bd, NUM * sizeof(float)));
HIP_CHECK(hipMalloc((void**)&C0d, NUM * sizeof(float)));
HIP_CHECK(hipMalloc((void**)&C1d, NUM * sizeof(float)));
// Memory transfer from host to device
HIP_CHECK(hipMemcpy(Ad,A,NUM * sizeof(float), hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Bd,B,NUM * sizeof(float), hipMemcpyHostToDevice));
//Kernel launch with manual/default block size
launchKernel(C0d, Ad, Bd, 1);
//Kernel launch with the block size suggested by hipOccupancyMaxPotentialBlockSize
//Kernel launch with the block size suggested by hipOccupancyMaxPotentialBlockSize
launchKernel(C1d, Ad, Bd, 0);
// Memory transfer from device to host
@@ -137,26 +137,26 @@ int main() {
//verify the results
double eps = 1.0E-6;
for (i = 0; i < NUM; i++) {
if (std::abs(C0[i] - cpuC[i]) > eps) {
errors++;
}
}
if (errors != 0){
printf("\nManual Test FAILED: %d errors\n", errors);
errors=0;
} else {
printf("\nManual Test PASSED!\n");
}
for (i = 0; i < NUM; i++) {
if (std::abs(C1[i] - cpuC[i]) > eps) {
errors++;
}
}
if (errors != 0){
printf("\n Automatic Test FAILED: %d errors\n", errors);
} else {
projects/hip/tests/catch/CMakeLists.txt
+33
مشاهده پرونده
@@ -54,6 +54,39 @@ if(POLICY CMP0037)
cmake_policy(SET CMP0037 OLD)
endif()
# Turn off CMAKE_HIP_ARCHITECTURES Feature if cmake version is 3.21+
if(CMAKE_VERSION VERSION_GREATER_EQUAL 3.21.0)
set(CMAKE_HIP_ARCHITECTURES OFF)
endif()
message(STATUS "CMAKE HIP ARCHITECTURES: ${CMAKE_HIP_ARCHITECTURES}")
# Identify the GPU Targets.
# This is done due to limitation of rocm_agent_enumerator
# While building test parallelly, rocm_agent_enumerator can fail and give out an empty target
# That results in hipcc building the test for gfx803 (the default target)
if(NOT DEFINED OFFLOAD_ARCH_STR AND EXISTS "${ROCM_PATH}/bin/rocm_agent_enumerator"
AND HIP_PLATFORM STREQUAL "amd" AND UNIX)
execute_process(COMMAND ${ROCM_PATH}/bin/rocm_agent_enumerator OUTPUT_VARIABLE HIP_GPU_ARCH
RESULT_VARIABLE ROCM_AGENT_ENUM_RESULT)
# Trim out gfx000
string(REPLACE "gfx000\n" "" HIP_GPU_ARCH ${HIP_GPU_ARCH})
string(LENGTH ${HIP_GPU_ARCH} HIP_GPU_ARCH_LEN)
# If string has more gfx target except gfx000
if(${HIP_GPU_ARCH_LEN} GREATER_EQUAL 1)
string(REGEX REPLACE "\n" ";" HIP_GPU_ARCH_LIST "${HIP_GPU_ARCH}")
set(OFFLOAD_ARCH_STR "")
foreach(_hip_gpu_arch ${HIP_GPU_ARCH_LIST})
set(OFFLOAD_ARCH_STR " ${OFFLOAD_ARCH_STR} --offload-arch=${_hip_gpu_arch} ")
endforeach()
message(STATUS "Using offload arch string: ${OFFLOAD_ARCH_STR}")
endif()
endif()
if(DEFINED OFFLOAD_ARCH_STR)
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${OFFLOAD_ARCH_STR} ")
endif()
# Use clang as host compiler with nvcc
if(HIP_COMPILER MATCHES "nvcc")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -ccbin clang")
projects/hip/tests/catch/hipTestMain/CMakeLists.txt
+21
مشاهده پرونده
@@ -1,3 +1,23 @@
# Copyright (c) 2021 Advanced Micro Devices, Inc. All Rights Reserved.
#
# Permission is hereby granted, free of charge, to any person obtaining a copy
# of this software and associated documentation files (the "Software"), to deal
# in the Software without restriction, including without limitation the rights
# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
# copies of the Software, and to permit persons to whom the Software is
# furnished to do so, subject to the following conditions:
#
# The above copyright notice and this permission notice shall be included in
# all copies or substantial portions of the Software.
#
# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
# THE SOFTWARE.
if(CMAKE_BUILD_TYPE MATCHES "^Debug$")
add_definitions(-DHT_LOG_ENABLE)
endif()
@@ -16,6 +36,7 @@ target_link_libraries(UnitTests PRIVATE UnitDeviceTests
OccupancyTest
DeviceTest
RTC
TextureTest
stdc++fs)
if(HIP_PLATFORM MATCHES "nvidia")
projects/hip/tests/catch/unit/CMakeLists.txt
+21
مشاهده پرونده
@@ -1,3 +1,23 @@
# Copyright (c) 2021 Advanced Micro Devices, Inc. All Rights Reserved.
#
# Permission is hereby granted, free of charge, to any person obtaining a copy
# of this software and associated documentation files (the "Software"), to deal
# in the Software without restriction, including without limitation the rights
# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
# copies of the Software, and to permit persons to whom the Software is
# furnished to do so, subject to the following conditions:
#
# The above copyright notice and this permission notice shall be included in
# all copies or substantial portions of the Software.
#
# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
# THE SOFTWARE.
add_subdirectory(memory)
add_subdirectory(deviceLib)
add_subdirectory(stream)
@@ -5,3 +25,4 @@ add_subdirectory(event)
add_subdirectory(occupancy)
add_subdirectory(device)
add_subdirectory(rtc)
add_subdirectory(texture)
projects/hip/tests/catch/unit/memory/CMakeLists.txt
+2
مشاهده پرونده
@@ -44,6 +44,7 @@ set(TEST_SRC
hipMemset3DFunctional.cc
hipMemset3DNegative.cc
hipMemset3DRegressMultiThread.cc
hipMallocManagedFlagsTst.cc
)
else()
set(TEST_SRC
@@ -88,6 +89,7 @@ set(TEST_SRC
hipMemset3DFunctional.cc
hipMemset3DNegative.cc
hipMemset3DRegressMultiThread.cc
hipMallocManagedFlagsTst.cc
)
endif()
# Create shared lib of all tests
projects/hip/tests/catch/unit/memory/hipMallocManagedFlagsTst.cc
+340
مشاهده پرونده
@@ -0,0 +1,340 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <atomic>
// Kernel function
__global__ void MallcMangdFlgTst(int n, float *x, float *y) {
int index = blockIdx.x * blockDim.x + threadIdx.x;
int stride = blockDim.x * gridDim.x;
for (int i = index; i < n; i += stride)
y[i] = x[i] * x[i];
}
// The following function prints info on attributes related to HMM
static int HmmAttrPrint() {
int managed = 0;
INFO("The following are the attribute values related to HMM for"
" device 0:\n");
HIP_CHECK(hipDeviceGetAttribute(&managed,
hipDeviceAttributeDirectManagedMemAccessFromHost, 0));
INFO("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed);
HIP_CHECK(hipDeviceGetAttribute(&managed,
hipDeviceAttributeConcurrentManagedAccess, 0));
INFO("hipDeviceAttributeConcurrentManagedAccess: " << managed);
HIP_CHECK(hipDeviceGetAttribute(&managed,
hipDeviceAttributePageableMemoryAccess, 0));
INFO("hipDeviceAttributePageableMemoryAccess: " << managed);
HIP_CHECK(hipDeviceGetAttribute(&managed,
hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0));
INFO("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:"
<< managed);
HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
0));
INFO("hipDeviceAttributeManagedMemory: " << managed);
return managed;
}
// The following section tests working of hipMallocManaged with flag parameters
TEST_CASE("Unit_hipMallocManaged_FlgParam") {
int managed = HmmAttrPrint();
if (managed == 1) {
std::atomic<int> DataMismatch{0};
bool IfTestPassed = true;
float *HmmAG = NULL, *HmmAH1 = NULL, *HmmAH2 = NULL, INIT_VAL = 2.5;
int NumDevs = 0, NUM_ELMS = 4096;
HIP_CHECK(hipGetDeviceCount(&NumDevs));
float *Ad = NULL, *Ah = NULL;
Ah = new float[NUM_ELMS];
// Testing hipMemAttachGlobal Flag
HIP_CHECK(hipMallocManaged(&HmmAG, NUM_ELMS * sizeof(float),
hipMemAttachGlobal));
// Initializing HmmAG memory
for (int i = 0; i < NUM_ELMS; i++) {
HmmAG[i] = INIT_VAL;
Ah[i] = 0;
}
int blockSize = 256;
int numBlocks = (NUM_ELMS + blockSize - 1) / blockSize;
dim3 dimGrid(numBlocks, 1, 1);
dim3 dimBlock(blockSize, 1, 1);
hipStream_t strm;
for (int i = 0; i < NumDevs; i++) {
HIP_CHECK(hipSetDevice(i));
HIP_CHECK(hipStreamCreate(&strm));
HIP_CHECK(hipMalloc(&Ad, NUM_ELMS * sizeof(float)));
HIP_CHECK(hipMemset(Ad, 0, NUM_ELMS * sizeof(float)));
MallcMangdFlgTst<<<dimGrid, dimBlock, 0, strm>>>(NUM_ELMS, HmmAG, Ad);
HIP_CHECK(hipStreamSynchronize(strm));
HIP_CHECK(hipMemcpy(Ah, Ad, NUM_ELMS * sizeof(float),
hipMemcpyDeviceToHost));
for (int j = 0; j < NUM_ELMS; ++j) {
if (Ah[j] != (INIT_VAL * INIT_VAL)) {
DataMismatch++;
}
}
if (DataMismatch != 0) {
WARN("Data Mismatch observed when kernel launched on");
WARN(" device: " << i);
IfTestPassed = false;
}
DataMismatch = 0;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipStreamDestroy(strm));
}
delete[] Ah;
HIP_CHECK(hipFree(HmmAG));
DataMismatch = 0;
HIP_CHECK(hipMallocManaged(&HmmAH1, NUM_ELMS * sizeof(float),
hipMemAttachHost));
HIP_CHECK(hipMallocManaged(&HmmAH2, NUM_ELMS * sizeof(float),
hipMemAttachHost));
// Initializing HmmAH memory
for (int i = 0; i < NUM_ELMS; i++) {
HmmAH1[i] = INIT_VAL;
HmmAH2[i] = 0;
}
for (int i = 0; i < NumDevs; i++) {
HIP_CHECK(hipSetDevice(i));
HIP_CHECK(hipStreamCreate(&strm));
HIP_CHECK(hipMemset(HmmAH2, 0, NUM_ELMS * sizeof(float)));
MallcMangdFlgTst<<<dimGrid, dimBlock, 0, strm>>>(NUM_ELMS,
HmmAH1, HmmAH2);
HIP_CHECK(hipStreamSynchronize(strm));
for (int j = 0; j < NUM_ELMS; ++j) {
if (HmmAH2[j] != (INIT_VAL * INIT_VAL)) {
DataMismatch++;
}
}
if (DataMismatch != 0) {
WARN("Data Mismatch observed when kernel launched on");
WARN(" device: " << i);
IfTestPassed = false;
}
HIP_CHECK(hipStreamDestroy(strm));
}
HIP_CHECK(hipFree(HmmAH1));
HIP_CHECK(hipFree(HmmAH2));
REQUIRE(IfTestPassed);
} else {
SUCCEED("Gpu doesnt support HMM! Hence skipping the test with PASS result");
}
}
// The following function tests Memory access allocated using hipMallocManaged
// in multiple streams
TEST_CASE("Unit_hipMallocManaged_AccessMultiStream") {
int managed = HmmAttrPrint();
if (managed == 1) {
std::atomic<int> DataMismatch{0};
bool IfTestPassed = true;
float *HmmAG = NULL, *HmmAH1 = NULL, *HmmAH2 = NULL, INIT_VAL = 2.5;
int NumStrms = 0, MultiDevice = 0, NUM_ELMS = 4096;
HIP_CHECK(hipGetDeviceCount(&MultiDevice));
if (MultiDevice >= 2) {
HIP_CHECK(hipGetDeviceCount(&NumStrms));
} else {
NumStrms = 4;
}
hipStream_t **Stream = new hipStream_t*[NumStrms];
for (int i = 0; i < NumStrms; ++i) {
Stream[i] = reinterpret_cast<hipStream_t*>(malloc(sizeof(hipStream_t)));
}
float *Ad = NULL, *Ah = NULL;
Ah = new float[NUM_ELMS];
for (int i = 0; i < NumStrms; ++i) {
if (MultiDevice >= 2) {
HIP_CHECK(hipSetDevice(i));
}
HIP_CHECK(hipStreamCreate(Stream[i]));
}
HIP_CHECK(hipSetDevice(0));
// Testing hipMemAttachGlobal Flag
HIP_CHECK(hipMallocManaged(&HmmAG, NUM_ELMS * sizeof(float),
hipMemAttachGlobal));
// Initializing HmmAG memory
for (int i = 0; i < NUM_ELMS; i++) {
HmmAG[i] = INIT_VAL;
Ah[i] = 0;
}
int blockSize = 256;
int numBlocks = (NUM_ELMS + blockSize - 1) / blockSize;
dim3 dimGrid(numBlocks, 1, 1);
dim3 dimBlock(blockSize, 1, 1);
for (int i = 0; i < NumStrms; i++) {
if (MultiDevice >= 2) {
HIP_CHECK(hipSetDevice(i));
}
HIP_CHECK(hipMalloc(&Ad, NUM_ELMS * sizeof(float)));
HIP_CHECK(hipMemset(Ad, 0, NUM_ELMS * sizeof(float)));
MallcMangdFlgTst<<<dimGrid, dimBlock, 0, *(Stream[i])>>>(NUM_ELMS,
HmmAG, Ad);
HIP_CHECK(hipStreamSynchronize(*(Stream[i])));
// Validating the results
HIP_CHECK(hipMemcpy(Ah, Ad, NUM_ELMS * sizeof(float),
hipMemcpyDeviceToHost));
for (int j = 0; j < NUM_ELMS; ++j) {
if (Ah[j] != (INIT_VAL * INIT_VAL)) {
DataMismatch++;
}
}
if (DataMismatch != 0) {
WARN("Data Mismatch observed when kernel launched on");
WARN(" device: " << i);
IfTestPassed = false;
}
DataMismatch = 0;
HIP_CHECK(hipFree(Ad));
}
delete[] Ah;
HIP_CHECK(hipFree(HmmAG));
DataMismatch = 0;
HIP_CHECK(hipMallocManaged(&HmmAH1, NUM_ELMS * sizeof(float),
hipMemAttachHost));
HIP_CHECK(hipMallocManaged(&HmmAH2, NUM_ELMS * sizeof(float),
hipMemAttachHost));
// Initializing HmmAH memory
for (int i = 0; i < NUM_ELMS; i++) {
HmmAH1[i] = INIT_VAL;
HmmAH2[i] = 0;
}
for (int i = 0; i < NumStrms; i++) {
if (MultiDevice >= 2) {
HIP_CHECK(hipSetDevice(i));
}
HIP_CHECK(hipMemset(HmmAH2, 0, NUM_ELMS * sizeof(float)));
MallcMangdFlgTst<<<dimGrid, dimBlock, 0, *(Stream[i])>>>(NUM_ELMS,
HmmAH1, HmmAH2);
HIP_CHECK(hipStreamSynchronize(*(Stream[i])));
for (int j = 0; j < NUM_ELMS; ++j) {
if (HmmAH2[j] != (INIT_VAL * INIT_VAL)) {
DataMismatch++;
break;
}
}
if (DataMismatch != 0) {
WARN("Data Mismatch observed when kernel launched on");
WARN(" device: " << i);
IfTestPassed = false;
}
}
HIP_CHECK(hipFree(HmmAH1));
HIP_CHECK(hipFree(HmmAH2));
for (int i = 0; i < NumStrms; ++i) {
HIP_CHECK(hipStreamDestroy(*(Stream[i])));
}
REQUIRE(IfTestPassed);
} else {
SUCCEED("Gpu doesnt support HMM! Hence skipping the test with PASS result");
}
}
TEST_CASE("Unit_hipMallocManaged_ExtremeSizes") {
int managed = HmmAttrPrint();
if (managed == 1) {
bool IfTestPassed = true;
hipError_t err;
void *Hmm = NULL;
size_t totalDevMem = 0, freeDevMem = 0;
int NumDevs = 0;
HIP_CHECK(hipGetDeviceCount(&NumDevs));
// Testing allocation of extreme and unusual mem values
for (int i = 0; i < NumDevs; i++) {
HIP_CHECK(hipSetDevice(i));
HIP_CHECK(hipMemGetInfo(&freeDevMem, &totalDevMem));
err = hipMallocManaged(&Hmm, 1, hipMemAttachGlobal);
if (hipSuccess == err) {
HIP_CHECK(hipFree(Hmm));
} else {
WARN("Observed error while allocating memory on GPU: " << i);
WARN(" size 1 with");
WARN(" hipMallocManaged() api with flag 'hipMemAttachGlobal'\n");
WARN("Error: " << hipGetErrorString(err));
IfTestPassed = false;
}
err = hipMallocManaged(&Hmm, freeDevMem, hipMemAttachGlobal);
if (hipSuccess == err) {
HIP_CHECK(hipFree(Hmm));
} else {
WARN("Observed error while allocating max free memory on GPU: " << i);
WARN(" with hipMallocManaged() api with flag 'hipMemAttachGlobal'\n");
WARN("Error: " << hipGetErrorString(err));
IfTestPassed = false;
}
err = hipMallocManaged(&Hmm, (freeDevMem - 1), hipMemAttachGlobal);
if (hipSuccess == err) {
HIP_CHECK(hipFree(Hmm));
} else {
WARN("Observed error while allocating max (free - 1) memory on ");
WARN("GPU: " << i);
WARN(" using hipMallocManaged() api with flag 'hipMemAttachGlobal'\n");
WARN("Error: " << hipGetErrorString(err));
IfTestPassed = false;
}
err = hipMallocManaged(&Hmm, 1, hipMemAttachHost);
if (hipSuccess == err) {
HIP_CHECK(hipFree(Hmm));
} else {
WARN("Observed error while allocating memory size 1 on GPU: " << i);
WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n");
WARN("Error: " << hipGetErrorString(err));
IfTestPassed = false;
}
err = hipMallocManaged(&Hmm, freeDevMem, hipMemAttachHost);
if (hipSuccess == err) {
HIP_CHECK(hipFree(Hmm));
} else {
WARN("Observed error while allocating max free memory on GPU: " << i);
WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n");
WARN("Error: " << hipGetErrorString(err));
IfTestPassed = false;
}
err = hipMallocManaged(&Hmm, (freeDevMem - 1), hipMemAttachHost);
if (hipSuccess == err) {
HIP_CHECK(hipFree(Hmm));
} else {
WARN("Observed error while allocating max (freeDevMem - 1) memory"
" on GPU: " << i);
WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n");
WARN("Error: " << hipGetErrorString(err));
IfTestPassed = false;
}
}
REQUIRE(IfTestPassed);
} else {
SUCCEED("Gpu doesnt support HMM! Hence skipping the test with PASS result");
}
}
projects/hip/tests/catch/unit/memory/hipMemset2DAsyncMultiThreadAndKernel.cc
+1 -5
مشاهده پرونده
@@ -112,9 +112,6 @@ TEST_CASE("Unit_hipMemset2DAsync_WithKernel") {
* hipMemSet2DAsync execution in multiple threads.
*/
TEST_CASE("Unit_hipMemset2DAsync_MultiThread") {
constexpr auto N = 4 * 1024 * 1024;
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
constexpr auto memPerThread = 200;
constexpr int memsetval = 0x22;
char *A_d, *A_h, *B_d, *B_h, *C_d;
@@ -122,11 +119,10 @@ TEST_CASE("Unit_hipMemset2DAsync_MultiThread") {
size_t width = NUM_W * sizeof(char);
size_t sizeElements = width * NUM_H;
size_t elements = NUM_W * NUM_H;
unsigned blocks{};
int validateCount{};
hipStream_t stream;
blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
auto thread_count = HipTest::getHostThreadCount(memPerThread, NUM_THREADS);
if (thread_count == 0) {
WARN("Resources not available for thread creation");
projects/hip/tests/catch/unit/stream/hipStreamWithCUMask.cc
+1 -2
مشاهده پرونده
@@ -95,10 +95,9 @@ void createDisabledCUMask(std::vector<uint32_t> *pdisabledCUMask,
}
}
void Callback(hipStream_t stream, hipError_t status,
void Callback(hipStream_t, hipError_t status,
void* userData) {
isPassed = true;
stream = 0;
HIP_CHECK(status);
REQUIRE(userData == nullptr);
for (size_t i = 0; i < N; i++) {
projects/hip/tests/catch/unit/texture/CMakeLists.txt
+33
مشاهده پرونده
@@ -0,0 +1,33 @@
# Copyright (c) 2021 Advanced Micro Devices, Inc. All Rights Reserved.
#
# Permission is hereby granted, free of charge, to any person obtaining a copy
# of this software and associated documentation files (the "Software"), to deal
# in the Software without restriction, including without limitation the rights
# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
# copies of the Software, and to permit persons to whom the Software is
# furnished to do so, subject to the following conditions:
#
# The above copyright notice and this permission notice shall be included in
# all copies or substantial portions of the Software.
#
# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
# THE SOFTWARE.
# Common Tests - Test independent of all platforms
set(TEST_SRC
hipCreateTextureObject_ArgValidation.cc
hipCreateTextureObject_Linear.cc
hipCreateTextureObject_Pitch2D.cc
hipCreateTextureObject_Array.cc
)
# Create shared lib of all tests
add_library(TextureTest SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
# Add dependency on build_tests to build it on this custom target
add_dependencies(build_tests TextureTest)
projects/hip/tests/catch/unit/texture/hipCreateTextureObject_ArgValidation.cc
+81
مشاهده پرونده
@@ -0,0 +1,81 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#define N 512
/*
* Validate argument list of texture object api.
*/
TEST_CASE("Unit_hipCreateTextureObject_ArgValidation") {
float *texBuf;
hipError_t ret;
constexpr int xsize = 32;
hipResourceDesc resDesc;
hipTextureDesc texDesc;
hipTextureObject_t texObj;
/** Initialization */
HIP_CHECK(hipMalloc(&texBuf, N * sizeof(float)));
// Populate resource descriptor
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = hipResourceTypeLinear;
resDesc.res.linear.devPtr = texBuf;
resDesc.res.linear.desc = hipCreateChannelDesc(xsize, 0, 0, 0,
hipChannelFormatKindFloat);
resDesc.res.linear.sizeInBytes = N * sizeof(float);
// Populate texture descriptor
memset(&texDesc, 0, sizeof(texDesc));
texDesc.readMode = hipReadModeElementType;
/** Sections */
SECTION("TextureObject as nullptr") {
ret = hipCreateTextureObject(nullptr, &resDesc, &texDesc, nullptr);
REQUIRE(ret != hipSuccess);
}
SECTION("Resouce Descriptor as nullptr") {
ret = hipCreateTextureObject(&texObj, nullptr, &texDesc, nullptr);
REQUIRE(ret != hipSuccess);
}
SECTION("Texture Descriptor as nullptr") {
if ((TestContext::get()).isAmd()) {
ret = hipCreateTextureObject(&texObj, &resDesc, nullptr, nullptr);
REQUIRE(ret != hipSuccess);
} else {
// API expected to return failure. Test skipped
// on nvidia as api returns success and would lead
// to unexpected behavior with app.
WARN("Texture Desc(nullptr) skipped on nvidia");
}
}
SECTION("Destroy TextureObject with nullptr") {
ret = hipDestroyTextureObject((hipTextureObject_t)nullptr);
// api to return success and no crash seen.
REQUIRE(ret == hipSuccess);
}
/** De-Initialization */
HIP_CHECK(hipFree(texBuf));
}
projects/hip/tests/catch/unit/texture/hipCreateTextureObject_Array.cc
+67
مشاهده پرونده
@@ -0,0 +1,67 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
/*
* Validates Array Resource texture object with negative/functional tests.
*/
TEST_CASE("Unit_hipCreateTextureObject_ArrayResource") {
hipError_t ret;
hipResourceDesc resDesc;
hipTextureDesc texDesc;
hipTextureObject_t texObj;
/* set resource type as hipResourceTypeArray and array(nullptr) */
// Populate resource descriptor
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = hipResourceTypeArray;
resDesc.res.array.array = nullptr;
// Populate texture descriptor
memset(&texDesc, 0, sizeof(texDesc));
texDesc.readMode = hipReadModeElementType;
ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr);
REQUIRE(ret != hipSuccess);
}
/*
* Validates MipMappedArray Resource texture object
* with negative/functional tests.
*/
TEST_CASE("Unit_hipCreateTextureObject_MmArrayResource") {
hipError_t ret;
hipResourceDesc resDesc;
hipTextureDesc texDesc;
hipTextureObject_t texObj;
/* set resource type as hipResourceTypeMipmappedArray and mipmap(nullptr) */
// Populate resource descriptor
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = hipResourceTypeMipmappedArray;
resDesc.res.mipmap.mipmap = nullptr;
// Populate texture descriptor
memset(&texDesc, 0, sizeof(texDesc));
texDesc.readMode = hipReadModeElementType;
ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr);
REQUIRE(ret != hipSuccess);
}
projects/hip/tests/catch/unit/texture/hipCreateTextureObject_Linear.cc
+133
مشاهده پرونده
@@ -0,0 +1,133 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#define UNALIGN_OFFSET 1
#define N 512
/*
* Validates Linear Resource texture object with negative/functional tests.
*/
TEST_CASE("Unit_hipCreateTextureObject_LinearResource") {
float *texBuf;
hipError_t ret;
constexpr int xsize = 32;
hipResourceDesc resDesc;
hipTextureDesc texDesc;
hipResourceViewDesc resViewDesc;
hipTextureObject_t texObj;
hipDeviceProp_t devProp;
/** Initialization */
HIP_CHECK(hipMalloc(&texBuf, N * sizeof(float)));
HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
memset(&resDesc, 0, sizeof(resDesc));
memset(&texDesc, 0, sizeof(texDesc));
resDesc.resType = hipResourceTypeLinear;
/** Sections */
SECTION("hipResourceTypeLinear and devPtr(nullptr)") {
// Populate resource descriptor
resDesc.res.linear.devPtr = nullptr;
resDesc.res.linear.desc = hipCreateChannelDesc(xsize, 0, 0, 0,
hipChannelFormatKindFloat);
resDesc.res.linear.sizeInBytes = N * sizeof(float);
// Populate texture descriptor
texDesc.readMode = hipReadModeElementType;
ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr);
REQUIRE(ret != hipSuccess);
}
SECTION("hipResourceTypeLinear and sizeInBytes(0)") {
if ((TestContext::get()).isAmd()) {
// Populate resource descriptor
resDesc.res.linear.devPtr = texBuf;
resDesc.res.linear.desc = hipCreateChannelDesc(xsize, 0, 0, 0,
hipChannelFormatKindFloat);
resDesc.res.linear.sizeInBytes = 0;
// Populate texture descriptor
texDesc.readMode = hipReadModeElementType;
ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr);
REQUIRE(ret != hipSuccess);
} else {
// API expected to return failure. Test skipped
// on nvidia as api returns success and would lead
// to unexpected behavior with app.
WARN("Resource type Linear/sizeInBytes(0) skipped on nvidia");
}
}
SECTION("hipResourceTypeLinear and sizeInBytes(max(size_t))") {
// Populate resource descriptor
resDesc.res.linear.devPtr = texBuf;
resDesc.res.linear.desc = hipCreateChannelDesc(xsize, 0, 0, 0,
hipChannelFormatKindFloat);
resDesc.res.linear.sizeInBytes = std::numeric_limits<std::size_t>::max();
// Populate texture descriptor
texDesc.readMode = hipReadModeElementType;
ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr);
REQUIRE(ret != hipSuccess);
}
SECTION("hipResourceTypeLinear and valid resource view descriptor") {
#if HT_AMD
// Populate resource descriptor
resDesc.res.linear.devPtr = texBuf;
resDesc.res.linear.desc = hipCreateChannelDesc(xsize, 0, 0, 0,
hipChannelFormatKindFloat);
resDesc.res.linear.sizeInBytes = N * sizeof(float);
// Populate texture descriptor
texDesc.readMode = hipReadModeElementType;
// Populate resourceview descriptor
memset(&resViewDesc, 0, sizeof(resViewDesc));
resViewDesc.format = hipResViewFormatFloat1;
resViewDesc.width = N * sizeof(float);
ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, &resViewDesc);
REQUIRE(ret != hipSuccess);
#else
// API expected to return error according to cuda documentation.
WARN("Resource view descriptor test skipped on nvidia");
#endif
}
SECTION("hipResourceTypeLinear and devicePtr un-aligned") {
if (devProp.textureAlignment > UNALIGN_OFFSET) {
// Populate resource descriptor
resDesc.res.linear.devPtr = reinterpret_cast<char *>(texBuf)
+ UNALIGN_OFFSET;
resDesc.res.linear.desc = hipCreateChannelDesc(xsize, 0, 0, 0,
hipChannelFormatKindFloat);
resDesc.res.linear.sizeInBytes = N * sizeof(float);
// Populate texture descriptor
texDesc.readMode = hipReadModeElementType;
ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr);
REQUIRE(ret != hipSuccess);
}
}
/** De-Initialization */
HIP_CHECK(hipFree(texBuf));
}
projects/hip/tests/catch/unit/texture/hipCreateTextureObject_Pitch2D.cc
+216
مشاهده پرونده
@@ -0,0 +1,216 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#define UNALIGN_OFFSET 1
#define SIZE_H 20
#define SIZE_W 30
#define N 512
/*
* Validates Pitch2D Resource texture object with negative and functional tests
*/
TEST_CASE("Unit_hipCreateTextureObject_Pitch2DResource") {
hipError_t ret;
hipResourceDesc resDesc;
hipTextureDesc texDesc;
hipTextureObject_t texObj;
hipDeviceProp_t devProp;
size_t devPitchA;
float *devPtrA;
/** Initialization */
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&devPtrA), &devPitchA,
SIZE_W*sizeof(float), SIZE_H));
HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
memset(&resDesc, 0, sizeof(resDesc));
memset(&texDesc, 0, sizeof(texDesc));
resDesc.resType = hipResourceTypePitch2D;
/** Sections */
SECTION("hipResourceTypePitch2D and devPtr(nullptr)") {
// Populate resource descriptor
resDesc.res.pitch2D.devPtr = nullptr;
resDesc.res.pitch2D.height = SIZE_H;
resDesc.res.pitch2D.width = SIZE_W;
resDesc.res.pitch2D.pitchInBytes = devPitchA;
resDesc.res.pitch2D.desc = hipCreateChannelDesc<float>();
// Populate texture descriptor
texDesc.readMode = hipReadModeElementType;
ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr);
REQUIRE(ret != hipSuccess);
}
SECTION("hipResourceTypePitch2D and devicePtr un-aligned") {
if (devProp.textureAlignment > UNALIGN_OFFSET) {
// Populate resource descriptor
resDesc.res.pitch2D.devPtr = reinterpret_cast<char *>(devPtrA)
+ UNALIGN_OFFSET;
resDesc.res.pitch2D.height = SIZE_H;
resDesc.res.pitch2D.width = SIZE_W;
resDesc.res.pitch2D.pitchInBytes = devPitchA;
resDesc.res.pitch2D.desc = hipCreateChannelDesc<float>();
// Populate texture descriptor
memset(&texDesc, 0, sizeof(texDesc));
texDesc.readMode = hipReadModeElementType;
ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr);
REQUIRE(ret != hipSuccess);
}
}
SECTION("hipResourceTypePitch2D and pitch(un-aligned)") {
// Populate resource descriptor
resDesc.res.pitch2D.devPtr = devPtrA;
resDesc.res.pitch2D.height = SIZE_H;
resDesc.res.pitch2D.width = SIZE_W;
resDesc.res.pitch2D.pitchInBytes = UNALIGN_OFFSET;
resDesc.res.pitch2D.desc = hipCreateChannelDesc<float>();
// Populate texture descriptor
texDesc.readMode = hipReadModeElementType;
ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr);
REQUIRE(ret != hipSuccess);
}
SECTION("hipResourceTypePitch2D and height(0)") {
if ((TestContext::get()).isAmd()) {
// Populate resource descriptor
resDesc.res.pitch2D.devPtr = devPtrA;
resDesc.res.pitch2D.height = 0;
resDesc.res.pitch2D.width = SIZE_W;
resDesc.res.pitch2D.pitchInBytes = devPitchA;
resDesc.res.pitch2D.desc = hipCreateChannelDesc<float>();
// Populate texture descriptor
texDesc.readMode = hipReadModeElementType;
ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr);
REQUIRE(ret != hipSuccess);
} else {
// Test expected to return error with height(0).
WARN("Resourcetype Pitch2D/height(0) skipped on nvidia");
}
}
SECTION("hipResourceTypePitch2D and height(0)/devptr(nullptr)") {
// Populate resource descriptor
resDesc.res.pitch2D.devPtr = nullptr;
resDesc.res.pitch2D.height = 0;
resDesc.res.pitch2D.width = SIZE_W;
resDesc.res.pitch2D.pitchInBytes = devPitchA;
resDesc.res.pitch2D.desc = hipCreateChannelDesc<float>();
// Populate texture descriptor
texDesc.readMode = hipReadModeElementType;
ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr);
REQUIRE(ret != hipSuccess);
}
SECTION("hipResourceTypePitch2D and height(max(size_t))") {
// Populate resource descriptor
resDesc.res.pitch2D.devPtr = devPtrA;
resDesc.res.pitch2D.height = std::numeric_limits<std::size_t>::max();
resDesc.res.pitch2D.width = SIZE_W;
resDesc.res.pitch2D.pitchInBytes = devPitchA;
resDesc.res.pitch2D.desc = hipCreateChannelDesc<float>();
// Populate texture descriptor
texDesc.readMode = hipReadModeElementType;
ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr);
REQUIRE(ret != hipSuccess);
}
SECTION("hipResourceTypePitch2D and width(0)") {
if ((TestContext::get()).isAmd()) {
// Populate resource descriptor
resDesc.resType = hipResourceTypePitch2D;
resDesc.res.pitch2D.devPtr = devPtrA;
resDesc.res.pitch2D.height = SIZE_H;
resDesc.res.pitch2D.width = 0;
resDesc.res.pitch2D.pitchInBytes = devPitchA;
resDesc.res.pitch2D.desc = hipCreateChannelDesc<float>();
// Populate texture descriptor
texDesc.readMode = hipReadModeElementType;
ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr);
REQUIRE(ret != hipSuccess);
} else {
// api expected to return failure when width(0) is passed.
WARN("ResourceType Pitch2D/width(0) skipped on nvidia");
}
}
SECTION("hipResourceTypePitch2D and width(0)/devPtr(nullptr)") {
// Populate resource descriptor
resDesc.res.pitch2D.devPtr = nullptr;
resDesc.res.pitch2D.height = SIZE_H;
resDesc.res.pitch2D.width = 0;
resDesc.res.pitch2D.pitchInBytes = devPitchA;
resDesc.res.pitch2D.desc = hipCreateChannelDesc<float>();
// Populate texture descriptor
texDesc.readMode = hipReadModeElementType;
ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr);
REQUIRE(ret != hipSuccess);
}
SECTION("hipResourceTypePitch2D and width(max(size_t))") {
// Populate resource descriptor
resDesc.res.pitch2D.devPtr = devPtrA;
resDesc.res.pitch2D.height = SIZE_H;
resDesc.res.pitch2D.width = std::numeric_limits<std::size_t>::max();
resDesc.res.pitch2D.pitchInBytes = devPitchA;
resDesc.res.pitch2D.desc = hipCreateChannelDesc<float>();
// Populate texture descriptor
texDesc.readMode = hipReadModeElementType;
ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr);
REQUIRE(ret != hipSuccess);
}
SECTION("hipResourceTypePitch2D and pitch(max(size_t))") {
// Populate resource descriptor
resDesc.res.pitch2D.devPtr = devPtrA;
resDesc.res.pitch2D.height = SIZE_H;
resDesc.res.pitch2D.width = SIZE_W;
resDesc.res.pitch2D.pitchInBytes = std::numeric_limits<std::size_t>::max();
resDesc.res.pitch2D.desc = hipCreateChannelDesc<float>();
// Populate texture descriptor
texDesc.readMode = hipReadModeElementType;
ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr);
REQUIRE(ret != hipSuccess);
}
/** De-Initialization */
HIP_CHECK(hipFree(devPtrA));
}
projects/hip/tests/hit/HIT.cmake
+39 -1
مشاهده پرونده
@@ -29,12 +29,50 @@ if (NOT ${BUILD_SHARED_LIBS})
endif()
message(STATUS "HIP runtime lib type - ${HIP_LIB_TYPE}")
message(STATUS "CMAKE_TESTING_TOOL: ${CMAKE_TESTING_TOOL}")
# Turn off CMAKE_HIP_ARCHITECTURES Feature if cmake version is 3.21+
if(CMAKE_VERSION VERSION_GREATER_EQUAL 3.21.0)
set(CMAKE_HIP_ARCHITECTURES OFF)
endif()
message(STATUS "CMAKE HIP ARCHITECTURES: ${CMAKE_HIP_ARCHITECTURES}")
# Identify the GPU Targets.
# This is done due to limitation of rocm_agent_enumerator
# While building test parallelly, rocm_agent_enumerator can fail and give out an empty target
# That results in hipcc building the test for gfx803 (the default target)
if(NOT DEFINED OFFLOAD_ARCH_STR AND EXISTS "${ROCM_PATH}/bin/rocm_agent_enumerator"
AND HIP_PLATFORM STREQUAL "amd" AND UNIX)
execute_process(COMMAND ${ROCM_PATH}/bin/rocm_agent_enumerator OUTPUT_VARIABLE HIP_GPU_ARCH
RESULT_VARIABLE ROCM_AGENT_ENUM_RESULT)
message(STATUS "ROCm Agent Enumurator Result: ${ROCM_AGENT_ENUM_RESULT}")
# Trim out gfx000
string(REPLACE "gfx000\n" "" HIP_GPU_ARCH ${HIP_GPU_ARCH})
string(LENGTH ${HIP_GPU_ARCH} HIP_GPU_ARCH_LEN)
# If string has more gfx target except gfx000
if(${HIP_GPU_ARCH_LEN} GREATER_EQUAL 1)
string(REGEX REPLACE "\n" ";" HIP_GPU_ARCH_LIST "${HIP_GPU_ARCH}")
set(OFFLOAD_ARCH_STR "")
foreach(_hip_gpu_arch ${HIP_GPU_ARCH_LIST})
set(OFFLOAD_ARCH_STR " ${OFFLOAD_ARCH_STR} --offload-arch=${_hip_gpu_arch} ")
endforeach()
message(STATUS "Using offload arch string: ${OFFLOAD_ARCH_STR}")
endif()
else()
message(STATUS "ROCm Agent Enumurator Not Found")
endif()
#-------------------------------------------------------------------------------
# Helper macro to parse BUILD instructions
macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _clang_options _nvcc_options _link_options _exclude_platforms _exclude_runtime _exclude_compiler _exclude_lib_type _depends _dir)
set(${_target})
set(${_sources})
set(${_hipcc_options})
if(DEFINED OFFLOAD_ARCH_STR)
set(${_hipcc_options} "${OFFLOAD_ARCH_STR}")
else()
set(${_hipcc_options})
endif()
set(${_clang_options})
set(${_nvcc_options})
set(${_link_options})
projects/hip/tests/src/Negative/memory/hipMemcpyFromSymbol.cpp
+3 -3
مشاهده پرونده
@@ -36,11 +36,11 @@ int main(){
e = hipMemcpyFromSymbol(S, HIP_SYMBOL(Sd), SIZE, 0, hipMemcpyDeviceToHost);
HIPASSERT(e==hipErrorInvalidSymbol);
e = hipMemcpyFromSymbol(S, NULL, SIZE, 0, hipMemcpyDeviceToHost);
HIPASSERT(e==hipErrorInvalidSymbol);
HIPCHECK(hipFree(Sd));
HIPCHECK(hipFree(Sd));
passed();
}
projects/hip/tests/src/Negative/memory/hipMemcpyFromSymbolAsync.cpp
+3 -3
مشاهده پرونده
@@ -39,11 +39,11 @@ int main(){
e = hipMemcpyFromSymbolAsync(S, HIP_SYMBOL(Sd), SIZE, 0, hipMemcpyDeviceToHost, stream);
HIPASSERT(e==hipErrorInvalidSymbol);
e = hipMemcpyFromSymbolAsync(S, NULL, SIZE, 0, hipMemcpyDeviceToHost, stream);
HIPASSERT(e==hipErrorInvalidSymbol);
HIPCHECK(hipFree(Sd));
HIPCHECK(hipFree(Sd));
passed();
}
projects/hip/tests/src/Negative/memory/hipMemcpyToSymbol.cpp
+3 -3
مشاهده پرونده
@@ -36,11 +36,11 @@ int main(){
e = hipMemcpyToSymbol(HIP_SYMBOL(Sd), S, SIZE, 0, hipMemcpyHostToDevice);
HIPASSERT(e==hipErrorInvalidSymbol);
e = hipMemcpyToSymbol(NULL, S, SIZE, 0, hipMemcpyHostToDevice);
HIPASSERT(e==hipErrorInvalidSymbol);
HIPCHECK(hipFree(Sd));
HIPCHECK(hipFree(Sd));
passed();
}
projects/hip/tests/src/Negative/memory/hipMemcpyToSymbolAsync.cpp
+5 -5
مشاهده پرونده
@@ -31,7 +31,7 @@ int main(){
void *Sd;
hipError_t e;
char S[SIZE]="This is not a device symbol";
HIPCHECK(hipMalloc(&Sd,SIZE));
hipStream_t stream;
@@ -39,11 +39,11 @@ int main(){
e = hipMemcpyToSymbolAsync(HIP_SYMBOL(Sd), S, SIZE, 0, hipMemcpyHostToDevice, stream);
HIPASSERT(e==hipErrorInvalidSymbol);
e = hipMemcpyToSymbolAsync(NULL, S, SIZE, 0, hipMemcpyHostToDevice, stream);
HIPASSERT(e==hipErrorInvalidSymbol);
HIPCHECK(hipFree(Sd));
HIPCHECK(hipFree(Sd));
passed();
}
projects/hip/tests/src/Negative/memory/hipMemory.cpp
+2 -2
مشاهده پرونده
@@ -1,4 +1,4 @@
/*
/*
Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
@@ -32,7 +32,7 @@ int main(){
e = hipMemcpy(0, str, SIZE, hipMemcpyHostToDevice);
HIPASSERT(e==hipErrorInvalidValue);
e = hipMemcpy(NULL, str, SIZE, hipMemcpyHostToDevice);
HIPASSERT(e==hipErrorInvalidValue);
projects/hip/tests/src/compiler/hipClassKernel.cpp
+8 -8
مشاهده پرونده
@@ -45,7 +45,7 @@ void HipClassTests::TestForOverride(void){
0,
0,
result_ecd);
HipClassTests::VerifyResult(result_ech,result_ecd);
HipClassTests::FreeMem(result_ech,result_ecd);
}
@@ -70,13 +70,13 @@ void HipClassTests::TestForOverload(void){
0,
0,
result_ecd);
HipClassTests::VerifyResult(result_ech,result_ecd);
HipClassTests::FreeMem(result_ech,result_ecd);
}
#endif
#ifdef ENABLE_FRIEND_TEST
#ifdef ENABLE_FRIEND_TEST
// check for friend
__global__ void
friendClassKernel(bool* result_ecd){
@@ -106,7 +106,7 @@ void HipClassTests::TestForEmptyClass(void){
0,
0,
result_ecd);
HipClassTests::VerifyResult(result_ech,result_ecd);
HipClassTests::FreeMem(result_ech,result_ecd);
}
@@ -157,7 +157,7 @@ void HipClassTests::TestForClassSize(void){
0,
0,
result_ecd);
HipClassTests::VerifyResult(result_ech,result_ecd);
HipClassTests::FreeMem(result_ech,result_ecd);
}
@@ -217,7 +217,7 @@ void HipClassTests::TestForPassByValue(void){
HipClassTests::VerifyResult(result_ech,result_ecd);
HipClassTests::FreeMem(result_ech,result_ecd);
}
// check obj created with hipMalloc
__global__ void
mallocObjKernel(testPassByValue *obj, bool* result_ecd) {
@@ -292,7 +292,7 @@ bool* HipClassTests::AllocateHostMemory(void){
}
bool* HipClassTests::AllocateDeviceMemory(void){
bool* result_ecd;
bool* result_ecd;
HIPCHECK(hipMalloc(&result_ecd,
NBOOL));
HIPCHECK(hipMemset(result_ecd,
@@ -351,5 +351,5 @@ int main(){
#ifdef ENABLE_DESTRUCTOR_TEST
classTests.TestForConsrtDesrt();
test_passed(TestForConsrtDesrt);
#endif
#endif
}
projects/hip/tests/src/deviceLib/hipBfloat16.cpp
+6 -6
مشاهده پرونده
@@ -55,17 +55,17 @@ __host__ __device__ void testOperations(float &fa, float &fb) {
hip_bfloat16 bf_a(fa);
hip_bfloat16 bf_b(fb);
float fc = float(bf_a);
float fd = float(bf_b);
float fd = float(bf_b);
assert(testRelativeAccuracy(fa, bf_a));
assert(testRelativeAccuracy(fb, bf_b));
assert(testRelativeAccuracy(fc + fd, bf_a + bf_b));
//when checked as above for add, operation sub fails on GPU
//when checked as above for add, operation sub fails on GPU
assert(hip_bfloat16(fc - fd) == (bf_a - bf_b));
assert(testRelativeAccuracy(fc * fd, bf_a * bf_b));
assert(testRelativeAccuracy(fc / fd, bf_a / bf_b));
hip_bfloat16 bf_opNegate = -bf_a;
assert(bf_opNegate == -bf_a);
@@ -75,7 +75,7 @@ __host__ __device__ void testOperations(float &fa, float &fb) {
bf_x--;
++bf_x;
--bf_x;
//hip_bfloat16 is converted to float and then inc/decremented, hence check with reduced precision
//hip_bfloat16 is converted to float and then inc/decremented, hence check with reduced precision
assert(testRelativeAccuracy(bf_x,bf_a));
bf_x = bf_a;
@@ -95,7 +95,7 @@ __host__ __device__ void testOperations(float &fa, float &fb) {
if (isnan(bf_rounded)) {
assert(isnan(bf_rounded) || isinf(bf_rounded));
}
}
}
__global__ void testOperationsGPU(float* d_a, float* d_b)
{
@@ -126,7 +126,7 @@ int main(){
hipLaunchKernelGGL(testOperationsGPU, 1, SIZE, 0, 0, d_fa, d_fb);
hipDeviceSynchronize();
cout<<"Device bfloat16 Operations Successful!!"<<endl;
cout<<"Device bfloat16 Operations Successful!!"<<endl;
delete[] h_fa;
delete[] h_fb;
projects/hip/tests/src/deviceLib/hipMathFunctions.cpp
+14 -14
مشاهده پرونده
@@ -56,21 +56,21 @@ __global__ void kernel_lgamma_double(double *input, double *output) {
void check_lgamma_double() {
using datatype_t = double;
const int NUM_INPUTS = 8;
auto memsize = NUM_INPUTS * sizeof(datatype_t);
// allocate memories
datatype_t *inputCPU = (datatype_t *) malloc(memsize);
datatype_t *outputCPU = (datatype_t *) malloc(memsize);
datatype_t *inputGPU = nullptr; hipMalloc((void**)&inputGPU, memsize);
datatype_t *outputGPU = nullptr; hipMalloc((void**)&outputGPU, memsize);
// populate input
for (int i=0; i<NUM_INPUTS; i++) {
inputCPU[i] = -3.5 + i;
}
// copy inputs to device
hipMemcpy(inputGPU, inputCPU, memsize, hipMemcpyHostToDevice);
@@ -84,13 +84,13 @@ void check_lgamma_double() {
for (int i=0; i<NUM_INPUTS; i++) {
CHECK_LGAMMA_DOUBLE(inputCPU[i], outputCPU[i], lgamma(inputCPU[i]));
}
// free memories
hipFree(inputGPU);
hipFree(outputGPU);
free(inputCPU);
free(outputCPU);
// done
return;
}
@@ -102,15 +102,15 @@ void check_abs_int64() {
const int NUM_INPUTS = 8;
auto memsize = NUM_INPUTS * sizeof(datatype_t);
// allocate memories
datatype_t *inputCPU = (datatype_t *) malloc(memsize);
datatype_t *outputCPU = (datatype_t *) malloc(memsize);
datatype_t *inputGPU = nullptr; hipMalloc((void**)&inputGPU, memsize);
datatype_t *outputGPU = nullptr; hipMalloc((void**)&outputGPU, memsize);
// populate input
inputCPU[0] = -81985529216486895ll;
inputCPU[0] = -81985529216486895ll;
inputCPU[1] = 81985529216486895ll;
inputCPU[2] = -1250999896491ll;
inputCPU[3] = 1250999896491ll;
@@ -118,7 +118,7 @@ void check_abs_int64() {
inputCPU[5] = 19088743ll;
inputCPU[6] = -291ll;
inputCPU[7] = 291ll;
// copy inputs to device
hipMemcpy(inputGPU, inputCPU, memsize, hipMemcpyHostToDevice);
@@ -137,17 +137,17 @@ void check_abs_int64() {
CHECK_ABS_INT64(inputCPU[5], outputCPU[5], outputCPU[5]);
CHECK_ABS_INT64(inputCPU[6], outputCPU[6], outputCPU[7]);
CHECK_ABS_INT64(inputCPU[7], outputCPU[7], outputCPU[7]);
// free memories
hipFree(inputGPU);
hipFree(outputGPU);
free(inputCPU);
free(outputCPU);
// done
return;
}
template<class T, class F>
__global__ void kernel_simple(F f, T *out) {
@@ -191,7 +191,7 @@ int main(int argc, char* argv[]) {
check_abs_int64();
// check_lgamma_double();
test_fp16();
test_pown();
projects/hip/tests/src/deviceLib/hipStdComplex.cpp
+4 -4
مشاهده پرونده
@@ -82,7 +82,7 @@ __device__ __host__ complex<FloatT> calc(complex<FloatT> A,
return A * B;
case CK_div:
return A / B;
ONE_ARG(abs)
ONE_ARG(arg)
ONE_ARG(sin)
@@ -111,7 +111,7 @@ void test() {
hipMalloc((void**)&Ad, sizeof(ComplexT)*LEN);
hipMalloc((void**)&Bd, sizeof(ComplexT)*LEN);
hipMalloc((void**)&Cd, sizeof(ComplexT)*LEN);
for (uint32_t i = 0; i < LEN; i++) {
A[i] = ComplexT((i + 1) * 1.0f, (i + 2) * 1.0f);
B[i] = A[i];
@@ -119,7 +119,7 @@ void test() {
}
hipMemcpy(Ad, A, sizeof(ComplexT)*LEN, hipMemcpyHostToDevice);
hipMemcpy(Bd, B, sizeof(ComplexT)*LEN, hipMemcpyHostToDevice);
// Run kernel for a calculation kind and verify by comparing with host
// calculation result. Returns false if fails.
auto test_fun = [&](enum CalcKind CK) {
@@ -145,7 +145,7 @@ void test() {
}
return true;
};
#define OP(x) assert(test_fun(CK_##x));
ALL_FUN
#undef OP
projects/hip/tests/src/deviceLib/hipTestHalf.cpp
+1 -1
مشاهده پرونده
@@ -84,7 +84,7 @@ void kernel_hisnan(__half* input, int* output) {
}
__global__
void kernel_hisinf(__half* input, int* output) {
void kernel_hisinf(__half* input, int* output) {
int tx = threadIdx.x;
output[tx] = __hisinf(input[tx]);
}
projects/hip/tests/src/deviceLib/hipTestNew.cpp
+1 -1
مشاهده پرونده
@@ -41,7 +41,7 @@ THE SOFTWARE.
private:
int a;
};
static __global__ void kernel(int* Ad) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
new(Ad+tid) A();
projects/hip/tests/src/hipEnvVarDriver.cpp
+2 -2
مشاهده پرونده
@@ -41,7 +41,7 @@ int readHipEnvVar(string flags, char* buff){
std::cout << "\nFinding hipEnvVar in " << directed_dir << "...\n";
FILE* directed_in = popen((directed_dir + flags).c_str(), "r");
if(fgets(buff, 512, directed_in) == NULL){
std::cout << "Finding hipEnvVar in " << dir << "...\n";
FILE* in = popen((dir + flags).c_str(), "r");
@@ -74,7 +74,7 @@ int getDeviceNumber(bool print_err=true) {
}
// Query the current device ID remotely to hipEnvVar
void getDevicePCIBusNumRemote(int deviceID, char* pciBusID) {
void getDevicePCIBusNumRemote(int deviceID, char* pciBusID) {
std::this_thread::sleep_for(std::chrono::milliseconds(10));
if (readHipEnvVar((" -d " + std::to_string(deviceID)), pciBusID)){
std::cerr << "The system cannot find hipEnvVar\n";
projects/hip/tests/src/ipc/hipMultiProcIpcEvent.cpp
Executable → Regular
مشاهده پرونده
projects/hip/tests/src/ipc/hipMultiProcIpcMem.cpp
Executable → Regular
مشاهده پرونده
projects/hip/tests/src/kernel/hipLaunchParmFunctor.cpp
Executable → Regular
مشاهده پرونده
projects/hip/tests/src/printf/hipPrintfManyDevices.cpp
Executable → Regular
مشاهده پرونده
projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaBaseType.cpp
Executable → Regular
مشاهده پرونده
projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaPublicApi.cpp
Executable → Regular
مشاهده پرونده
projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupType.cpp
Executable → Regular
مشاهده پرونده
projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockType.cpp
Executable → Regular
مشاهده پرونده
projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockTypeViaBaseType.cpp
Executable → Regular
مشاهده پرونده
projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockTypeViaPublicApi.cpp
Executable → Regular
مشاهده پرونده
projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoD.cpp
Executable → Regular
مشاهده پرونده
projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp
Executable → Regular
مشاهده پرونده
projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeer.cpp
Executable → Regular
مشاهده پرونده
projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeerAsync.cpp
Executable → Regular
مشاهده پرونده
projects/hip/tests/src/runtimeApi/memory/hipMemcpyWithStream.cpp
Executable → Regular
مشاهده پرونده
projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp
Executable → Regular
مشاهده پرونده
projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp
Executable → Regular
مشاهده پرونده