1

我正在尝试使用 JCuda 中的共享内存来实现简单的矩阵乘法程序。

以下是我的 JCudaSharedMatrixMul.java 代码：

import static jcuda.driver.JCudaDriver.cuCtxCreate;
import static jcuda.driver.JCudaDriver.cuCtxSynchronize;
import static jcuda.driver.JCudaDriver.cuDeviceGet;
import static jcuda.driver.JCudaDriver.cuInit;
import static jcuda.driver.JCudaDriver.cuLaunchKernel;
import static jcuda.driver.JCudaDriver.cuMemAlloc;
import static jcuda.driver.JCudaDriver.cuMemFree;
import static jcuda.driver.JCudaDriver.cuMemcpyDtoH;
import static jcuda.driver.JCudaDriver.cuMemcpyHtoD;
import static jcuda.driver.JCudaDriver.cuModuleGetFunction;
import static jcuda.driver.JCudaDriver.cuModuleLoad;
import static jcuda.runtime.JCuda.cudaEventCreate;
import static jcuda.runtime.JCuda.cudaEventRecord;
import static jcuda.runtime.JCuda.*;

import java.io.ByteArrayOutputStream;
import java.io.File;
import java.io.IOException;
import java.io.InputStream;
import java.util.Scanner;

import jcuda.Pointer;
import jcuda.Sizeof;
import jcuda.driver.CUcontext;
import jcuda.driver.CUdevice;
import jcuda.driver.CUdeviceptr;
import jcuda.driver.CUfunction;
import jcuda.driver.CUmodule;
import jcuda.driver.JCudaDriver;
import jcuda.runtime.cudaEvent_t;



public class JCudaSharedMatrixMul
{

    public static void main(String[] args) throws IOException 
    {
        // Enable exceptions and omit all subsequent error checks
        JCudaDriver.setExceptionsEnabled(true);

        // Create the PTX file by calling the NVCC
        String ptxFilename = preparePtxFile("JCudaSharedMatrixMulKernel.cu");

        //Initialize the driver and create a context for the first device.
        cuInit(0);
        CUdevice device = new CUdevice();
        cuDeviceGet (device, 0);
        CUcontext context = new CUcontext();
        cuCtxCreate(context, 0, device);

        //Load PTX file
        CUmodule module = new CUmodule();
        cuModuleLoad(module,ptxFilename);

        //Obtain a function pointer to the Add function
        CUfunction function = new CUfunction();
        cuModuleGetFunction(function, module, "jCudaSharedMatrixMulKernel");

        int numRows = 16;
        int numCols = 16;

        //Allocate and fill Host input Matrices:
        float hostMatrixA[] = new float[numRows*numCols];
        float hostMatrixB[] = new float[numRows*numCols];
        float hostMatrixC[] = new float[numRows*numCols];


        for(int i = 0; i<numRows; i++)

        {
            for(int j = 0; j<numCols; j++)
            {
                hostMatrixA[i*numCols+j] = (float) 1;
                hostMatrixB[i*numCols+j] = (float) 1;
            }
        }
        // Allocate the device input data, and copy the
        // host input data to the device
        CUdeviceptr devMatrixA = new CUdeviceptr();
        cuMemAlloc(devMatrixA, numRows * numCols * Sizeof.FLOAT);

        //This is the part where it gives me the error
        cuMemcpyHtoD(devMatrixA, Pointer.to(hostMatrixA), numRows * numCols * Sizeof.FLOAT);

        CUdeviceptr devMatrixB = new CUdeviceptr();
        cuMemAlloc(devMatrixB, numRows * numCols * Sizeof.FLOAT);

        //This is the part where it gives me the error
        cuMemcpyHtoD(devMatrixB, Pointer.to(hostMatrixB ), numRows * numCols * Sizeof.FLOAT);

        //Allocate device matrix C to store output
        CUdeviceptr devMatrixC = new CUdeviceptr();
        cuMemAlloc(devMatrixC, numRows * numCols * Sizeof.FLOAT);

        // Set up the kernel parameters: A pointer to an array
        // of pointers which point to the actual values.

        Pointer kernelParameters = Pointer.to(
                                Pointer.to(new int[]{numCols}), 
                                Pointer.to(devMatrixA),
                                Pointer.to(devMatrixB),
                                Pointer.to(devMatrixC));

        //Kernel thread configuration
        int blockSize = 16;
        int gridSize = 1;

        cudaEvent_t start = new cudaEvent_t();
        cudaEvent_t stop = new cudaEvent_t();
        cudaEventCreate(start);
        cudaEventCreate(stop);
        long start_nano=System.nanoTime();
        cudaEventRecord(start, null);

        cuLaunchKernel(function, 
                       gridSize, 1, 1,
                       blockSize, 16, 1,
                       250, null, kernelParameters, null);

        cuCtxSynchronize();
        cudaEventRecord(stop, null);
        long end_nano=System.nanoTime();
        float elapsedTimeMsArray[] = { Float.NaN };
        cudaEventElapsedTime(elapsedTimeMsArray, start, stop);
        float elapsedTimeMs = elapsedTimeMsArray[0];
        System.out.println("Time Required (Using cudaevent elapsed time) = " + " " +elapsedTimeMs+
                "Time Required (Using nanotime)= "+(end_nano-start_nano)/1000000);
        // Allocate host output memory and copy the device output
        // to the host.

        //This is the part where it gives me the error
        cuMemcpyDtoH(Pointer.to(hostMatrixC), devMatrixC, numRows * numCols * Sizeof.FLOAT);



        //verify the result
        for (int i =0; i<numRows; i++)
        {
            for (int j =0; j<numRows; j++)
            {
                System.out.print("   "+ hostMatrixC[i*numCols+j]);
            }
            System.out.println("");
        }

        cuMemFree(devMatrixA);
        cuMemFree(devMatrixB);
        cuMemFree(devMatrixC);
    }

    private static String preparePtxFile(String cuFileName) throws IOException
    {
        int endIndex = cuFileName.lastIndexOf('.');
        if (endIndex == -1)
            endIndex = cuFileName.length()-1;
        {
        }
        String ptxFileName = cuFileName.substring(0, endIndex+1)+"ptx";
        File ptxFile = new File(ptxFileName);
        if (ptxFile.exists())
        {
            return ptxFileName;
        }

        File cuFile = new File(cuFileName);
        if (!cuFile.exists())
        {
            throw new IOException("Input file not found: "+cuFileName);
        }
        String modelString = "-m"+System.getProperty("sun.arch.data.model");
        String command = "nvcc " + modelString + " -ptx "+ cuFile.getPath()+" -o "+ptxFileName;

        System.out.println("Executing\n"+command);
        Process process = Runtime.getRuntime().exec(command);

        String errorMessage = new String(toByteArray(process.getErrorStream()));
        String outputMessage = new String(toByteArray(process.getInputStream()));
        int exitValue = 0;
        try
        {
            exitValue = process.waitFor();
        }
        catch (InterruptedException e)
        {
            Thread.currentThread().interrupt();
            throw new IOException(
                "Interrupted while waiting for nvcc output", e);
        }

        if (exitValue != 0)
        {
            System.out.println("nvcc process exitValue "+exitValue);
            System.out.println("errorMessage:\n"+errorMessage);
            System.out.println("outputMessage:\n"+outputMessage);
            throw new IOException(
                "Could not create .ptx file: "+errorMessage);
        }

        System.out.println("Finished creating PTX file");
        return ptxFileName;

    }

    private static byte[] toByteArray(InputStream inputStream) throws IOException
    {
        ByteArrayOutputStream baos = new ByteArrayOutputStream();
        byte buffer[] = new byte[8192];
        while (true)
        {
            int read = inputStream.read(buffer);
            if (read == -1)
            {
                break;
            }
            baos.write(buffer, 0, read);
        }
        return baos.toByteArray();
    }


}

以下是我的 JCudaSharedMatrixMulKernel.cu 代码：

extern "C"
__global__ void jCudaSharedMatrixMulKernel(int N,float *ad,float *bd,float *cd)
{
    float pvalue=0;
    int TILE=blockDim.x;
    int ty=threadIdx.y;
    int tx=threadIdx.x;

    __shared__ float ads[4][4];
    __shared__ float bds[4][4];

    int Row = blockIdx.y * blockDim.y + threadIdx.y;
    int Col = blockIdx.x * blockDim.x + threadIdx.x;

    for(int i=0;i< N/TILE;++i)
    {
        ads[ty][tx] = ad[Row * N + (i * TILE) + tx];
        bds[ty][tx] = bd[(i * TILE + ty) * N + Col];

        __syncthreads();

        for(int k=0;k<TILE;k++)
                pvalue += ads[ty][k] * bds[k][tx];

        __syncthreads();  
    }

    cd[Row * N + Col] = pvalue;
}

在我上面的示例中，每个块使用的总共享内存为 2*4*4*4 = 128 字节。在 cuLaunchKernel 中，当我将 sharedMemBytes参数定义为 0（零）时，它会给我以下错误：

**Exception in thread "main" jcuda.CudaException: CUDA_ERROR_LAUNCH_FAILED
    at jcuda.driver.JCudaDriver.checkResult(JCudaDriver.java:282)
    at jcuda.driver.JCudaDriver.cuCtxSynchronize(JCudaDriver.java:1795)
    at JCudaSharedMatrixMul.main(JCudaSharedMatrixMul.java:121)**

当我将其定义为 128 时，它会给出与上述相同的错误。但是当我把它设为 129 时，它会给我正确的输出！当我给出 129 到 49024 之间的任何值时，它会给我正确的结果。我的问题是为什么我将其定义为 128 时无法获得正确的输出？另外可以定义的最大共享内存是多少？为什么这个 129-49024 范围在这里工作？

4

1 回答 1

3

您正在启动 16x16 线程块：

    cuLaunchKernel(function, 
                   gridSize, 1, 1,
                   blockSize, 16, 1,  <-- the first two params are block.x and block.y
                   250, null, kernelParameters, null);

所以__shared__ float ads[4][4];根本不应该工作。例如，这些内核代码行对于某些线程来说会越界访问这些共享数组：

    ads[ty][tx] = ad[Row * N + (i * TILE) + tx];
    bds[ty][tx] = bd[(i * TILE + ty) * N + Col];
         ^   ^
         |   tx goes from 0..15 for a 16x16 threadblock
         ty goes from 0..15 for a 16x16 threadblock

您的代码在这方面被破坏了。如果您使用cuda-memcheck它运行代码，即使在您“通过”的情况下，也可能会捕获这些越界访问。查看matrixMulDrv cuda 示例代码，将具有指导意义，您会看到共享内存分配是 ，2*block_size*block_size因为它也应该适用于您的情况，但是您的共享内存定义不[16][16]应该[4][4]当您超过 128 个字节时恰好可以工作，但是您的代码存在缺陷。

您的共享定义应该是：

__shared__ float ads[16][16];
__shared__ float bds[16][16];

由于上述分配是静态分配，并且sharedMemBytes参数定义为动态共享内存分配，因此对于本示例您不需要分配任何（0 为 OK）动态共享内存，它仍然有效。此处介绍了静态和动态之间的区别。

每个块的最大共享内存在文档中可用，或者如果您运行 cudadeviceQuery示例代码。对于 cc2.0 和更新的设备，它是 48K 字节。

于 2014-08-18T11:15:14.167 回答