Commit 5131d286 authored by BRAMAS Berenger's avatar BRAMAS Berenger
Browse files

Use an alignement of 8

parent 620bd90d
...@@ -15,7 +15,9 @@ static void FCudaCheckCore(cudaError_t code, const char *file, int line) { ...@@ -15,7 +15,9 @@ static void FCudaCheckCore(cudaError_t code, const char *file, int line) {
} }
#define FCudaCheck( test ) { FCudaCheckCore((test), __FILE__, __LINE__); } #define FCudaCheck( test ) { FCudaCheckCore((test), __FILE__, __LINE__); }
#define FCudaCheckAfterCall() { FCudaCheckCore((cudaGetLastError()), __FILE__, __LINE__); } #define FCudaCheckAfterCall() { FCudaCheckCore((cudaGetLastError()), __FILE__, __LINE__); }
#define FCudaAssertLF(ARGS) ARGS; #define FCudaAssertLF(ARGS) if(!(ARGS)){\
printf("Error line %d\n", __LINE__);\
}
#define FMGetOppositeNeighIndex(index) (27-(index)-1) #define FMGetOppositeNeighIndex(index) (27-(index)-1)
#define FMGetOppositeInterIndex(index) (343-(index)-1) #define FMGetOppositeInterIndex(index) (343-(index)-1)
......
...@@ -4,13 +4,15 @@ ...@@ -4,13 +4,15 @@
#include "FCudaGlobal.hpp" #include "FCudaGlobal.hpp"
#include "../FStarPUDefaultAlign.hpp"
/** /**
* @brief The FCudaGroupOfCells class manages the cells in block allocation. * @brief The FCudaGroupOfCells class manages the cells in block allocation.
*/ */
template <const size_t CellClassSize> template <const size_t CellClassSize>
class FCudaGroupOfCells { class FCudaGroupOfCells {
/** One header is allocated at the beginning of each block */ /** One header is allocated at the beginning of each block */
struct alignas(1) BlockHeader{ struct alignas(FStarPUDefaultAlign::StructAlign) BlockHeader{
MortonIndex startingIndex; MortonIndex startingIndex;
MortonIndex endingIndex; MortonIndex endingIndex;
int numberOfCellsInBlock; int numberOfCellsInBlock;
......
...@@ -3,12 +3,12 @@ ...@@ -3,12 +3,12 @@
#define FCUDAGROUPOFPARTICLES_HPP #define FCUDAGROUPOFPARTICLES_HPP
#include "FCudaGlobal.hpp" #include "FCudaGlobal.hpp"
#include "../FStarPUDefaultAlign.hpp"
template <unsigned NbAttributesPerParticle, class AttributeClass = FReal> template <unsigned NbAttributesPerParticle, class AttributeClass = FReal>
class FCudaGroupOfParticles { class FCudaGroupOfParticles {
/** One header is allocated at the beginning of each block */ /** One header is allocated at the beginning of each block */
struct alignas(1) BlockHeader{ struct alignas(FStarPUDefaultAlign::StructAlign) BlockHeader{
MortonIndex startingIndex; MortonIndex startingIndex;
MortonIndex endingIndex; MortonIndex endingIndex;
int numberOfLeavesInBlock; int numberOfLeavesInBlock;
...@@ -25,7 +25,7 @@ class FCudaGroupOfParticles { ...@@ -25,7 +25,7 @@ class FCudaGroupOfParticles {
}; };
/** Information about a leaf */ /** Information about a leaf */
struct alignas(1) LeafHeader { struct alignas(FStarPUDefaultAlign::StructAlign) LeafHeader {
int nbParticles; int nbParticles;
size_t offSet; size_t offSet;
}; };
......
...@@ -7,18 +7,18 @@ ...@@ -7,18 +7,18 @@
#include "../Utils/FAssert.hpp" #include "../Utils/FAssert.hpp"
#include "../Utils/FAlignedMemory.hpp" #include "../Utils/FAlignedMemory.hpp"
#include "../Containers/FTreeCoordinate.hpp" #include "../Containers/FTreeCoordinate.hpp"
#include "FStarPUDefaultAlign.hpp"
#include <list> #include <list>
#include <functional> #include <functional>
/** /**
* @brief The FGroupOfCells class manages the cells in block allocation. * @brief The FGroupOfCells class manages the cells in block allocation.
*/ */
template <class CellClass> template <class CellClass>
class FGroupOfCells { class FGroupOfCells {
/** One header is allocated at the beginning of each block */ /** One header is allocated at the beginning of each block */
struct alignas(1) BlockHeader{ struct alignas(FStarPUDefaultAlign::StructAlign) BlockHeader{
MortonIndex startingIndex; MortonIndex startingIndex;
MortonIndex endingIndex; MortonIndex endingIndex;
int numberOfCellsInBlock; int numberOfCellsInBlock;
......
...@@ -8,6 +8,7 @@ ...@@ -8,6 +8,7 @@
#include "../Utils/FAssert.hpp" #include "../Utils/FAssert.hpp"
#include "../Containers/FTreeCoordinate.hpp" #include "../Containers/FTreeCoordinate.hpp"
#include "../Utils/FAlignedMemory.hpp" #include "../Utils/FAlignedMemory.hpp"
#include "FStarPUDefaultAlign.hpp"
#include <list> #include <list>
#include <functional> #include <functional>
...@@ -19,7 +20,7 @@ ...@@ -19,7 +20,7 @@
template <unsigned NbAttributesPerParticle, class AttributeClass = FReal> template <unsigned NbAttributesPerParticle, class AttributeClass = FReal>
class FGroupOfParticles { class FGroupOfParticles {
/** One header is allocated at the beginning of each block */ /** One header is allocated at the beginning of each block */
struct alignas(1) BlockHeader{ struct alignas(FStarPUDefaultAlign::StructAlign) BlockHeader{
MortonIndex startingIndex; MortonIndex startingIndex;
MortonIndex endingIndex; MortonIndex endingIndex;
int numberOfLeavesInBlock; int numberOfLeavesInBlock;
...@@ -36,7 +37,7 @@ class FGroupOfParticles { ...@@ -36,7 +37,7 @@ class FGroupOfParticles {
}; };
/** Information about a leaf */ /** Information about a leaf */
struct alignas(1) LeafHeader { struct alignas(FStarPUDefaultAlign::StructAlign) LeafHeader {
int nbParticles; int nbParticles;
size_t offSet; size_t offSet;
}; };
......
...@@ -3,7 +3,9 @@ ...@@ -3,7 +3,9 @@
#include "../Utils/FGlobal.hpp" #include "../Utils/FGlobal.hpp"
struct alignas(1) OutOfBlockInteraction{ #include "FStarPUDefaultAlign.hpp"
struct alignas(FStarPUDefaultAlign::StructAlign) OutOfBlockInteraction{
MortonIndex outIndex; MortonIndex outIndex;
MortonIndex insideIndex; MortonIndex insideIndex;
int outPosition; int outPosition;
......
...@@ -3,17 +3,19 @@ ...@@ -3,17 +3,19 @@
typedef long long int MortonIndex; typedef long long int MortonIndex;
#define DefaultStructAlign ___DefaultStructAlign___
typedef struct OutOfBlockInteraction{ typedef struct OutOfBlockInteraction{
MortonIndex outIndex; MortonIndex outIndex;
MortonIndex insideIndex; MortonIndex insideIndex;
int outPosition; int outPosition;
} __attribute__ ((aligned (1))); } __attribute__ ((aligned (DefaultStructAlign)));
struct Uptr9{ struct Uptr9{
__global unsigned char* ptrs[9]; __global unsigned char* ptrs[9];
} __attribute__ ((aligned (1))); } __attribute__ ((aligned (DefaultStructAlign)));
struct size_t9{ struct size_t9{
size_t v[9]; size_t v[9];
}__attribute__ ((aligned (1))); }__attribute__ ((aligned (DefaultStructAlign)));
__kernel void FOpenCL__bottomPassPerform(__global unsigned char* leafCellsPtr, size_t leafCellsSize, __kernel void FOpenCL__bottomPassPerform(__global unsigned char* leafCellsPtr, size_t leafCellsSize,
__global unsigned char* containersPtr, size_t containersSize, __global unsigned char* containersPtr, size_t containersSize,
__global void* userkernel ){ __global void* userkernel ){
......
...@@ -4,6 +4,8 @@ ...@@ -4,6 +4,8 @@
// Return the same thing as FEmptyKernel.cl // Return the same thing as FEmptyKernel.cl
#include "../FStarPUDefaultAlign.hpp"
class FEmptyOpenCLCode{ class FEmptyOpenCLCode{
public: public:
...@@ -13,18 +15,18 @@ public: ...@@ -13,18 +15,18 @@ public:
const char* getKernelCode(const int /*inDevId*/){ const char* getKernelCode(const int /*inDevId*/){
const char* kernelcode = const char* kernelcode =
"typedef long long int MortonIndex; \ "typedef long long int MortonIndex; \
\ #define DefaultStructAlign " FStarPUDefaultAlignStr "\
typedef struct OutOfBlockInteraction{ \ typedef struct OutOfBlockInteraction{ \
MortonIndex outIndex; \ MortonIndex outIndex; \
MortonIndex insideIndex; \ MortonIndex insideIndex; \
int outPosition; \ int outPosition; \
} __attribute__ ((aligned (1))); \ } __attribute__ ((aligned (DefaultStructAlign))); \
struct Uptr9{ \ struct Uptr9{ \
__global unsigned char* ptrs[9]; \ __global unsigned char* ptrs[9]; \
} __attribute__ ((aligned (1))); \ } __attribute__ ((aligned (DefaultStructAlign))); \
struct size_t9{ \ struct size_t9{ \
size_t v[9]; \ size_t v[9]; \
}__attribute__ ((aligned (1))); \ }__attribute__ ((aligned (DefaultStructAlign))); \
__kernel void FOpenCL__bottomPassPerform(__global unsigned char* leafCellsPtr, size_t leafCellsSize, \ __kernel void FOpenCL__bottomPassPerform(__global unsigned char* leafCellsPtr, size_t leafCellsSize, \
__global unsigned char* containersPtr, size_t containersSize, \ __global unsigned char* containersPtr, size_t containersSize, \
__global void* userkernel ){ \ __global void* userkernel ){ \
......
...@@ -16,17 +16,19 @@ ...@@ -16,17 +16,19 @@
#include "FEmptyOpenCLCode.hpp" #include "FEmptyOpenCLCode.hpp"
#include "../FStarPUDefaultAlign.hpp"
#include <starpu.h> #include <starpu.h>
template <class OriginalKernelClass, class KernelFilenameClass = FEmptyOpenCLCode> template <class OriginalKernelClass, class KernelFilenameClass = FEmptyOpenCLCode>
class FOpenCLDeviceWrapper { class FOpenCLDeviceWrapper {
protected: protected:
struct alignas(1) Uptr9{ struct alignas(FStarPUDefaultAlign::StructAlign) Uptr9{
cl_mem ptrs[9]; cl_mem ptrs[9];
}; };
struct alignas(1) size_t9{ struct alignas(FStarPUDefaultAlign::StructAlign) size_t9{
size_t v[9]; size_t v[9];
}; };
...@@ -175,8 +177,10 @@ public: ...@@ -175,8 +177,10 @@ public:
FAssertLF(handle); FAssertLF(handle);
//starpu_data_acquire(handle, STARPU_RW); //starpu_data_acquire(handle, STARPU_RW);
//starpu_data_release(handle); //starpu_data_release(handle);
//uintptr_t dataTest = ((starpu_variable_interface2*)(((_starpu_data_replicate2*)(handle->per_node))->data_interface))->ptr;//((starpu_variable_interface*)(handle->per_node[0]->data_interface))->ptr;
void* data = starpu_data_get_local_ptr( handle); void* data = starpu_data_get_local_ptr( handle);
FAssertLF(data && starpu_data_get_size(handle) == leafCellsSize); FAssertLF(data && starpu_data_get_size(handle) == leafCellsSize);
std::cout << "currentCells data " << data << "\n";
//const int errcode_ret = clEnqueueReadBuffer(queue_bottomPassPerform, leafCellsPtr, //const int errcode_ret = clEnqueueReadBuffer(queue_bottomPassPerform, leafCellsPtr,
// CL_TRUE, 0, leafCellsSize, data, 0, NULL, NULL); // CL_TRUE, 0, leafCellsSize, data, 0, NULL, NULL);
//FAssertLF(errcode_ret == CL_SUCCESS, "OpenCL error code " , errcode_ret); //FAssertLF(errcode_ret == CL_SUCCESS, "OpenCL error code " , errcode_ret);
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment