cms::rocmtest Namespace Reference

Functions
__device__ void	add_vectors_d (const double *__restrict__ in1, const double *__restrict__ in2, double *__restrict__ out, size_t size)
__device__ void	add_vectors_f (const float *__restrict__ in1, const float *__restrict__ in2, float *__restrict__ out, size_t size)
__global__ void	kernel_add_vectors_d (const double *__restrict__ in1, const double *__restrict__ in2, double *__restrict__ out, size_t size)
__global__ void	kernel_add_vectors_f (const float *__restrict__ in1, const float *__restrict__ in2, float *__restrict__ out, size_t size)
void	opqaue_add_vectors_d (const double *in1, const double *in2, double *out, size_t size)
void	opqaue_add_vectors_f (const float *in1, const float *in2, float *out, size_t size)
void	requireDevices ()
	Print message and exit if there are no ROCm devices. More...
bool	testDevices ()
	In presence of ROCm devices, return true; otherwise print message and return false. More...
void	wrapper_add_vectors_d (const double *__restrict__ in1, const double *__restrict__ in2, double *__restrict__ out, size_t size)
void	wrapper_add_vectors_f (const float *__restrict__ in1, const float *__restrict__ in2, float *__restrict__ out, size_t size)

Function Documentation

add_vectors_d()

__device__ void cms::rocmtest::add_vectors_d	(	const double *__restrict__	in1,
		const double *__restrict__	in2,
		double *__restrict__	out,
		size_t	size
	)

Definition at line 22 of file DeviceAddition.hip.cc.

References cms::cudacompat::blockDim, cms::cudacompat::blockIdx, cms::cudacompat::gridDim, mps_fire::i, MillePedeFileConverter_cfg::out, gpuPixelDoublets::stride, and cms::cudacompat::threadIdx.

Referenced by kernel_add_vectors_d().

                                             {
    uint32_t thread = threadIdx.x + blockIdx.x * blockDim.x;
    uint32_t stride = blockDim.x * gridDim.x;

    for (size_t i = thread; i < size; i += stride) {
      out[i] = in1[i] + in2[i];
    }
  }

add_vectors_f()

__device__ void cms::rocmtest::add_vectors_f	(	const float *__restrict__	in1,
		const float *__restrict__	in2,
		float *__restrict__	out,
		size_t	size
	)

Definition at line 10 of file DeviceAddition.hip.cc.

References cms::cudacompat::blockDim, cms::cudacompat::blockIdx, cms::cudacompat::gridDim, mps_fire::i, MillePedeFileConverter_cfg::out, gpuPixelDoublets::stride, and cms::cudacompat::threadIdx.

Referenced by kernel_add_vectors_f(), and HeterogeneousCoreROCmTestDevicePlugins::kernel_add_vectors_f().

                                             {
    uint32_t thread = threadIdx.x + blockIdx.x * blockDim.x;
    uint32_t stride = blockDim.x * gridDim.x;

    for (size_t i = thread; i < size; i += stride) {
      out[i] = in1[i] + in2[i];
    }
  }

kernel_add_vectors_d()

__global__ void cms::rocmtest::kernel_add_vectors_d	(	const double *__restrict__	in1,
		const double *__restrict__	in2,
		double *__restrict__	out,
		size_t	size
	)

Definition at line 17 of file DeviceAdditionKernel.hip.cc.

References add_vectors_d(), and MillePedeFileConverter_cfg::out.

                                                    {
    add_vectors_d(in1, in2, out, size);
  }

kernel_add_vectors_f()

__global__ void cms::rocmtest::kernel_add_vectors_f	(	const float *__restrict__	in1,
		const float *__restrict__	in2,
		float *__restrict__	out,
		size_t	size
	)

Definition at line 10 of file DeviceAdditionKernel.hip.cc.

References add_vectors_f(), and MillePedeFileConverter_cfg::out.

                                                    {
    add_vectors_f(in1, in2, out, size);
  }

opqaue_add_vectors_d()

void cms::rocmtest::opqaue_add_vectors_d	(	const double *	in1,
		const double *	in2,
		double *	out,
		size_t	size
	)

Definition at line 42 of file DeviceAdditionOpaque.cc.

References hipCheck, and wrapper_add_vectors_d().

                                                                                                  {
    // allocate input and output buffers on the device
    double* in1_d;
    double* in2_d;
    double* out_d;
    hipCheck(hipMalloc(&in1_d, size * sizeof(double)));
    hipCheck(hipMalloc(&in2_d, size * sizeof(double)));
    hipCheck(hipMalloc(&out_d, size * sizeof(double)));

    // copy the input data to the device
    hipCheck(hipMemcpy(in1_d, in1_h, size * sizeof(double), hipMemcpyHostToDevice));
    hipCheck(hipMemcpy(in2_d, in2_h, size * sizeof(double), hipMemcpyHostToDevice));

    // fill the output buffer with zeros
    hipCheck(hipMemset(out_d, 0, size * sizeof(double)));

    // launch the 1-dimensional kernel for vector addition
    wrapper_add_vectors_d(in1_d, in2_d, out_d, size);

    // copy the results from the device to the host
    hipCheck(hipMemcpy(out_h, out_d, size * sizeof(double), hipMemcpyDeviceToHost));

    // wait for all the operations to complete
    hipCheck(hipDeviceSynchronize());

    // free the input and output buffers on the device
    hipCheck(hipFree(in1_d));
    hipCheck(hipFree(in2_d));
    hipCheck(hipFree(out_d));
  }

opqaue_add_vectors_f()

void cms::rocmtest::opqaue_add_vectors_f	(	const float *	in1,
		const float *	in2,
		float *	out,
		size_t	size
	)

Definition at line 11 of file DeviceAdditionOpaque.cc.

References hipCheck, and wrapper_add_vectors_f().

Referenced by ROCmTestOpaqueAdditionModule::analyze().

                                                                                               {
    // allocate input and output buffers on the device
    float* in1_d;
    float* in2_d;
    float* out_d;
    hipCheck(hipMalloc(&in1_d, size * sizeof(float)));
    hipCheck(hipMalloc(&in2_d, size * sizeof(float)));
    hipCheck(hipMalloc(&out_d, size * sizeof(float)));

    // copy the input data to the device
    hipCheck(hipMemcpy(in1_d, in1_h, size * sizeof(float), hipMemcpyHostToDevice));
    hipCheck(hipMemcpy(in2_d, in2_h, size * sizeof(float), hipMemcpyHostToDevice));

    // fill the output buffer with zeros
    hipCheck(hipMemset(out_d, 0, size * sizeof(float)));

    // launch the 1-dimensional kernel for vector addition
    wrapper_add_vectors_f(in1_d, in2_d, out_d, size);

    // copy the results from the device to the host
    hipCheck(hipMemcpy(out_h, out_d, size * sizeof(float), hipMemcpyDeviceToHost));

    // wait for all the operations to complete
    hipCheck(hipDeviceSynchronize());

    // free the input and output buffers on the device
    hipCheck(hipFree(in1_d));
    hipCheck(hipFree(in2_d));
    hipCheck(hipFree(out_d));
  }

requireDevices()

void cms::rocmtest::requireDevices

(

)

Print message and exit if there are no ROCm devices.

Definition at line 24 of file requireDevices.cc.

References beamvalidation::exit(), and testDevices().

                        {
    if (not testDevices()) {
      exit(EXIT_SUCCESS);
    }
  }

testDevices()

bool cms::rocmtest::testDevices

(

)

In presence of ROCm devices, return true; otherwise print message and return false.

Definition at line 10 of file requireDevices.cc.

References DMR_cfg::cerr, cms::alpakatools::devices(), and mps_update::status.

Referenced by requireDevices().

                     {
    int devices = 0;
    auto status = hipGetDeviceCount(&devices);
    if (status != hipSuccess) {
      std::cerr << "Failed to initialise the ROCm runtime, the test will be skipped.\n";
      return false;
    }
    if (devices == 0) {
      std::cerr << "No ROCm devices available, the test will be skipped.\n";
      return false;
    }
    return true;
  }

wrapper_add_vectors_d()

void cms::rocmtest::wrapper_add_vectors_d	(	const double *__restrict__	in1,
		const double *__restrict__	in2,
		double *__restrict__	out,
		size_t	size
	)

Definition at line 20 of file DeviceAdditionWrapper.hip.cc.

References hipCheck, and MillePedeFileConverter_cfg::out.

Referenced by cms::cudatest::opaque_add_vectors_d(), and opqaue_add_vectors_d().

                                          {
    // launch the 1-dimensional kernel for vector addition
    kernel_add_vectors_d<<<32, 32>>>(in1, in2, out, size);
    hipCheck(hipGetLastError());
  }

wrapper_add_vectors_f()

void cms::rocmtest::wrapper_add_vectors_f	(	const float *__restrict__	in1,
		const float *__restrict__	in2,
		float *__restrict__	out,
		size_t	size
	)

Definition at line 11 of file DeviceAdditionWrapper.hip.cc.

References hipCheck, and MillePedeFileConverter_cfg::out.

Referenced by CUDATestWrapperAdditionModule::analyze(), ROCmTestWrapperAdditionModule::analyze(), CUDATestKernelAdditionModule::analyze(), CUDATestDeviceAdditionModule::analyze(), cms::cudatest::opaque_add_vectors_f(), and opqaue_add_vectors_f().

                                          {
    // launch the 1-dimensional kernel for vector addition
    kernel_add_vectors_f<<<32, 32>>>(in1, in2, out, size);
    hipCheck(hipGetLastError());
  }