t

GnibChen · Jan 30, 2022 · df7e227 · df7e227
1 parent 7e328ce
commit df7e227
Show file tree

Hide file tree

Showing 2 changed files with 15 additions and 14 deletions.
diff --git a/09/01_texture/01/main.cu b/09/01_texture/01/main.cu
@@ -14,22 +14,23 @@ __global__ void kernel(cudaTextureObject_t texVel, cudaSurfaceObject_t sufLoc, u
     unsigned int z = threadIdx.z + blockDim.z * blockIdx.z;
     if (x >= n || y >= n || z >= n) return;
     float4 vel = tex3D<float4>(texVel, x, y, z);
-    float4 loc = make_float4(x + 0.5f, y + 0.5f, z + 0.5f, 1.f) - vel;
-    surf3Dwrite<float4>(loc, sufLoc, x, y, z, cudaBoundaryModeTrap);
+    float4 loc = make_float4(x + 0.5f, y + 0.5f, z + 0.5f, 42.f) - vel;
+    surf3Dwrite<float4>(loc, sufLoc, 1, 1, 1, cudaBoundaryModeTrap);
 }
 
 int main() {
     unsigned int n = 2;
 
-    auto arrLoc = CudaArray<float4>::make({{n, n, n}});
+    auto arrLoc = CudaArray<float4>::make({{n, n, n}});//, cudaCreateChannelDesc(16, 16, 16, 16, cudaChannelFormatKindFloat)});
     auto sufLoc = CudaSurface<float4>::make(arrLoc);
-    auto arrVel = CudaArray<float4>::make({{n, n, n}});
+    auto arrVel = CudaArray<float4>::make({{n, n, n}});//, cudaCreateChannelDesc(16, 16, 16, 16, cudaChannelFormatKindFloat)});
+    auto sufVel = CudaSurface<float4>::make(arrVel);
     auto texVel = CudaTexture<float4>::make(arrVel);
 
     std::vector<float4> cpuVel(n * n * n);
-    for (int z = 0; z < n; z++) {
-        for (int y = 0; y < n; y++) {
-            for (int x = 0; x < n; x++) {
+    for (unsigned int z = 0; z < n; z++) {
+        for (unsigned int y = 0; y < n; y++) {
+            for (unsigned int x = 0; x < n; x++) {
                 cpuVel[x + n * (y + n * z)] = make_float4(1.f, 0.f, 0.f, 0.f);
             }
         }
@@ -40,11 +41,11 @@ int main() {
 
     std::vector<float4> cpuLoc(n * n * n);
     arrLoc.copyOut(cpuLoc.data());
-    for (int z = 0; z < n; z++) {
-        for (int y = 0; y < n; y++) {
-            for (int x = 0; x < n; x++) {
+    for (unsigned int z = 0; z < n; z++) {
+        for (unsigned int y = 0; y < n; y++) {
+            for (unsigned int x = 0; x < n; x++) {
                 float4 val = cpuLoc[x + n * (y + n * z)];
-                printf("%d,%d,%d: %f,%f,%f,%f\n", x, y, z, val.x, val.y, val.z, val.w);
+                printf("%u,%u,%u: %f,%f,%f,%f\n", x, y, z, val.x, val.y, val.z, val.w);
             }
         }
     }

diff --git a/09/include/CudaArray.cuh b/09/include/CudaArray.cuh
@@ -11,7 +11,7 @@ class CudaArray {
     struct BuildArgs {
         std::array<unsigned int, 3> const dim{};
         cudaChannelFormatDesc desc{cudaCreateChannelDesc<T>()};  // or cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsigned)
-        int flags{0}; // or cudaArraySurfaceLoadStore
+        int flags{cudaArraySurfaceLoadStore}; // or 0
     };
 
     struct Impl {
@@ -25,7 +25,7 @@ class CudaArray {
 
         void copyIn(T const *_data) {
             cudaMemcpy3DParms copy3DParams{};
-            copy3DParams.srcPtr = make_cudaPitchedPtr((void *)_data, m_dim[0] * sizeof(T), m_dim[1], m_dim[2]);
+            copy3DParams.srcPtr = make_cudaPitchedPtr((void *)_data, m_dim[0] * sizeof(T), m_dim[0], m_dim[1]);
             copy3DParams.dstArray = m_cuArray;
             copy3DParams.extent = make_cudaExtent(m_dim[0], m_dim[1], m_dim[2]);
             copy3DParams.kind = cudaMemcpyHostToDevice;
@@ -35,7 +35,7 @@ class CudaArray {
         void copyOut(T *_data) {
             cudaMemcpy3DParms copy3DParams{};
             copy3DParams.srcArray = m_cuArray;
-            copy3DParams.dstPtr = make_cudaPitchedPtr((void *)_data, m_dim[0] * sizeof(T), m_dim[1], m_dim[2]);
+            copy3DParams.dstPtr = make_cudaPitchedPtr((void *)_data, m_dim[0] * sizeof(T), m_dim[0], m_dim[1]);
             copy3DParams.extent = make_cudaExtent(m_dim[0], m_dim[1], m_dim[2]);
             copy3DParams.kind = cudaMemcpyDeviceToHost;
             checkCudaErrors(cudaMemcpy3D(&copy3DParams));