cory-master/BabelStream
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[HIP] Fixes to work with latest HIP

Browse Source 
- Remove hipLaunchParm lp from __global__ kernels
- Replace hipLaunchKernel with hipLaunchKernelGGL
- Pass on template parameters to kernels
This commit is contained in:
Maneesh Gupta 2018-03-19 15:18:57 +05:30
parent 6958f070b1
commit d664544afd
1 changed files with 12 additions and 12 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

24
HIPStream.cpp
View File

@ -86,7 +86,7 @@ HIPStream<T>::~HIPStream()
template <typename T>
__global__ void init_kernel(hipLaunchParm lp, T * a, T * b, T * c, T initA, T initB, T initC)
__global__ void init_kernel(T * a, T * b, T * c, T initA, T initB, T initC)
{
const int i = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
a[i] = initA;
@ -97,7 +97,7 @@ __global__ void init_kernel(hipLaunchParm lp, T * a, T * b, T * c, T initA, T in
template <class T>
void HIPStream<T>::init_arrays(T initA, T initB, T initC)
{
hipLaunchKernel(HIP_KERNEL_NAME(init_kernel), dim3(array_size/TBSIZE), dim3(TBSIZE), 0, 0, d_a, d_b, d_c, initA, initB, initC);
hipLaunchKernelGGL(HIP_KERNEL_NAME(init_kernel<T>), dim3(array_size/TBSIZE), dim3(TBSIZE), 0, 0, d_a, d_b, d_c, initA, initB, initC);
check_error();
hipDeviceSynchronize();
check_error();
@ -117,7 +117,7 @@ void HIPStream<T>::read_arrays(std::vector<T>& a, std::vector<T>& b, std::vector
template <typename T>
__global__ void copy_kernel(hipLaunchParm lp, const T * a, T * c)
__global__ void copy_kernel(const T * a, T * c)
{
const int i = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
c[i] = a[i];
@ -126,14 +126,14 @@ __global__ void copy_kernel(hipLaunchParm lp, const T * a, T * c)
template <class T>
void HIPStream<T>::copy()
{
hipLaunchKernel(HIP_KERNEL_NAME(copy_kernel), dim3(array_size/TBSIZE), dim3(TBSIZE), 0, 0, d_a, d_c);
hipLaunchKernelGGL(HIP_KERNEL_NAME(copy_kernel<T>), dim3(array_size/TBSIZE), dim3(TBSIZE), 0, 0, d_a, d_c);
check_error();
hipDeviceSynchronize();
check_error();
}
template <typename T>
__global__ void mul_kernel(hipLaunchParm lp, T * b, const T * c)
__global__ void mul_kernel(T * b, const T * c)
{
const T scalar = startScalar;
const int i = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
@ -143,14 +143,14 @@ __global__ void mul_kernel(hipLaunchParm lp, T * b, const T * c)
template <class T>
void HIPStream<T>::mul()
{
hipLaunchKernel(HIP_KERNEL_NAME(mul_kernel), dim3(array_size/TBSIZE), dim3(TBSIZE), 0, 0, d_b, d_c);
hipLaunchKernelGGL(HIP_KERNEL_NAME(mul_kernel<T>), dim3(array_size/TBSIZE), dim3(TBSIZE), 0, 0, d_b, d_c);
check_error();
hipDeviceSynchronize();
check_error();
}
template <typename T>
__global__ void add_kernel(hipLaunchParm lp, const T * a, const T * b, T * c)
__global__ void add_kernel(const T * a, const T * b, T * c)
{
const int i = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
c[i] = a[i] + b[i];
@ -159,14 +159,14 @@ __global__ void add_kernel(hipLaunchParm lp, const T * a, const T * b, T * c)
template <class T>
void HIPStream<T>::add()
{
hipLaunchKernel(HIP_KERNEL_NAME(add_kernel), dim3(array_size/TBSIZE), dim3(TBSIZE), 0, 0, d_a, d_b, d_c);
hipLaunchKernelGGL(HIP_KERNEL_NAME(add_kernel<T>), dim3(array_size/TBSIZE), dim3(TBSIZE), 0, 0, d_a, d_b, d_c);
check_error();
hipDeviceSynchronize();
check_error();
}
template <typename T>
__global__ void triad_kernel(hipLaunchParm lp, T * a, const T * b, const T * c)
__global__ void triad_kernel(T * a, const T * b, const T * c)
{
const T scalar = startScalar;
const int i = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
@ -176,14 +176,14 @@ __global__ void triad_kernel(hipLaunchParm lp, T * a, const T * b, const T * c)
template <class T>
void HIPStream<T>::triad()
{
hipLaunchKernel(HIP_KERNEL_NAME(triad_kernel), dim3(array_size/TBSIZE), dim3(TBSIZE), 0, 0, d_a, d_b, d_c);
hipLaunchKernelGGL(HIP_KERNEL_NAME(triad_kernel<T>), dim3(array_size/TBSIZE), dim3(TBSIZE), 0, 0, d_a, d_b, d_c);
check_error();
hipDeviceSynchronize();
check_error();
}
template <class T>
__global__ void dot_kernel(hipLaunchParm lp, const T * a, const T * b, T * sum, unsigned int array_size)
__global__ void dot_kernel(const T * a, const T * b, T * sum, unsigned int array_size)
{
__shared__ T tb_sum[TBSIZE];
@ -210,7 +210,7 @@ __global__ void dot_kernel(hipLaunchParm lp, const T * a, const T * b, T * sum,
template <class T>
T HIPStream<T>::dot()
{
hipLaunchKernel(HIP_KERNEL_NAME(dot_kernel), dim3(DOT_NUM_BLOCKS), dim3(TBSIZE), 0, 0, d_a, d_b, d_sum, array_size);
hipLaunchKernelGGL(HIP_KERNEL_NAME(dot_kernel<T>), dim3(DOT_NUM_BLOCKS), dim3(TBSIZE), 0, 0, d_a, d_b, d_sum, array_size);
check_error();
hipMemcpy(sums, d_sum, DOT_NUM_BLOCKS*sizeof(T), hipMemcpyDeviceToHost);